OpenMP: Add support for 'close' in map clause

Message ID f541a0ac-90b7-0985-0cc2-34ae6d9c7dfa@codesourcery.com
State New
Headers show
Series
  • OpenMP: Add support for 'close' in map clause
Related show

Commit Message

Marcel Vollweiler May 10, 2021, 2:11 p.m.
Hi,

This patch adds handling for the map-type-modifier 'close' in the map
clause that was introduced with OpenMP 5.0: "The close map-type-modifier
is a hint to the runtime to allocate memory close to the target device."

In OpenMP 5.0 'close' can be used beside/together with 'always' in a
list of map-type-modifiers.

With this patch, 'close' will be parsed and ignored for C and C++. A
patch for Fortran will be provided separately.

Marcel
-----------------
Mentor Graphics (Deutschland) GmbH, Arnulfstrasse 201, 80634 München Registergericht München HRB 106955, Geschäftsführer: Thomas Heurung, Frank Thürauf
OpenMP: Add support for 'close' in map clause

gcc/c/ChangeLog:

	* c-parser.c (c_parser_omp_clause_map): Support map-type-modifier 
	'close'.

gcc/cp/ChangeLog:

	* parser.c (cp_parser_omp_clause_map): Support map-type-modifier 
	'close'.

gcc/testsuite/ChangeLog:

	* c-c++-common/gomp/map-6.c: New test.
	* c-c++-common/gomp/map-7.c: New test.

Comments

Andre Vehreschild via Gcc-patches May 10, 2021, 6:34 p.m. | #1
On Mon, May 10, 2021 at 04:11:39PM +0200, Marcel Vollweiler wrote:
> @@ -15660,37 +15665,54 @@ c_parser_omp_clause_map (c_parser *parser, tree list)

>    if (!parens.require_open (parser))

>      return list;

>  

> -  if (c_parser_next_token_is (parser, CPP_NAME))

> +  int always = 0;

> +  int close = 0;

> +  int pos = 1;

> +  while (c_parser_peek_nth_token_raw (parser, pos)->type == CPP_NAME)


Nice, totally missed that Joseph has added this.

>      {

> -      c_token *tok = c_parser_peek_token (parser);

> +      c_token *tok = c_parser_peek_nth_token_raw (parser, pos);

>        const char *p = IDENTIFIER_POINTER (tok->value);

> -      always_id_kind = tok->id_kind;

> -      always_loc = tok->location;

> -      always_id = tok->value;

>        if (strcmp ("always", p) == 0)

>  	{

> -	  c_token *sectok = c_parser_peek_2nd_token (parser);

> -	  if (sectok->type == CPP_COMMA)

> +	  if (always)

>  	    {

> -	      c_parser_consume_token (parser);

> -	      c_parser_consume_token (parser);

> -	      always = 2;

> +	      c_parser_error (parser, "expected modifier %<always%> only once");


The usual wording would be
"too many %<always%> modifiers"

> +	      parens.skip_until_found_close (parser);

> +	      return list;

> +	    }

> +

> +	  always_id_kind = tok->id_kind;

> +	  always_loc = tok->location;

> +	  always_id = tok->value;


But you don't need any of the always_{id_kind,loc,id} variables anymore,
so they should be removed and everything that touches them too.

> +

> +	  always++;

> +	}

> +      else if (strcmp ("close", p) == 0)

> +	{

> +	  if (close)

> +	    {

> +	      c_parser_error (parser, "expected modifier %<close%> only once");


Similarly.

> +	      parens.skip_until_found_close (parser);

> +	      return list;

>  	    }

> -	  else if (sectok->type == CPP_NAME)

> +

> +	  close++;

> +	}

> +      else if (c_parser_peek_nth_token_raw (parser, pos + 1)->type == CPP_COLON)


IMHO you should at least check that tok->type == CPP_NAME before
checking pos + 1 token's type, you don't want to skip over CPP_EOF,
CPP_PRAGMA_EOF, or even CPP_CLOSE_PAREN etc.
Perhaps by adding
      if (tok->type != CPP_NAME)
	break;
right after c_token *tok = c_parser_peek_nth_token_raw (parser, pos); ?

> +	{

> +	  for (int i = 1; i < pos; ++i)

>  	    {

> -	      p = IDENTIFIER_POINTER (sectok->value);

> -	      if (strcmp ("alloc", p) == 0

> -		  || strcmp ("to", p) == 0

> -		  || strcmp ("from", p) == 0

> -		  || strcmp ("tofrom", p) == 0

> -		  || strcmp ("release", p) == 0

> -		  || strcmp ("delete", p) == 0)

> -		{

> -		  c_parser_consume_token (parser);

> -		  always = 1;

> -		}

> +	      c_parser_peek_token(parser);


Formatting, space before (

> +	      c_parser_consume_token (parser);

>  	    }

> +	  break;


And, IMHO something should clear always and close (btw, might be better
to use close_modifier as variable name and for consistency always_modifier)
unless we reach the CPP_COLON case.

Because we don't want
  map (always, close)
to imply
  map (always, close, tofrom: always, close)
but
  map (tofrom: always, close)
and my reading of your changes suggests that we actually use the
*_ALWAYS* kinds in that case.

> +	      cp_parser_error (parser,

> +			       "expected modifier %<always%> only once");


See above.

> +	      cp_parser_skip_to_closing_parenthesis (parser,

> +						     /*recovering=*/true,

> +						     /*or_comma=*/false,

> +						     /*consume_paren=*/true);

> +	      return list;

> +	    }

> +

> +	  always = true;

> +	}

> +      else if (strcmp ("close", p) == 0)

> +	{

> +	  if (close)

> +	    {

> +	      cp_parser_error (parser,

> +			       "expected modifier %<close%> only once");


Likewise.

> +      else if (cp_lexer_peek_nth_token (parser->lexer, pos + 1)->type

> +	       == CPP_COLON)

> +	{

> +	  for (int i = 1; i < pos; ++i)

> +	    cp_lexer_consume_token (parser->lexer);

> +	  break;

> +	}

> +      else

> +	break;

> +

> +      if (cp_lexer_peek_nth_token (parser->lexer, pos + 1)->type == CPP_COMMA)

> +	pos++;

> +      pos++;

>      }


Again, I don't see anything that would clear always/close if it didn't reach
the CPP_COLON case.

And it should be covered in the testcase.

	Jakub
Marcel Vollweiler May 11, 2021, 2:27 p.m. | #2
Am 10.05.2021 um 20:34 schrieb Jakub Jelinek:
> On Mon, May 10, 2021 at 04:11:39PM +0200, Marcel Vollweiler wrote:

>> @@ -15660,37 +15665,54 @@ c_parser_omp_clause_map (c_parser *parser, tree list)

>>     if (!parens.require_open (parser))

>>       return list;

>>

>> -  if (c_parser_next_token_is (parser, CPP_NAME))

>> +  int always = 0;

>> +  int close = 0;

>> +  int pos = 1;

>> +  while (c_parser_peek_nth_token_raw (parser, pos)->type == CPP_NAME)

>

> Nice, totally missed that Joseph has added this.

>

>>       {

>> -      c_token *tok = c_parser_peek_token (parser);

>> +      c_token *tok = c_parser_peek_nth_token_raw (parser, pos);

>>         const char *p = IDENTIFIER_POINTER (tok->value);

>> -      always_id_kind = tok->id_kind;

>> -      always_loc = tok->location;

>> -      always_id = tok->value;

>>         if (strcmp ("always", p) == 0)

>>      {

>> -      c_token *sectok = c_parser_peek_2nd_token (parser);

>> -      if (sectok->type == CPP_COMMA)

>> +      if (always)

>>          {

>> -          c_parser_consume_token (parser);

>> -          c_parser_consume_token (parser);

>> -          always = 2;

>> +          c_parser_error (parser, "expected modifier %<always%> only once");

>

> The usual wording would be

> "too many %<always%> modifiers"

>


Changed for 'always' and 'close' for C and C++.

>> +          parens.skip_until_found_close (parser);

>> +          return list;

>> +        }

>> +

>> +      always_id_kind = tok->id_kind;

>> +      always_loc = tok->location;

>> +      always_id = tok->value;

>

> But you don't need any of the always_{id_kind,loc,id} variables anymore,

> so they should be removed and everything that touches them too.

>


That's true. I removed them.

>> +

>> +      always++;

>> +    }

>> +      else if (strcmp ("close", p) == 0)

>> +    {

>> +      if (close)

>> +        {

>> +          c_parser_error (parser, "expected modifier %<close%> only once");

>

> Similarly.

>

>> +          parens.skip_until_found_close (parser);

>> +          return list;

>>          }

>> -      else if (sectok->type == CPP_NAME)

>> +

>> +      close++;

>> +    }

>> +      else if (c_parser_peek_nth_token_raw (parser, pos + 1)->type == CPP_COLON)

>

> IMHO you should at least check that tok->type == CPP_NAME before

> checking pos + 1 token's type, you don't want to skip over CPP_EOF,

> CPP_PRAGMA_EOF, or even CPP_CLOSE_PAREN etc.

> Perhaps by adding

>        if (tok->type != CPP_NAME)

>       break;

> right after c_token *tok = c_parser_peek_nth_token_raw (parser, pos); ?


The check of the token's type at position 'pos' is done in the condition
of the while loop, that means
    'c_parser_peek_nth_token_raw (parser, pos + 1)->type == CPP_COLON'
is only reached when
    'c_parser_peek_nth_token_raw (parser, pos)->type == CPP_NAME'
holds (since 'pos' is not changed in between).

>

>> +    {

>> +      for (int i = 1; i < pos; ++i)

>>          {

>> -          p = IDENTIFIER_POINTER (sectok->value);

>> -          if (strcmp ("alloc", p) == 0

>> -              || strcmp ("to", p) == 0

>> -              || strcmp ("from", p) == 0

>> -              || strcmp ("tofrom", p) == 0

>> -              || strcmp ("release", p) == 0

>> -              || strcmp ("delete", p) == 0)

>> -            {

>> -              c_parser_consume_token (parser);

>> -              always = 1;

>> -            }

>> +          c_parser_peek_token(parser);

>

> Formatting, space before (

>


Corrected.

>> +          c_parser_consume_token (parser);

>>          }

>> +      break;

>

> And, IMHO something should clear always and close (btw, might be better

> to use close_modifier as variable name and for consistency always_modifier)

> unless we reach the CPP_COLON case.

>


Good point, I agree with both. Cleared and renamed :)

> Because we don't want

>    map (always, close)

> to imply

>    map (always, close, tofrom: always, close)

> but

>    map (tofrom: always, close)

> and my reading of your changes suggests that we actually use the

> *_ALWAYS* kinds in that case.

>

>> +          cp_parser_error (parser,

>> +                           "expected modifier %<always%> only once");

>

> See above.

>

>> +          cp_parser_skip_to_closing_parenthesis (parser,

>> +                                                 /*recovering=*/true,

>> +                                                 /*or_comma=*/false,

>> +                                                 /*consume_paren=*/true);

>> +          return list;

>> +        }

>> +

>> +      always = true;

>> +    }

>> +      else if (strcmp ("close", p) == 0)

>> +    {

>> +      if (close)

>> +        {

>> +          cp_parser_error (parser,

>> +                           "expected modifier %<close%> only once");

>

> Likewise.

>

>> +      else if (cp_lexer_peek_nth_token (parser->lexer, pos + 1)->type

>> +           == CPP_COLON)

>> +    {

>> +      for (int i = 1; i < pos; ++i)

>> +        cp_lexer_consume_token (parser->lexer);

>> +      break;

>> +    }

>> +      else

>> +    break;

>> +

>> +      if (cp_lexer_peek_nth_token (parser->lexer, pos + 1)->type == CPP_COMMA)

>> +    pos++;

>> +      pos++;

>>       }

>

> Again, I don't see anything that would clear always/close if it didn't reach

> the CPP_COLON case.

>

> And it should be covered in the testcase.


I added a test case similar to your example above:
        '#pragma omp target map (always, close, to: always, close, b7)'

>

>       Jakub

>


The new version of the patch including the changes is attached.

Thanks,

Marcel
-----------------
Mentor Graphics (Deutschland) GmbH, Arnulfstrasse 201, 80634 München Registergericht München HRB 106955, Geschäftsführer: Thomas Heurung, Frank Thürauf
OpenMP: Add support for 'close' in map clause

gcc/c/ChangeLog:

	* c-parser.c (c_parser_omp_clause_map): Support map-type-modifier 
	'close'.

gcc/cp/ChangeLog:

	* parser.c (cp_parser_omp_clause_map): Support map-type-modifier 
	'close'.

gcc/testsuite/ChangeLog:

	* c-c++-common/gomp/map-6.c: New test.
	* c-c++-common/gomp/map-7.c: New test.

diff --git a/gcc/c/c-parser.c b/gcc/c/c-parser.c
index 5cdeb21..310df89 100644
--- a/gcc/c/c-parser.c
+++ b/gcc/c/c-parser.c
@@ -15643,54 +15643,73 @@ c_parser_omp_clause_depend (c_parser *parser, tree list)
    map-kind:
      alloc | to | from | tofrom | release | delete
 
-   map ( always [,] map-kind: variable-list ) */
+   map ( always [,] map-kind: variable-list )
+
+   OpenMP 5.0:
+   map ( [map-type-modifier[,] ...] map-kind: variable-list )
+
+   map-type-modifier:
+     always | close */
 
 static tree
 c_parser_omp_clause_map (c_parser *parser, tree list)
 {
   location_t clause_loc = c_parser_peek_token (parser)->location;
   enum gomp_map_kind kind = GOMP_MAP_TOFROM;
-  int always = 0;
-  enum c_id_kind always_id_kind = C_ID_NONE;
-  location_t always_loc = UNKNOWN_LOCATION;
-  tree always_id = NULL_TREE;
   tree nl, c;
 
   matching_parens parens;
   if (!parens.require_open (parser))
     return list;
 
-  if (c_parser_next_token_is (parser, CPP_NAME))
+  int always_modifier = 0;
+  int close_modifier = 0;
+  int pos = 1;
+  while (c_parser_peek_nth_token_raw (parser, pos)->type == CPP_NAME)
     {
-      c_token *tok = c_parser_peek_token (parser);
+      c_token *tok = c_parser_peek_nth_token_raw (parser, pos);
       const char *p = IDENTIFIER_POINTER (tok->value);
-      always_id_kind = tok->id_kind;
-      always_loc = tok->location;
-      always_id = tok->value;
       if (strcmp ("always", p) == 0)
 	{
-	  c_token *sectok = c_parser_peek_2nd_token (parser);
-	  if (sectok->type == CPP_COMMA)
+	  if (always_modifier)
 	    {
-	      c_parser_consume_token (parser);
-	      c_parser_consume_token (parser);
-	      always = 2;
+	      c_parser_error (parser, "too many %<always%> modifiers");
+	      parens.skip_until_found_close (parser);
+	      return list;
+	    }
+
+	  always_modifier++;
+	}
+      else if (strcmp ("close", p) == 0)
+	{
+	  if (close_modifier)
+	    {
+	      c_parser_error (parser, "too many %<close%> modifiers");
+	      parens.skip_until_found_close (parser);
+	      return list;
 	    }
-	  else if (sectok->type == CPP_NAME)
+
+	  close_modifier++;
+	}
+      else if (c_parser_peek_nth_token_raw (parser, pos + 1)->type == CPP_COLON)
+	{
+	  for (int i = 1; i < pos; ++i)
 	    {
-	      p = IDENTIFIER_POINTER (sectok->value);
-	      if (strcmp ("alloc", p) == 0
-		  || strcmp ("to", p) == 0
-		  || strcmp ("from", p) == 0
-		  || strcmp ("tofrom", p) == 0
-		  || strcmp ("release", p) == 0
-		  || strcmp ("delete", p) == 0)
-		{
-		  c_parser_consume_token (parser);
-		  always = 1;
-		}
+	      c_parser_peek_token (parser);
+	      c_parser_consume_token (parser);
 	    }
+	  break;
 	}
+      else
+	{
+	  always_modifier = 0;
+	  close_modifier = 0;
+	  break;
+	}
+
+      if (c_parser_peek_nth_token_raw (parser, pos + 1)->type == CPP_COMMA)
+	pos++;
+      pos++;
     }
 
   if (c_parser_next_token_is (parser, CPP_NAME)
@@ -15700,11 +15719,11 @@ c_parser_omp_clause_map (c_parser *parser, tree list)
       if (strcmp ("alloc", p) == 0)
 	kind = GOMP_MAP_ALLOC;
       else if (strcmp ("to", p) == 0)
-	kind = always ? GOMP_MAP_ALWAYS_TO : GOMP_MAP_TO;
+	kind = always_modifier ? GOMP_MAP_ALWAYS_TO : GOMP_MAP_TO;
       else if (strcmp ("from", p) == 0)
-	kind = always ? GOMP_MAP_ALWAYS_FROM : GOMP_MAP_FROM;
+	kind = always_modifier ? GOMP_MAP_ALWAYS_FROM : GOMP_MAP_FROM;
       else if (strcmp ("tofrom", p) == 0)
-	kind = always ? GOMP_MAP_ALWAYS_TOFROM : GOMP_MAP_TOFROM;
+	kind = always_modifier ? GOMP_MAP_ALWAYS_TOFROM : GOMP_MAP_TOFROM;
       else if (strcmp ("release", p) == 0)
 	kind = GOMP_MAP_RELEASE;
       else if (strcmp ("delete", p) == 0)
@@ -15719,35 +15738,6 @@ c_parser_omp_clause_map (c_parser *parser, tree list)
       c_parser_consume_token (parser);
       c_parser_consume_token (parser);
     }
-  else if (always)
-    {
-      if (always_id_kind != C_ID_ID)
-	{
-	  c_parser_error (parser, "expected identifier");
-	  parens.skip_until_found_close (parser);
-	  return list;
-	}
-
-      tree t = lookup_name (always_id);
-      if (t == NULL_TREE)
-	{
-	  undeclared_variable (always_loc, always_id);
-	  t = error_mark_node;
-	}
-      if (t != error_mark_node)
-	{
-	  tree u = build_omp_clause (clause_loc, OMP_CLAUSE_MAP);
-	  OMP_CLAUSE_DECL (u) = t;
-	  OMP_CLAUSE_CHAIN (u) = list;
-	  OMP_CLAUSE_SET_MAP_KIND (u, kind);
-	  list = u;
-	}
-      if (always == 1)
-	{
-	  parens.skip_until_found_close (parser);
-	  return list;
-	}
-    }
 
   nl = c_parser_omp_variable_list (parser, clause_loc, OMP_CLAUSE_MAP, list);
 
diff --git a/gcc/cp/parser.c b/gcc/cp/parser.c
index 99eccf0..55a59ed 100644
--- a/gcc/cp/parser.c
+++ b/gcc/cp/parser.c
@@ -37840,40 +37840,77 @@ cp_parser_omp_clause_depend (cp_parser *parser, tree list, location_t loc)
    map-kind:
      alloc | to | from | tofrom | release | delete
 
-   map ( always [,] map-kind: variable-list ) */
+   map ( always [,] map-kind: variable-list )
+
+   OpenMP 5.0:
+   map ( [map-type-modifier[,] ...] map-kind: variable-list )
+
+   map-type-modifier:
+     always | close */
 
 static tree
 cp_parser_omp_clause_map (cp_parser *parser, tree list)
 {
   tree nlist, c;
   enum gomp_map_kind kind = GOMP_MAP_TOFROM;
-  bool always = false;
+  bool always_modifier = false;
+  bool close_modifier = false;
+  int pos = 1;
 
   if (!cp_parser_require (parser, CPP_OPEN_PAREN, RT_OPEN_PAREN))
     return list;
 
-  if (cp_lexer_next_token_is (parser->lexer, CPP_NAME))
+  while (cp_lexer_peek_nth_token (parser->lexer, pos)->type == CPP_NAME
+	 || cp_lexer_peek_nth_token (parser->lexer, pos)->keyword == RID_DELETE)
     {
-      tree id = cp_lexer_peek_token (parser->lexer)->u.value;
+      tree id = cp_lexer_peek_nth_token (parser->lexer, pos)->u.value;
       const char *p = IDENTIFIER_POINTER (id);
 
       if (strcmp ("always", p) == 0)
 	{
-	  int nth = 2;
-	  if (cp_lexer_peek_nth_token (parser->lexer, 2)->type == CPP_COMMA)
-	    nth++;
-	  if ((cp_lexer_peek_nth_token (parser->lexer, nth)->type == CPP_NAME
-	       || (cp_lexer_peek_nth_token (parser->lexer, nth)->keyword
-		   == RID_DELETE))
-	      && (cp_lexer_peek_nth_token (parser->lexer, nth + 1)->type
-		  == CPP_COLON))
+	  if (always_modifier)
 	    {
-	      always = true;
-	      cp_lexer_consume_token (parser->lexer);
-	      if (nth == 3)
-		cp_lexer_consume_token (parser->lexer);
+	      cp_parser_error (parser, "too many %<always%> modifiers");
+	      cp_parser_skip_to_closing_parenthesis (parser,
+						     /*recovering=*/true,
+						     /*or_comma=*/false,
+						     /*consume_paren=*/true);
+	      return list;
+	    }
+
+	  always_modifier = true;
+	}
+      else if (strcmp ("close", p) == 0)
+	{
+	  if (close_modifier)
+	    {
+	      cp_parser_error (parser, "too many %<close%> modifiers");
+	      cp_parser_skip_to_closing_parenthesis (parser,
+						     /*recovering=*/true,
+						     /*or_comma=*/false,
+						     /*consume_paren=*/true);
+	      return list;
 	    }
+
+	  close_modifier = true;
 	}
+      else if (cp_lexer_peek_nth_token (parser->lexer, pos + 1)->type
+	       == CPP_COLON)
+	{
+	  for (int i = 1; i < pos; ++i)
+	    cp_lexer_consume_token (parser->lexer);
+	  break;
+	}
+      else
+	{
+	  always_modifier = 0;
+	  close_modifier = 0;
+	  break;
+	}
+
+      if (cp_lexer_peek_nth_token (parser->lexer, pos + 1)->type == CPP_COMMA)
+	pos++;
+      pos++;
     }
 
   if (cp_lexer_next_token_is (parser->lexer, CPP_NAME)
@@ -37885,11 +37922,11 @@ cp_parser_omp_clause_map (cp_parser *parser, tree list)
       if (strcmp ("alloc", p) == 0)
 	kind = GOMP_MAP_ALLOC;
       else if (strcmp ("to", p) == 0)
-	kind = always ? GOMP_MAP_ALWAYS_TO : GOMP_MAP_TO;
+	kind = always_modifier ? GOMP_MAP_ALWAYS_TO : GOMP_MAP_TO;
       else if (strcmp ("from", p) == 0)
-	kind = always ? GOMP_MAP_ALWAYS_FROM : GOMP_MAP_FROM;
+	kind = always_modifier ? GOMP_MAP_ALWAYS_FROM : GOMP_MAP_FROM;
       else if (strcmp ("tofrom", p) == 0)
-	kind = always ? GOMP_MAP_ALWAYS_TOFROM : GOMP_MAP_TOFROM;
+	kind = always_modifier ? GOMP_MAP_ALWAYS_TOFROM : GOMP_MAP_TOFROM;
       else if (strcmp ("release", p) == 0)
 	kind = GOMP_MAP_RELEASE;
       else
diff --git a/gcc/testsuite/c-c++-common/gomp/map-6.c b/gcc/testsuite/c-c++-common/gomp/map-6.c
new file mode 100644
index 0000000..b524e30
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/gomp/map-6.c
@@ -0,0 +1,126 @@
+/* { dg-do compile } */
+/* { dg-additional-options "-fdump-tree-original" } */
+
+void
+foo (void)
+{
+  /* Test to ensure that the close modifier is parsed and ignored in map clauses. */
+  int a, b, b1, b2, b3, b4, b5, b6, b7;
+
+  #pragma omp target map (a)
+  ;
+
+  #pragma omp target map (to:a)
+  ;
+
+  #pragma omp target map (a to: b) /* { dg-error "expected '\\)' before 'to'" } */
+  ;
+
+  #pragma omp target map (close a) /* { dg-error "'close' undeclared" "" { target c } } */ 
+  /* { dg-error "'close' has not been declared" "" { target c++ } .-1 } */ 
+  /* { dg-error "expected '\\)' before 'a'" "" { target *-*-* } .-2 } */
+  ;
+
+  #pragma omp target map (always a) /* { dg-error "'always' undeclared" "" { target c } } */
+  /* { dg-error "'always' has not been declared" "" { target c++ } .-1 } */ 
+  /* { dg-error "expected '\\)' before 'a'" "" { target *-*-* } .-2 } */
+  ;
+
+  #pragma omp target map (close to:a)
+  ;
+
+  #pragma omp target map (close, to:a)
+  ;
+
+  #pragma omp target map (close delete:a) /* { dg-error "'#pragma omp target' with map-type other than 'to', 'from', 'tofrom' or 'alloc' on 'map' clause" } */
+  ;
+
+  #pragma omp target map (close always to:b1)
+  ;
+
+  #pragma omp target map (close, always to:b2)
+  ;
+
+  #pragma omp target map (close, always, to:b3)
+  ;
+
+  #pragma omp target map (always close to:b4)
+  ;
+
+  #pragma omp target map (always, close to:b5)
+  ;
+
+  #pragma omp target map (always, close, to:b6)
+  ;
+
+  #pragma omp target map (always, always, to:a) /* { dg-error "too many 'always' modifiers" } */
+  ;
+
+  #pragma omp target map (always always, to:a) /* { dg-error "too many 'always' modifiers" } */
+  ;
+
+  #pragma omp target map (always, always to:a) /* { dg-error "too many 'always' modifiers" } */
+  ;
+
+  #pragma omp target map (always always to:a) /* { dg-error "too many 'always' modifiers" } */
+  ;
+
+  #pragma omp target map (close, close, to:a) /* { dg-error "too many 'close' modifiers" } */
+  ;
+
+  #pragma omp target map (close close, to:a) /* { dg-error "too many 'close' modifiers" } */
+  ;
+
+  #pragma omp target map (close, close to:a) /* { dg-error "too many 'close' modifiers" } */
+  ;
+
+  #pragma omp target map (close close to:a) /* { dg-error "too many 'close' modifiers" } */
+  ;
+
+  #pragma omp target map (always to : a) map (close to : b)
+  ;
+
+  int close = 0;
+  #pragma omp target map (close) 
+  ;
+
+  #pragma omp target map (close a) /* { dg-error "expected '\\)' before 'a'" } */ 
+  ;
+
+  int always = 0;
+  #pragma omp target map (always)
+  ;
+
+  #pragma omp target map (always a) /* { dg-error "expected '\\)' before 'a'" } */
+  ;
+
+  #pragma omp target map (always, close)
+  ;
+
+  #pragma omp target map (always, close, to: always, close, b7)
+  ;
+
+  int to = 0;
+  #pragma omp target map (always, close, to)
+  ;
+
+  #pragma omp target map (to, always, close)
+    {
+      to = always = close = 1;
+    }
+  if (to != 1 || always != 1 || close != 1)
+    __builtin_abort ();
+  ;
+}
+
+/* { dg-final { scan-tree-dump-not "map\\(\[^\n\r)]*close\[^\n\r)]*to:" "original" } } */
+
+/* { dg-final { scan-tree-dump-times "pragma omp target map\\(always,to:" 7 "original" } } */
+
+/* { dg-final { scan-tree-dump "pragma omp target map\\(always,to:b1" "original" } } */
+/* { dg-final { scan-tree-dump "pragma omp target map\\(always,to:b2" "original" } } */
+/* { dg-final { scan-tree-dump "pragma omp target map\\(always,to:b3" "original" } } */
+/* { dg-final { scan-tree-dump "pragma omp target map\\(always,to:b4" "original" } } */
+/* { dg-final { scan-tree-dump "pragma omp target map\\(always,to:b5" "original" } } */
+/* { dg-final { scan-tree-dump "pragma omp target map\\(always,to:b6" "original" } } */
+/* { dg-final { scan-tree-dump "pragma omp target map\\(always,to:b7\\) map\\(always,to:close\\) map\\(always,to:always\\)" "original" } } */
diff --git a/gcc/testsuite/c-c++-common/gomp/map-7.c b/gcc/testsuite/c-c++-common/gomp/map-7.c
new file mode 100644
index 0000000..3f1e972
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/gomp/map-7.c
@@ -0,0 +1,20 @@
+/* { dg-do compile } */
+
+void
+foo (void)
+{
+  /* Test to ensure that the close modifier is parsed and ignored in map clauses. */
+
+  #define N 1024
+  int always[N];
+  int close;
+
+  #pragma omp target map(always[:N]) 
+  ;
+
+  #pragma omp target map(close, always[:N]) 
+  ;
+
+  #pragma omp target map(always[:N], close) 
+  ;
+}
Andre Vehreschild via Gcc-patches May 11, 2021, 3:20 p.m. | #3
On Tue, May 11, 2021 at 04:27:55PM +0200, Marcel Vollweiler wrote:
> > The usual wording would be

> > "too many %<always%> modifiers"

> > 

> 

> Changed for 'always' and 'close' for C and C++.


One extra thing, sorry, forgot to mention, for the translators it might be
better to use "too many %qs modifiers", "always" (or, "close").
That way they can translate it just once instead of twice.

> > IMHO you should at least check that tok->type == CPP_NAME before

> > checking pos + 1 token's type, you don't want to skip over CPP_EOF,

> > CPP_PRAGMA_EOF, or even CPP_CLOSE_PAREN etc.

> > Perhaps by adding

> >        if (tok->type != CPP_NAME)

> >       break;

> > right after c_token *tok = c_parser_peek_nth_token_raw (parser, pos); ?

> 

> The check of the token's type at position 'pos' is done in the condition

> of the while loop, that means

>    'c_parser_peek_nth_token_raw (parser, pos + 1)->type == CPP_COLON'

> is only reached when

>    'c_parser_peek_nth_token_raw (parser, pos)->type == CPP_NAME'

> holds (since 'pos' is not changed in between).


You're right.

> > And, IMHO something should clear always and close (btw, might be better

> > to use close_modifier as variable name and for consistency always_modifier)

> > unless we reach the CPP_COLON case.

> > 

> 

> Good point, I agree with both. Cleared and renamed :)


I think the clearing is still insufficient.
It will clear on e.g. map (always, close, foobar)
but not on map (always) or map (always, close)
because in that case the loop terminates by the while loop condition
no longer being true.

And there is another thing I have missed (and again should be in the
testsuite):
map (always, always)
or
map (always, close, close)
etc. will with the patch diagnose that too many 'always' modifiers
(or 'close'), but that isn't correct diagnostic, there aren't any
modifiers, but the same variable is mapped multiple times.

So, one possibility is to remember details like:
potential always modifier has been seen
potential always modifier has been seen more than once
potential close modifier has been seen
potential close modifier has been seen more than once
and only when seeing the colon enact them and diagnose too many modifiers
(but then not with cp_parser_error but error with a location_t of one of the
modifiers), e.g. always_modifier == -1 would mean 1 potential has been seen,
== -2 more than one potential and == 1 it was modifier.

Or another one is not to do much in the first raw token lookup loop,
just check if it is a sequence of
always
close
,
tokens followed by
CPP_NAME (other than always, close) + CPP_CLONE combo
and in that case just set a bool flag that map-kind is present,
but don't consume any tokens.
And then in another loop if that bool flag is set, lookup non-raw
tokens and parse them, setting flags, doing diagnostics etc.
Basically do the look-ahead only to check if it is
map (var1, var2, ...)
case
or
map (modifiers map-kind: var1, var2, ...)
case.

	Jakub
Tobias Burnus May 11, 2021, 3:30 p.m. | #4
On 11.05.21 17:20, Jakub Jelinek via Gcc-patches wrote:
> One extra thing, sorry, forgot to mention, for the translators it might be

> better to use "too many %qs modifiers", "always" (or, "close").

> That way they can translate it just once instead of twice.


That won't work for

c_parser_error (parser, "expected modifier %<always%> only once");

as the error function is not variadic, contrary to warning_at/error_at; namely

c/c-parser.h:extern bool c_parser_error (c_parser *parser, const char *gmsgid);

Tobias

-----------------
Mentor Graphics (Deutschland) GmbH, Arnulfstrasse 201, 80634 München Registergericht München HRB 106955, Geschäftsführer: Thomas Heurung, Frank Thürauf
Andre Vehreschild via Gcc-patches May 11, 2021, 3:33 p.m. | #5
On Tue, May 11, 2021 at 05:30:19PM +0200, Tobias Burnus wrote:
> On 11.05.21 17:20, Jakub Jelinek via Gcc-patches wrote:

> > One extra thing, sorry, forgot to mention, for the translators it might be

> > better to use "too many %qs modifiers", "always" (or, "close").

> > That way they can translate it just once instead of twice.

> 

> That won't work for

> 

> c_parser_error (parser, "expected modifier %<always%> only once");

> 

> as the error function is not variadic, contrary to warning_at/error_at; namely

> 

> c/c-parser.h:extern bool c_parser_error (c_parser *parser, const char *gmsgid);


Ok, ignore that part then.  Sorry.

	Jakub
Marcel Vollweiler May 12, 2021, 2:36 p.m. | #6
Am 11.05.2021 um 17:20 schrieb Jakub Jelinek:
> On Tue, May 11, 2021 at 04:27:55PM +0200, Marcel Vollweiler wrote:

>>> The usual wording would be

>>> "too many %<always%> modifiers"

>>>

>>

>> Changed for 'always' and 'close' for C and C++.

>

> One extra thing, sorry, forgot to mention, for the translators it might be

> better to use "too many %qs modifiers", "always" (or, "close").

> That way they can translate it just once instead of twice.

>

>>> IMHO you should at least check that tok->type == CPP_NAME before

>>> checking pos + 1 token's type, you don't want to skip over CPP_EOF,

>>> CPP_PRAGMA_EOF, or even CPP_CLOSE_PAREN etc.

>>> Perhaps by adding

>>>         if (tok->type != CPP_NAME)

>>>        break;

>>> right after c_token *tok = c_parser_peek_nth_token_raw (parser, pos); ?

>>

>> The check of the token's type at position 'pos' is done in the condition

>> of the while loop, that means

>>     'c_parser_peek_nth_token_raw (parser, pos + 1)->type == CPP_COLON'

>> is only reached when

>>     'c_parser_peek_nth_token_raw (parser, pos)->type == CPP_NAME'

>> holds (since 'pos' is not changed in between).

>

> You're right.

>

>>> And, IMHO something should clear always and close (btw, might be better

>>> to use close_modifier as variable name and for consistency always_modifier)

>>> unless we reach the CPP_COLON case.

>>>

>>

>> Good point, I agree with both. Cleared and renamed :)

>

> I think the clearing is still insufficient.

> It will clear on e.g. map (always, close, foobar)

> but not on map (always) or map (always, close)

> because in that case the loop terminates by the while loop condition

> no longer being true.


That's true. This is modified together with the diagnostics (see below).

>

> And there is another thing I have missed (and again should be in the

> testsuite):

> map (always, always)

> or

> map (always, close, close)

> etc. will with the patch diagnose that too many 'always' modifiers

> (or 'close'), but that isn't correct diagnostic, there aren't any

> modifiers, but the same variable is mapped multiple times.

>

> So, one possibility is to remember details like:

> potential always modifier has been seen

> potential always modifier has been seen more than once

> potential close modifier has been seen

> potential close modifier has been seen more than once

> and only when seeing the colon enact them and diagnose too many modifiers

> (but then not with cp_parser_error but error with a location_t of one of the

> modifiers), e.g. always_modifier == -1 would mean 1 potential has been seen,

> == -2 more than one potential and == 1 it was modifier.

>

> Or another one is not to do much in the first raw token lookup loop,

> just check if it is a sequence of

> always

> close

> ,

> tokens followed by

> CPP_NAME (other than always, close) + CPP_CLONE combo

> and in that case just set a bool flag that map-kind is present,

> but don't consume any tokens.

> And then in another loop if that bool flag is set, lookup non-raw

> tokens and parse them, setting flags, doing diagnostics etc.

> Basically do the look-ahead only to check if it is

> map (var1, var2, ...)

> case

> or

> map (modifiers map-kind: var1, var2, ...)

> case.


I changed the patch similar to your second approach. I.e., use a first
loop to check if and where a potential map-type is given. In the second
loop (which is only entered if a potential map-type exists) the tokens
are consumed and diagnostic is applied.

I avoided the diagnostic for the variables (besides the modifier) at
this place, since this should continue to be done in
'c_parser_omp_variable_list' / 'cp_parser_omp_var_list_no_open' from my
point of view.

>

>       Jakub

>


The new version of the patch is attached.

Thanks,

Marcel
-----------------
Mentor Graphics (Deutschland) GmbH, Arnulfstrasse 201, 80634 München Registergericht München HRB 106955, Geschäftsführer: Thomas Heurung, Frank Thürauf
OpenMP: Add support for 'close' in map clause

gcc/c/ChangeLog:

	* c-parser.c (c_parser_omp_clause_map): Support map-type-modifier 
	'close'.

gcc/cp/ChangeLog:

	* parser.c (cp_parser_omp_clause_map): Support map-type-modifier 
	'close'.

gcc/testsuite/ChangeLog:

	* c-c++-common/gomp/map-6.c: New test.
	* c-c++-common/gomp/map-7.c: New test.

diff --git a/gcc/c/c-parser.c b/gcc/c/c-parser.c
index 5cdeb21..c7f3f18 100644
--- a/gcc/c/c-parser.c
+++ b/gcc/c/c-parser.c
@@ -15643,54 +15643,83 @@ c_parser_omp_clause_depend (c_parser *parser, tree list)
    map-kind:
      alloc | to | from | tofrom | release | delete
 
-   map ( always [,] map-kind: variable-list ) */
+   map ( always [,] map-kind: variable-list )
+
+   OpenMP 5.0:
+   map ( [map-type-modifier[,] ...] map-kind: variable-list )
+
+   map-type-modifier:
+     always | close */
 
 static tree
 c_parser_omp_clause_map (c_parser *parser, tree list)
 {
   location_t clause_loc = c_parser_peek_token (parser)->location;
   enum gomp_map_kind kind = GOMP_MAP_TOFROM;
-  int always = 0;
-  enum c_id_kind always_id_kind = C_ID_NONE;
-  location_t always_loc = UNKNOWN_LOCATION;
-  tree always_id = NULL_TREE;
   tree nl, c;
 
   matching_parens parens;
   if (!parens.require_open (parser))
     return list;
 
-  if (c_parser_next_token_is (parser, CPP_NAME))
+  int pos = 1;
+  int map_kind_pos = 0;
+  while (c_parser_peek_nth_token_raw (parser, pos)->type == CPP_NAME)
+    {
+      if (c_parser_peek_nth_token_raw (parser, pos + 1)->type == CPP_COLON)
+	{
+	  map_kind_pos = pos;
+	  break;
+	}
+
+      if (c_parser_peek_nth_token_raw (parser, pos + 1)->type == CPP_COMMA)
+	pos++;
+      pos++;
+    }
+
+  int always_modifier = 0;
+  int close_modifier = 0;
+  for (int pos = 1; pos < map_kind_pos; ++pos)
     {
       c_token *tok = c_parser_peek_token (parser);
+
+      if (tok->type == CPP_COMMA)
+	{
+	  c_parser_consume_token (parser);
+	  continue;
+	}
+
       const char *p = IDENTIFIER_POINTER (tok->value);
-      always_id_kind = tok->id_kind;
-      always_loc = tok->location;
-      always_id = tok->value;
       if (strcmp ("always", p) == 0)
 	{
-	  c_token *sectok = c_parser_peek_2nd_token (parser);
-	  if (sectok->type == CPP_COMMA)
+	  if (always_modifier)
 	    {
-	      c_parser_consume_token (parser);
-	      c_parser_consume_token (parser);
-	      always = 2;
+	      c_parser_error (parser, "too many %<always%> modifiers");
+	      parens.skip_until_found_close (parser);
+	      return list;
 	    }
-	  else if (sectok->type == CPP_NAME)
+	  always_modifier++;
+	}
+      else if (strcmp ("close", p) == 0)
+	{
+	  if (close_modifier)
 	    {
-	      p = IDENTIFIER_POINTER (sectok->value);
-	      if (strcmp ("alloc", p) == 0
-		  || strcmp ("to", p) == 0
-		  || strcmp ("from", p) == 0
-		  || strcmp ("tofrom", p) == 0
-		  || strcmp ("release", p) == 0
-		  || strcmp ("delete", p) == 0)
-		{
-		  c_parser_consume_token (parser);
-		  always = 1;
-		}
+	      c_parser_error (parser, "too many %<close%> modifiers");
+	      parens.skip_until_found_close (parser);
+	      return list;
 	    }
+	  close_modifier++;
 	}
+      else
+	{
+	  c_parser_error (parser, "%<#pragma omp target%> with map-type-"
+				  "modifier other than %<always%> or %<close%>"
+				  "on %<map%> clause");
+	  parens.skip_until_found_close (parser);
+	  return list;
+	}
+
+	c_parser_consume_token (parser);
     }
 
   if (c_parser_next_token_is (parser, CPP_NAME)
@@ -15700,11 +15729,11 @@ c_parser_omp_clause_map (c_parser *parser, tree list)
       if (strcmp ("alloc", p) == 0)
 	kind = GOMP_MAP_ALLOC;
       else if (strcmp ("to", p) == 0)
-	kind = always ? GOMP_MAP_ALWAYS_TO : GOMP_MAP_TO;
+	kind = always_modifier ? GOMP_MAP_ALWAYS_TO : GOMP_MAP_TO;
       else if (strcmp ("from", p) == 0)
-	kind = always ? GOMP_MAP_ALWAYS_FROM : GOMP_MAP_FROM;
+	kind = always_modifier ? GOMP_MAP_ALWAYS_FROM : GOMP_MAP_FROM;
       else if (strcmp ("tofrom", p) == 0)
-	kind = always ? GOMP_MAP_ALWAYS_TOFROM : GOMP_MAP_TOFROM;
+	kind = always_modifier ? GOMP_MAP_ALWAYS_TOFROM : GOMP_MAP_TOFROM;
       else if (strcmp ("release", p) == 0)
 	kind = GOMP_MAP_RELEASE;
       else if (strcmp ("delete", p) == 0)
@@ -15719,35 +15748,6 @@ c_parser_omp_clause_map (c_parser *parser, tree list)
       c_parser_consume_token (parser);
       c_parser_consume_token (parser);
     }
-  else if (always)
-    {
-      if (always_id_kind != C_ID_ID)
-	{
-	  c_parser_error (parser, "expected identifier");
-	  parens.skip_until_found_close (parser);
-	  return list;
-	}
-
-      tree t = lookup_name (always_id);
-      if (t == NULL_TREE)
-	{
-	  undeclared_variable (always_loc, always_id);
-	  t = error_mark_node;
-	}
-      if (t != error_mark_node)
-	{
-	  tree u = build_omp_clause (clause_loc, OMP_CLAUSE_MAP);
-	  OMP_CLAUSE_DECL (u) = t;
-	  OMP_CLAUSE_CHAIN (u) = list;
-	  OMP_CLAUSE_SET_MAP_KIND (u, kind);
-	  list = u;
-	}
-      if (always == 1)
-	{
-	  parens.skip_until_found_close (parser);
-	  return list;
-	}
-    }
 
   nl = c_parser_omp_variable_list (parser, clause_loc, OMP_CLAUSE_MAP, list);
 
diff --git a/gcc/cp/parser.c b/gcc/cp/parser.c
index 99eccf0..7f7342f 100644
--- a/gcc/cp/parser.c
+++ b/gcc/cp/parser.c
@@ -37840,40 +37840,90 @@ cp_parser_omp_clause_depend (cp_parser *parser, tree list, location_t loc)
    map-kind:
      alloc | to | from | tofrom | release | delete
 
-   map ( always [,] map-kind: variable-list ) */
+   map ( always [,] map-kind: variable-list )
+
+   OpenMP 5.0:
+   map ( [map-type-modifier[,] ...] map-kind: variable-list )
+
+   map-type-modifier:
+     always | close */
 
 static tree
 cp_parser_omp_clause_map (cp_parser *parser, tree list)
 {
   tree nlist, c;
   enum gomp_map_kind kind = GOMP_MAP_TOFROM;
-  bool always = false;
 
   if (!cp_parser_require (parser, CPP_OPEN_PAREN, RT_OPEN_PAREN))
     return list;
 
-  if (cp_lexer_next_token_is (parser->lexer, CPP_NAME))
+  int pos = 1;
+  int map_kind_pos = 0;
+  while (cp_lexer_peek_nth_token (parser->lexer, pos)->type == CPP_NAME
+	 || cp_lexer_peek_nth_token (parser->lexer, pos)->keyword == RID_DELETE)
     {
-      tree id = cp_lexer_peek_token (parser->lexer)->u.value;
-      const char *p = IDENTIFIER_POINTER (id);
+      if (cp_lexer_peek_nth_token (parser->lexer, pos + 1)->type == CPP_COLON)
+	{
+	  map_kind_pos = pos;
+	  break;
+	}
+
+      if (cp_lexer_peek_nth_token (parser->lexer, pos + 1)->type == CPP_COMMA)
+	pos++;
+      pos++;
+    }
 
+  bool always_modifier = false;
+  bool close_modifier = false;
+  for (int pos = 1; pos < map_kind_pos; ++pos)
+    {
+      cp_token *tok = cp_lexer_peek_token (parser->lexer);
+      if (tok->type == CPP_COMMA)
+	{
+	  cp_lexer_consume_token (parser->lexer);
+	  continue;
+	}
+
+      const char *p = IDENTIFIER_POINTER (tok->u.value);
       if (strcmp ("always", p) == 0)
 	{
-	  int nth = 2;
-	  if (cp_lexer_peek_nth_token (parser->lexer, 2)->type == CPP_COMMA)
-	    nth++;
-	  if ((cp_lexer_peek_nth_token (parser->lexer, nth)->type == CPP_NAME
-	       || (cp_lexer_peek_nth_token (parser->lexer, nth)->keyword
-		   == RID_DELETE))
-	      && (cp_lexer_peek_nth_token (parser->lexer, nth + 1)->type
-		  == CPP_COLON))
+	  if (always_modifier)
 	    {
-	      always = true;
-	      cp_lexer_consume_token (parser->lexer);
-	      if (nth == 3)
-		cp_lexer_consume_token (parser->lexer);
+	      cp_parser_error (parser, "too many %<always%> modifiers");
+	      cp_parser_skip_to_closing_parenthesis (parser,
+						     /*recovering=*/true,
+						     /*or_comma=*/false,
+						     /*consume_paren=*/true);
+	      return list;
+	    }
+	  always_modifier = true;
+	}
+      else if (strcmp ("close", p) == 0)
+	{
+	  if (close_modifier)
+	    {
+	      cp_parser_error (parser, "too many %<close%> modifiers");
+	      cp_parser_skip_to_closing_parenthesis (parser,
+						     /*recovering=*/true,
+						     /*or_comma=*/false,
+						     /*consume_paren=*/true);
+	      return list;
 	    }
+	  close_modifier = true;
+	}
+      else
+	{
+	  cp_parser_error (parser, "%<#pragma omp target%> with map-type-"
+				   "modifier other than %<always%> or %<close%>"
+				   "on %<map%> clause");
+	  cp_parser_skip_to_closing_parenthesis (parser,
+						 /*recovering=*/true,
+						 /*or_comma=*/false,
+						 /*consume_paren=*/true);
+	      return list;
 	}
+
+	cp_lexer_consume_token (parser->lexer);
     }
 
   if (cp_lexer_next_token_is (parser->lexer, CPP_NAME)
@@ -37885,11 +37935,11 @@ cp_parser_omp_clause_map (cp_parser *parser, tree list)
       if (strcmp ("alloc", p) == 0)
 	kind = GOMP_MAP_ALLOC;
       else if (strcmp ("to", p) == 0)
-	kind = always ? GOMP_MAP_ALWAYS_TO : GOMP_MAP_TO;
+	kind = always_modifier ? GOMP_MAP_ALWAYS_TO : GOMP_MAP_TO;
       else if (strcmp ("from", p) == 0)
-	kind = always ? GOMP_MAP_ALWAYS_FROM : GOMP_MAP_FROM;
+	kind = always_modifier ? GOMP_MAP_ALWAYS_FROM : GOMP_MAP_FROM;
       else if (strcmp ("tofrom", p) == 0)
-	kind = always ? GOMP_MAP_ALWAYS_TOFROM : GOMP_MAP_TOFROM;
+	kind = always_modifier ? GOMP_MAP_ALWAYS_TOFROM : GOMP_MAP_TOFROM;
       else if (strcmp ("release", p) == 0)
 	kind = GOMP_MAP_RELEASE;
       else
diff --git a/gcc/testsuite/c-c++-common/gomp/map-6.c b/gcc/testsuite/c-c++-common/gomp/map-6.c
new file mode 100644
index 0000000..e30bd14
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/gomp/map-6.c
@@ -0,0 +1,135 @@
+/* { dg-do compile } */
+/* { dg-additional-options "-fdump-tree-original" } */
+
+void
+foo (void)
+{
+  /* Test to ensure that the close modifier is parsed and ignored in map clauses. */
+  int a, b, b1, b2, b3, b4, b5, b6, b7;
+
+  #pragma omp target map (a)
+  ;
+
+  #pragma omp target map (to:a)
+  ;
+
+  #pragma omp target map (a to: b) /* { dg-error "'#pragma omp target' with map-type-modifier other than 'always' or 'close'" } */
+  ;
+
+  #pragma omp target map (close, a to: b) /* { dg-error "'#pragma omp target' with map-type-modifier other than 'always' or 'close'" } */
+  ;
+
+  #pragma omp target map (close a) /* { dg-error "'close' undeclared" "" { target c } } */ 
+  /* { dg-error "'close' has not been declared" "" { target c++ } .-1 } */ 
+  /* { dg-error "expected '\\)' before 'a'" "" { target *-*-* } .-2 } */
+  ;
+
+  #pragma omp target map (always a) /* { dg-error "'always' undeclared" "" { target c } } */
+  /* { dg-error "'always' has not been declared" "" { target c++ } .-1 } */ 
+  /* { dg-error "expected '\\)' before 'a'" "" { target *-*-* } .-2 } */
+  ;
+
+  #pragma omp target map (close to:a)
+  ;
+
+  #pragma omp target map (close, to:a)
+  ;
+
+  #pragma omp target map (close delete:a) /* { dg-error "'#pragma omp target' with map-type other than 'to', 'from', 'tofrom' or 'alloc' on 'map' clause" } */
+  ;
+
+  #pragma omp target map (close always to:b1)
+  ;
+
+  #pragma omp target map (close, always to:b2)
+  ;
+
+  #pragma omp target map (close, always, to:b3)
+  ;
+
+  #pragma omp target map (always close to:b4)
+  ;
+
+  #pragma omp target map (always, close to:b5)
+  ;
+
+  #pragma omp target map (always, close, to:b6)
+  ;
+
+  #pragma omp target map (always, always, to:a) /* { dg-error "too many 'always' modifiers" } */
+  ;
+
+  #pragma omp target map (always always, to:a) /* { dg-error "too many 'always' modifiers" } */
+  ;
+
+  #pragma omp target map (always, always to:a) /* { dg-error "too many 'always' modifiers" } */
+  ;
+
+  #pragma omp target map (always always to:a) /* { dg-error "too many 'always' modifiers" } */
+  ;
+
+  #pragma omp target map (close, close, to:a) /* { dg-error "too many 'close' modifiers" } */
+  ;
+
+  #pragma omp target map (close close, to:a) /* { dg-error "too many 'close' modifiers" } */
+  ;
+
+  #pragma omp target map (close, close to:a) /* { dg-error "too many 'close' modifiers" } */
+  ;
+
+  #pragma omp target map (close close to:a) /* { dg-error "too many 'close' modifiers" } */
+  ;
+
+  #pragma omp target map (always to : a) map (close to : b)
+  ;
+
+  int close = 0;
+  #pragma omp target map (close) 
+  ;
+
+  #pragma omp target map (close a) /* { dg-error "expected '\\)' before 'a'" } */ 
+  ;
+
+  int always = 0;
+  #pragma omp target map (always)
+  ;
+
+  #pragma omp target map (always a) /* { dg-error "expected '\\)' before 'a'" } */
+  ;
+
+  #pragma omp target map (always, close)
+  ;
+
+  #pragma omp target map (always, always)  /* { dg-error "'always' appears more than once in map clauses" } */
+  ;
+
+  #pragma omp target map (always, always, close)  /* { dg-error "'always' appears more than once in map clauses" } */
+  ;
+
+  #pragma omp target map (always, close, to: always, close, b7)
+  ;
+
+  int to = 0;
+  #pragma omp target map (always, close, to)
+  ;
+
+  #pragma omp target map (to, always, close)
+    {
+      to = always = close = 1;
+    }
+  if (to != 1 || always != 1 || close != 1)
+    __builtin_abort ();
+  ;
+}
+
+/* { dg-final { scan-tree-dump-not "map\\(\[^\n\r)]*close\[^\n\r)]*to:" "original" } } */
+
+/* { dg-final { scan-tree-dump-times "pragma omp target map\\(always,to:" 7 "original" } } */
+
+/* { dg-final { scan-tree-dump "pragma omp target map\\(always,to:b1" "original" } } */
+/* { dg-final { scan-tree-dump "pragma omp target map\\(always,to:b2" "original" } } */
+/* { dg-final { scan-tree-dump "pragma omp target map\\(always,to:b3" "original" } } */
+/* { dg-final { scan-tree-dump "pragma omp target map\\(always,to:b4" "original" } } */
+/* { dg-final { scan-tree-dump "pragma omp target map\\(always,to:b5" "original" } } */
+/* { dg-final { scan-tree-dump "pragma omp target map\\(always,to:b6" "original" } } */
+/* { dg-final { scan-tree-dump "pragma omp target map\\(always,to:b7\\) map\\(always,to:close\\) map\\(always,to:always\\)" "original" } } */
diff --git a/gcc/testsuite/c-c++-common/gomp/map-7.c b/gcc/testsuite/c-c++-common/gomp/map-7.c
new file mode 100644
index 0000000..3f1e972
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/gomp/map-7.c
@@ -0,0 +1,20 @@
+/* { dg-do compile } */
+
+void
+foo (void)
+{
+  /* Test to ensure that the close modifier is parsed and ignored in map clauses. */
+
+  #define N 1024
+  int always[N];
+  int close;
+
+  #pragma omp target map(always[:N]) 
+  ;
+
+  #pragma omp target map(close, always[:N]) 
+  ;
+
+  #pragma omp target map(always[:N], close) 
+  ;
+}
Andre Vehreschild via Gcc-patches May 12, 2021, 2:46 p.m. | #7
On Wed, May 12, 2021 at 04:36:00PM +0200, Marcel Vollweiler wrote:
> +	  c_parser_error (parser, "%<#pragma omp target%> with map-type-"

> +				  "modifier other than %<always%> or %<close%>"


Please just write modifier rather than map-type-modifier

> +	  cp_parser_error (parser, "%<#pragma omp target%> with map-type-"


Likewise.

> +				   "modifier other than %<always%> or %<close%>"

> +				   "on %<map%> clause");

> +	  cp_parser_skip_to_closing_parenthesis (parser,

> +						 /*recovering=*/true,

> +						 /*or_comma=*/false,

> +						 /*consume_paren=*/true);

> +	      return list;


And this line looks misindented, previous is tab + 2 spaces, this one
tab + 6 spaces, should be tab + 4 spaces.

Ok for trunk with those changes (and adjusting the testcases accordingly).

	Jakub

Patch

diff --git a/gcc/c/c-parser.c b/gcc/c/c-parser.c
index 5cdeb21..78cba7f 100644
--- a/gcc/c/c-parser.c
+++ b/gcc/c/c-parser.c
@@ -15643,14 +15643,19 @@  c_parser_omp_clause_depend (c_parser *parser, tree list)
    map-kind:
      alloc | to | from | tofrom | release | delete
 
-   map ( always [,] map-kind: variable-list ) */
+   map ( always [,] map-kind: variable-list )
+
+   OpenMP 5.0:
+   map ( [map-type-modifier[,] ...] map-kind: variable-list )
+
+   map-type-modifier:
+     always | close */
 
 static tree
 c_parser_omp_clause_map (c_parser *parser, tree list)
 {
   location_t clause_loc = c_parser_peek_token (parser)->location;
   enum gomp_map_kind kind = GOMP_MAP_TOFROM;
-  int always = 0;
   enum c_id_kind always_id_kind = C_ID_NONE;
   location_t always_loc = UNKNOWN_LOCATION;
   tree always_id = NULL_TREE;
@@ -15660,37 +15665,54 @@  c_parser_omp_clause_map (c_parser *parser, tree list)
   if (!parens.require_open (parser))
     return list;
 
-  if (c_parser_next_token_is (parser, CPP_NAME))
+  int always = 0;
+  int close = 0;
+  int pos = 1;
+  while (c_parser_peek_nth_token_raw (parser, pos)->type == CPP_NAME)
     {
-      c_token *tok = c_parser_peek_token (parser);
+      c_token *tok = c_parser_peek_nth_token_raw (parser, pos);
       const char *p = IDENTIFIER_POINTER (tok->value);
-      always_id_kind = tok->id_kind;
-      always_loc = tok->location;
-      always_id = tok->value;
       if (strcmp ("always", p) == 0)
 	{
-	  c_token *sectok = c_parser_peek_2nd_token (parser);
-	  if (sectok->type == CPP_COMMA)
+	  if (always)
 	    {
-	      c_parser_consume_token (parser);
-	      c_parser_consume_token (parser);
-	      always = 2;
+	      c_parser_error (parser, "expected modifier %<always%> only once");
+	      parens.skip_until_found_close (parser);
+	      return list;
+	    }
+
+	  always_id_kind = tok->id_kind;
+	  always_loc = tok->location;
+	  always_id = tok->value;
+
+	  always++;
+	}
+      else if (strcmp ("close", p) == 0)
+	{
+	  if (close)
+	    {
+	      c_parser_error (parser, "expected modifier %<close%> only once");
+	      parens.skip_until_found_close (parser);
+	      return list;
 	    }
-	  else if (sectok->type == CPP_NAME)
+
+	  close++;
+	}
+      else if (c_parser_peek_nth_token_raw (parser, pos + 1)->type == CPP_COLON)
+	{
+	  for (int i = 1; i < pos; ++i)
 	    {
-	      p = IDENTIFIER_POINTER (sectok->value);
-	      if (strcmp ("alloc", p) == 0
-		  || strcmp ("to", p) == 0
-		  || strcmp ("from", p) == 0
-		  || strcmp ("tofrom", p) == 0
-		  || strcmp ("release", p) == 0
-		  || strcmp ("delete", p) == 0)
-		{
-		  c_parser_consume_token (parser);
-		  always = 1;
-		}
+	      c_parser_peek_token(parser);
+	      c_parser_consume_token (parser);
 	    }
+	  break;
 	}
+      else
+	break;
+
+      if (c_parser_peek_nth_token_raw (parser, pos + 1)->type == CPP_COMMA)
+	pos++;
+      pos++;
     }
 
   if (c_parser_next_token_is (parser, CPP_NAME)
@@ -15719,35 +15741,6 @@  c_parser_omp_clause_map (c_parser *parser, tree list)
       c_parser_consume_token (parser);
       c_parser_consume_token (parser);
     }
-  else if (always)
-    {
-      if (always_id_kind != C_ID_ID)
-	{
-	  c_parser_error (parser, "expected identifier");
-	  parens.skip_until_found_close (parser);
-	  return list;
-	}
-
-      tree t = lookup_name (always_id);
-      if (t == NULL_TREE)
-	{
-	  undeclared_variable (always_loc, always_id);
-	  t = error_mark_node;
-	}
-      if (t != error_mark_node)
-	{
-	  tree u = build_omp_clause (clause_loc, OMP_CLAUSE_MAP);
-	  OMP_CLAUSE_DECL (u) = t;
-	  OMP_CLAUSE_CHAIN (u) = list;
-	  OMP_CLAUSE_SET_MAP_KIND (u, kind);
-	  list = u;
-	}
-      if (always == 1)
-	{
-	  parens.skip_until_found_close (parser);
-	  return list;
-	}
-    }
 
   nl = c_parser_omp_variable_list (parser, clause_loc, OMP_CLAUSE_MAP, list);
 
diff --git a/gcc/cp/parser.c b/gcc/cp/parser.c
index 99eccf0..f7fecf7 100644
--- a/gcc/cp/parser.c
+++ b/gcc/cp/parser.c
@@ -37840,7 +37840,13 @@  cp_parser_omp_clause_depend (cp_parser *parser, tree list, location_t loc)
    map-kind:
      alloc | to | from | tofrom | release | delete
 
-   map ( always [,] map-kind: variable-list ) */
+   map ( always [,] map-kind: variable-list )
+
+   OpenMP 5.0:
+   map ( [map-type-modifier[,] ...] map-kind: variable-list )
+
+   map-type-modifier:
+     always | close */
 
 static tree
 cp_parser_omp_clause_map (cp_parser *parser, tree list)
@@ -37848,32 +37854,61 @@  cp_parser_omp_clause_map (cp_parser *parser, tree list)
   tree nlist, c;
   enum gomp_map_kind kind = GOMP_MAP_TOFROM;
   bool always = false;
+  bool close = false;
+  int pos = 1;
 
   if (!cp_parser_require (parser, CPP_OPEN_PAREN, RT_OPEN_PAREN))
     return list;
 
-  if (cp_lexer_next_token_is (parser->lexer, CPP_NAME))
+  while (cp_lexer_peek_nth_token (parser->lexer, pos)->type == CPP_NAME
+	 || cp_lexer_peek_nth_token (parser->lexer, pos)->keyword == RID_DELETE)
     {
-      tree id = cp_lexer_peek_token (parser->lexer)->u.value;
+      tree id = cp_lexer_peek_nth_token (parser->lexer, pos)->u.value;
       const char *p = IDENTIFIER_POINTER (id);
 
       if (strcmp ("always", p) == 0)
 	{
-	  int nth = 2;
-	  if (cp_lexer_peek_nth_token (parser->lexer, 2)->type == CPP_COMMA)
-	    nth++;
-	  if ((cp_lexer_peek_nth_token (parser->lexer, nth)->type == CPP_NAME
-	       || (cp_lexer_peek_nth_token (parser->lexer, nth)->keyword
-		   == RID_DELETE))
-	      && (cp_lexer_peek_nth_token (parser->lexer, nth + 1)->type
-		  == CPP_COLON))
+	  if (always)
 	    {
-	      always = true;
-	      cp_lexer_consume_token (parser->lexer);
-	      if (nth == 3)
-		cp_lexer_consume_token (parser->lexer);
+	      cp_parser_error (parser,
+			       "expected modifier %<always%> only once");
+	      cp_parser_skip_to_closing_parenthesis (parser,
+						     /*recovering=*/true,
+						     /*or_comma=*/false,
+						     /*consume_paren=*/true);
+	      return list;
+	    }
+
+	  always = true;
+	}
+      else if (strcmp ("close", p) == 0)
+	{
+	  if (close)
+	    {
+	      cp_parser_error (parser,
+			       "expected modifier %<close%> only once");
+	      cp_parser_skip_to_closing_parenthesis (parser,
+						     /*recovering=*/true,
+						     /*or_comma=*/false,
+						     /*consume_paren=*/true);
+	      return list;
 	    }
+
+	  close = true;
 	}
+      else if (cp_lexer_peek_nth_token (parser->lexer, pos + 1)->type
+	       == CPP_COLON)
+	{
+	  for (int i = 1; i < pos; ++i)
+	    cp_lexer_consume_token (parser->lexer);
+	  break;
+	}
+      else
+	break;
+
+      if (cp_lexer_peek_nth_token (parser->lexer, pos + 1)->type == CPP_COMMA)
+	pos++;
+      pos++;
     }
 
   if (cp_lexer_next_token_is (parser->lexer, CPP_NAME)
diff --git a/gcc/testsuite/c-c++-common/gomp/map-6.c b/gcc/testsuite/c-c++-common/gomp/map-6.c
new file mode 100644
index 0000000..f7e22ed
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/gomp/map-6.c
@@ -0,0 +1,122 @@ 
+/* { dg-do compile } */
+/* { dg-additional-options "-fdump-tree-original" } */
+
+void
+foo (void)
+{
+  /* Test to ensure that the close modifier is parsed and ignored in map clauses. */
+  int a, b, b1, b2, b3, b4, b5, b6;
+
+  #pragma omp target map (a)
+  ;
+
+  #pragma omp target map (to:a)
+  ;
+
+  #pragma omp target map (a to: b) /* { dg-error "expected '\\)' before 'to'" } */
+  ;
+
+  #pragma omp target map (close a) /* { dg-error "'close' undeclared" "" { target c } } */ 
+  /* { dg-error "'close' has not been declared" "" { target c++ } .-1 } */ 
+  /* { dg-error "expected '\\)' before 'a'" "" { target *-*-* } .-2 } */
+  ;
+
+  #pragma omp target map (always a) /* { dg-error "'always' undeclared" "" { target c } } */
+  /* { dg-error "'always' has not been declared" "" { target c++ } .-1 } */ 
+  /* { dg-error "expected '\\)' before 'a'" "" { target *-*-* } .-2 } */
+  ;
+
+  #pragma omp target map (close to:a)
+  ;
+
+  #pragma omp target map (close, to:a)
+  ;
+
+  #pragma omp target map (close delete:a) /* { dg-error "'#pragma omp target' with map-type other than 'to', 'from', 'tofrom' or 'alloc' on 'map' clause" } */
+  ;
+
+  #pragma omp target map (close always to:b1)
+  ;
+
+  #pragma omp target map (close, always to:b2)
+  ;
+
+  #pragma omp target map (close, always, to:b3)
+  ;
+
+  #pragma omp target map (always close to:b4)
+  ;
+
+  #pragma omp target map (always, close to:b5)
+  ;
+
+  #pragma omp target map (always, close, to:b6)
+  ;
+
+  #pragma omp target map (always, always, to:a) /* { dg-error "expected modifier 'always' only once" } */
+  ;
+
+  #pragma omp target map (always always, to:a) /* { dg-error "expected modifier 'always' only once" } */
+  ;
+
+  #pragma omp target map (always, always to:a) /* { dg-error "expected modifier 'always' only once" } */
+  ;
+
+  #pragma omp target map (always always to:a) /* { dg-error "expected modifier 'always' only once" } */
+  ;
+
+  #pragma omp target map (close, close, to:a) /* { dg-error "expected modifier 'close' only once" } */
+  ;
+
+  #pragma omp target map (close close, to:a) /* { dg-error "expected modifier 'close' only once" } */
+  ;
+
+  #pragma omp target map (close, close to:a) /* { dg-error "expected modifier 'close' only once" } */
+  ;
+
+  #pragma omp target map (close close to:a) /* { dg-error "expected modifier 'close' only once" } */
+  ;
+
+  #pragma omp target map (always to : a) map (close to : b)
+  ;
+
+  int close = 0;
+  #pragma omp target map (close) 
+  ;
+
+  #pragma omp target map (close a) /* { dg-error "expected '\\)' before 'a'" } */ 
+  ;
+
+  int always = 0;
+  #pragma omp target map (always)
+  ;
+
+  #pragma omp target map (always a) /* { dg-error "expected '\\)' before 'a'" } */
+  ;
+
+  #pragma omp target map (always, close)
+  ;
+
+  int to = 0;
+  #pragma omp target map (always, close, to)
+  ;
+
+  #pragma omp target map (to, always, close)
+    {
+      to = always = close = 1;
+    }
+  if (to != 1 || always != 1 || close != 1)
+    __builtin_abort ();
+  ;
+}
+
+/* { dg-final { scan-tree-dump-not "map\\(\[^\n\r)]*close\[^\n\r)]*to:" "original" } } */
+
+/* { dg-final { scan-tree-dump-times "pragma omp target map\\(always,to:" 6 "original" } } */
+
+/* { dg-final { scan-tree-dump "pragma omp target map\\(always,to:b1" "original" } } */
+/* { dg-final { scan-tree-dump "pragma omp target map\\(always,to:b2" "original" } } */
+/* { dg-final { scan-tree-dump "pragma omp target map\\(always,to:b3" "original" } } */
+/* { dg-final { scan-tree-dump "pragma omp target map\\(always,to:b4" "original" } } */
+/* { dg-final { scan-tree-dump "pragma omp target map\\(always,to:b5" "original" } } */
+/* { dg-final { scan-tree-dump "pragma omp target map\\(always,to:b6" "original" } } */
diff --git a/gcc/testsuite/c-c++-common/gomp/map-7.c b/gcc/testsuite/c-c++-common/gomp/map-7.c
new file mode 100644
index 0000000..3f1e972
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/gomp/map-7.c
@@ -0,0 +1,20 @@ 
+/* { dg-do compile } */
+
+void
+foo (void)
+{
+  /* Test to ensure that the close modifier is parsed and ignored in map clauses. */
+
+  #define N 1024
+  int always[N];
+  int close;
+
+  #pragma omp target map(always[:N]) 
+  ;
+
+  #pragma omp target map(close, always[:N]) 
+  ;
+
+  #pragma omp target map(always[:N], close) 
+  ;
+}