[Patch v2] C, C++, Fortran, OpenMP: Add support for device-modifiers for 'omp target device'

Marcel Vollweiler marcel@codesourcery.com
Fri Aug 20 19:18:32 GMT 2021


Hi Jakub,

this is the second version of the patch for the device-modifiers for
'omp target device'.

Am 20.07.2021 um 15:30 schrieb Jakub Jelinek:
> On Wed, Jul 07, 2021 at 07:59:58PM +0200, Marcel Vollweiler wrote:
>> OpenMP: Add support for device-modifiers for 'omp target device'
>>
>> gcc/c/ChangeLog:
>>
>>      * c-parser.c (c_parser_omp_clause_device): Add support for
>>      device-modifiers for 'omp target device'.
>>
>> gcc/cp/ChangeLog:
>>
>>      * parser.c (cp_parser_omp_clause_device): Add support for
>>      device-modifiers for 'omp target device'.
>>
>> gcc/fortran/ChangeLog:
>>
>>      * openmp.c (gfc_match_omp_clauses): Add support for
>>      device-modifiers for 'omp target device'.
>>
>> gcc/testsuite/ChangeLog:
>>
>>      * c-c++-common/gomp/target-device-1.c: New test.
>>      * c-c++-common/gomp/target-device-2.c: New test.
>>      * gfortran.dg/gomp/target-device-1.f90: New test.
>>      * gfortran.dg/gomp/target-device-2.f90: New test.
>
>>   static tree
>>   c_parser_omp_clause_device (c_parser *parser, tree list)
>>   {
>>     location_t clause_loc = c_parser_peek_token (parser)->location;
>> +  location_t expr_loc;
>> +  c_expr expr;
>> +  tree c, t;
>> +
>>     matching_parens parens;
>> -  if (parens.require_open (parser))
>> +  if (!parens.require_open (parser))
>> +    return list;
>> +
>> +  int pos = 1;
>> +  int pos_colon = 0;
>> +  while (c_parser_peek_nth_token_raw (parser, pos)->type == CPP_NAME
>> +     || c_parser_peek_nth_token_raw (parser, pos)->type == CPP_COLON
>> +     || c_parser_peek_nth_token_raw (parser, pos)->type == CPP_COMMA)
>
> Why CPP_COMMA?  The OpenMP 5.0/5.1/5.2 grammar only supports a single device
> modifier.
> So please simplify it to just an
>    if (c_parser_next_token_is (parser, CPP_NAME)
>        && c_parser_peek_2nd_token (parser, 2)->type == CPP_COLON)
>     {
> and check there just for the two modifiers.
>        const char *p
>       = IDENTIFIER_POINTER (c_parser_peek_token (parser)->value);
>        if (strcmp ("ancestor", p) == 0)
>          ...
>        else if (strcmp ("device-num", p) == 0)
>       ;
>        else
>          error_at (..., "expected %<ancestor%> or %<device-num%>");
>      }
> Similarly for C++.

The parser files for C and C++ are simplyfied accordingly.

>
> Also, even if we sorry on device(ancestor: ...), it would be nice if you
> in tree.h define OMP_CLAUSE_DEVICE_ANCESTOR macro (with
>    (OMP_CLAUSE_SUBCODE_CHECK (NODE, OMP_CLAUSE_DEVICE)->base.public_flag)
> definition), set it, sorry later on it (e.g. omp-expand.c) only if it
> survived till then (wasn't removed because of other errors) and diagnose
> the various restrictions/requirements on device(ancestor:).

I changed it as you proposed. I marked the tests for "sorry,
unimplemented: 'ancestor' not yet supported" with xfail because a
previous sorry for "requires reverse_offload" suppresses the message for
'ancestor'. "reverse_offload" is explicitly needed due to the
specificated ancestor restrictions (OpenMP specification p. 175, l. 1).

> In particular:
> 1) that OMP_CLAUSE_DEVICE clauses with OMP_CLAUSE_DEVICE_ANCESTOR
>     only appear on OMP_TARGET and not on other constructs
>     (this can be easily tested e.g. during gimplification, when
>     gimplify_scan_omp_clauses sees OMP_CLAUSE_DEVICE with
>     OMP_CLAUSE_DEVICE_ANCESTOR and code != OMP_TARGET, diagnose)
> 2) that if after the usual fully folding the argument is INTEGER_CST,
>     it is equal to 1 (the spec says must evaluate to 1, but doesn't say
>     it has to be a constant, so it can evaluate to 1 at runtime but if it is
>     a constant other than 1, we know it will not evaluate to 1); this can be
>     done in *finish_omp_clauses
> 3) that omp_requires_mask has OMP_REQUIRES_REVERSE_OFFLOAD set; this should
>     be checked during the parsing
> 4) only the device, firstprivate, private, defaultmap, and map clauses may
>     appear on the construct; can be also done during gimplification, there is
>     at most one device clause, so walking all clauses when we see
>     OMP_CLAUSE_DEVICE_ANCESTOR is still linear complexity
> 5) no OpenMP constructs or calls to OpenMP API runtime routines are allowed inside
>     the corresponding target region (this is something that should be checked
>     in omp-low.c region nesting code, we already have similar restrictions
>     for e.g. the loop construct)
> Everything should be covered by testcases.

Tests were added for all cases.

>
>       Jakub
>

I tested on x86_64-linux with nvptx offloading with no regressions.

Marcel
-----------------
Siemens Electronic Design Automation GmbH; Anschrift: Arnulfstraße 201, 80634 München; Gesellschaft mit beschränkter Haftung; Geschäftsführer: Thomas Heurung, Frank Thürauf; Sitz der Gesellschaft: München; Registergericht München, HRB 106955
-------------- next part --------------
Add support for device-modifiers for 'omp target device'.

'device_num' and 'ancestor' are now parsed on target device constructs for C,
C++, and Fortran (see OpenMP specification 5.0, p. 170). When 'ancestor' is
 used, then 'sorry, not supported' is output. Moreover, the restrictions for
'ancestor' are implemented (see OpenMP specification 5.0, p. 174f).

gcc/c/ChangeLog:

	* c-parser.c (c_parser_omp_clause_device): Parse device-modifiers 'device_num'
	and 'ancestor' in 'target device' clauses.
	* c-typeck.c (c_finish_omp_clauses): Error handling. Constant device ids must
	evaluate to '1' if 'ancestor' is used.

gcc/cp/ChangeLog:

	* parser.c (cp_parser_omp_clause_device): Parse device-modifiers 'device_num'
	and 'ancestor' in 'target device' clauses.
	* semantics.c (finish_omp_clauses): Error handling. Constant device ids must
	evaluate to '1' if 'ancestor' is used.

gcc/fortran/ChangeLog:

	* gfortran.h: Add variable for 'ancestor' in struct gfc_omp_clauses.
	* openmp.c (gfc_match_omp_clauses): Parse device-modifiers 'device_num'
        and 'ancestor' in 'target device' clauses.
	* trans-openmp.c (gfc_trans_omp_clauses): Set OMP_CLAUSE_DEVICE_ANCESTOR.

gcc/ChangeLog:

	* gimplify.c (gimplify_scan_omp_clauses): Error handling. 'ancestor' only
	allowed on target constructs and only with particular other clauses.
	* omp-expand.c (expand_omp_target): Output of 'sorry, not supported' if
	'ancestor' is used.
	* omp-low.c (check_omp_nesting_restrictions): Error handling. No nested OpenMP
        structs when 'ancestor' is used.
	(scan_omp_1_stmt): No usage of OpenMP runtime routines in a target region when
	'ancestor' is used.
	* tree-pretty-print.c (dump_omp_clause): Append 'ancestor'.
	* tree.h (OMP_CLAUSE_DEVICE_ANCESTOR): Define macro.

gcc/testsuite/ChangeLog:

	* c-c++-common/gomp/target-device-1.c: New test.
	* c-c++-common/gomp/target-device-2.c: New test.
	* c-c++-common/gomp/target-device-ancestor-1.c: New test.
	* c-c++-common/gomp/target-device-ancestor-2.c: New test.
	* c-c++-common/gomp/target-device-ancestor-3.c: New test.
	* c-c++-common/gomp/target-device-ancestor-4.c: New test.
	* gfortran.dg/gomp/target-device-1.f90: New test.
	* gfortran.dg/gomp/target-device-2.f90: New test.
	* gfortran.dg/gomp/target-device-ancestor-1.f90: New test.
	* gfortran.dg/gomp/target-device-ancestor-2.f90: New test.
	* gfortran.dg/gomp/target-device-ancestor-3.f90: New test.
	* gfortran.dg/gomp/target-device-ancestor-4.f90: New test.

diff --git a/gcc/c/c-parser.c b/gcc/c/c-parser.c
index 9a56e0c..6c92d94 100644
--- a/gcc/c/c-parser.c
+++ b/gcc/c/c-parser.c
@@ -15864,37 +15864,81 @@ c_parser_omp_clause_map (c_parser *parser, tree list)
 }
 
 /* OpenMP 4.0:
-   device ( expression ) */
+>>>>>>>>>>>>>>>>>>>>>>>>>>>>>>>>>>>>>>>>>>>>>>>>>>>>>>>>>>>>>>>>>>>>>>>>>>>>>>>>>>>>>>>   device ( expression )
+
+   OpenMP 5.0:
+   device ( [device-modifier :] integer-expression )
+
+   device-modifier:
+     ancestor | device_num */
 
 static tree
 c_parser_omp_clause_device (c_parser *parser, tree list)
 {
   location_t clause_loc = c_parser_peek_token (parser)->location;
-  matching_parens parens;
-  if (parens.require_open (parser))
-    {
-      location_t expr_loc = c_parser_peek_token (parser)->location;
-      c_expr expr = c_parser_expr_no_commas (parser, NULL);
-      expr = convert_lvalue_to_rvalue (expr_loc, expr, false, true);
-      tree c, t = expr.value;
-      t = c_fully_fold (t, false, NULL);
+  location_t expr_loc;
+  c_expr expr;
+  tree c, t;
+  bool ancestor = false;
 
-      parens.skip_until_found_close (parser);
+  matching_parens parens;
+  if (!parens.require_open (parser))
+    return list;
 
-      if (!INTEGRAL_TYPE_P (TREE_TYPE (t)))
+  if (c_parser_next_token_is (parser, CPP_NAME)
+      && c_parser_peek_2nd_token (parser)->type == CPP_COLON)
+    {
+      c_token *tok = c_parser_peek_token (parser);
+      const char *p = IDENTIFIER_POINTER (tok->value);
+      if (strcmp ("ancestor", p) == 0)
 	{
-	  c_parser_error (parser, "expected integer expression");
+	  /* A requires directive with the reverse_offload clause must be
+	  specified.  */
+	  if ((omp_requires_mask & OMP_REQUIRES_REVERSE_OFFLOAD) == 0)
+	    {
+	      c_parser_error (parser, "a %<requires%> directive with the "
+				      "%<reverse_offload%> clause must be "
+				      "specified");
+	      parens.skip_until_found_close (parser);
+	      return list;
+	    }
+	  ancestor = true;
+	}
+      else if (strcmp ("device_num", p) == 0)
+	;
+      else
+	{
+	  error_at (tok->location, "expected %<ancestor%> or %<device_num%>");
+	  parens.skip_until_found_close (parser);
 	  return list;
 	}
+      c_parser_consume_token (parser);
+      c_parser_consume_token (parser);
+    }
 
-      check_no_duplicate_clause (list, OMP_CLAUSE_DEVICE, "device");
+  expr_loc = c_parser_peek_token (parser)->location;
+  expr = c_parser_expr_no_commas (parser, NULL);
+  expr = convert_lvalue_to_rvalue (expr_loc, expr, false, true);
+  t = expr.value;
+  t = c_fully_fold (t, false, NULL);
 
-      c = build_omp_clause (clause_loc, OMP_CLAUSE_DEVICE);
-      OMP_CLAUSE_DEVICE_ID (c) = t;
-      OMP_CLAUSE_CHAIN (c) = list;
-      list = c;
+  parens.skip_until_found_close (parser);
+
+  if (!INTEGRAL_TYPE_P (TREE_TYPE (t)))
+    {
+      c_parser_error (parser, "expected integer expression");
+      return list;
     }
 
+  check_no_duplicate_clause (list, OMP_CLAUSE_DEVICE, "device");
+
+  c = build_omp_clause (clause_loc, OMP_CLAUSE_DEVICE);
+
+  OMP_CLAUSE_DEVICE_ID (c) = t;
+  OMP_CLAUSE_CHAIN (c) = list;
+  OMP_CLAUSE_DEVICE_ANCESTOR (c) = ancestor;
+
+  list = c;
   return list;
 }
 
diff --git a/gcc/c/c-typeck.c b/gcc/c/c-typeck.c
index 5349ef1..b4d8d81 100644
--- a/gcc/c/c-typeck.c
+++ b/gcc/c/c-typeck.c
@@ -15139,6 +15139,22 @@ c_finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
 	case OMP_CLAUSE_COLLAPSE:
 	case OMP_CLAUSE_FINAL:
 	case OMP_CLAUSE_DEVICE:
+	  if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_DEVICE
+	      && OMP_CLAUSE_DEVICE_ANCESTOR (c))
+	    {
+	      t = OMP_CLAUSE_DEVICE_ID (c);
+	      if (TREE_CODE (t) == INTEGER_CST
+		  && wi::to_widest (t) != 1)
+		{
+		  error_at (OMP_CLAUSE_LOCATION (c),
+			    "the %<device%> clause expression must evaluate to "
+			    "%<1%>");
+		  remove = true;
+		  break;
+		}
+	    }
+	  /* FALLTHRU */
+
 	case OMP_CLAUSE_DIST_SCHEDULE:
 	case OMP_CLAUSE_PARALLEL:
 	case OMP_CLAUSE_FOR:
diff --git a/gcc/cp/parser.c b/gcc/cp/parser.c
index 93698aa..2d876ce 100644
--- a/gcc/cp/parser.c
+++ b/gcc/cp/parser.c
@@ -38536,18 +38536,57 @@ cp_parser_omp_clause_map (cp_parser *parser, tree list)
 }
 
 /* OpenMP 4.0:
-   device ( expression ) */
+   device ( expression )
+
+   OpenMP 5.0:
+   device ( [device-modifier :] integer-expression )
+
+   device-modifier:
+     ancestor | device_num */
 
 static tree
 cp_parser_omp_clause_device (cp_parser *parser, tree list,
 			     location_t location)
 {
   tree t, c;
+  bool ancestor = false;
 
   matching_parens parens;
   if (!parens.require_open (parser))
     return list;
 
+  if (cp_lexer_next_token_is (parser->lexer, CPP_NAME)
+      && cp_lexer_nth_token_is (parser->lexer, 2, CPP_COLON))
+    {
+      cp_token *tok = cp_lexer_peek_token (parser->lexer);
+      const char *p = IDENTIFIER_POINTER (tok->u.value);
+      if (strcmp ("ancestor", p) == 0)
+	{
+	  ancestor = true;
+
+	  /* A requires directive with the reverse_offload clause must be
+	  specified.  */
+	  if ((omp_requires_mask & OMP_REQUIRES_REVERSE_OFFLOAD) == 0)
+	    {
+	      error_at (tok->location, "a %<requires%> directive with the "
+				       "%<reverse_offload%> clause must be "
+				       "specified");
+	      cp_parser_skip_to_closing_parenthesis (parser, true, false, true);
+	      return list;
+	    }
+	}
+      else if (strcmp ("device_num", p) == 0)
+	;
+      else
+	{
+	  error_at (tok->location, "expected %<ancestor%> or %<device_num%>");
+	  cp_parser_skip_to_closing_parenthesis (parser, true, false, true);
+	  return list;
+	}
+      cp_lexer_consume_token (parser->lexer);
+      cp_lexer_consume_token (parser->lexer);
+    }
+
   t = cp_parser_assignment_expression (parser);
 
   if (t == error_mark_node
@@ -38562,6 +38601,7 @@ cp_parser_omp_clause_device (cp_parser *parser, tree list,
   c = build_omp_clause (location, OMP_CLAUSE_DEVICE);
   OMP_CLAUSE_DEVICE_ID (c) = t;
   OMP_CLAUSE_CHAIN (c) = list;
+  OMP_CLAUSE_DEVICE_ANCESTOR (c) = ancestor;
 
   return c;
 }
diff --git a/gcc/cp/semantics.c b/gcc/cp/semantics.c
index b080259..0651f5a 100644
--- a/gcc/cp/semantics.c
+++ b/gcc/cp/semantics.c
@@ -7334,6 +7334,15 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
 			"%<device%> id must be integral");
 	      remove = true;
 	    }
+	  else if (OMP_CLAUSE_DEVICE_ANCESTOR (c)
+		   && TREE_CODE (t) == INTEGER_CST
+		   && wi::to_widest (t) != 1)
+	    {
+	      error_at (OMP_CLAUSE_LOCATION (c),
+			"the %<device%> clause expression must evaluate to "
+			"%<1%>");
+	      remove = true;
+	    }
 	  else
 	    {
 	      t = mark_rvalue_use (t);
diff --git a/gcc/fortran/gfortran.h b/gcc/fortran/gfortran.h
index f4a50d7..b428f06 100644
--- a/gcc/fortran/gfortran.h
+++ b/gcc/fortran/gfortran.h
@@ -1473,6 +1473,7 @@ typedef struct gfc_omp_clauses
   enum gfc_omp_sched_kind dist_sched_kind;
   struct gfc_expr *dist_chunk_size;
   const char *critical_name;
+  bool ancestor;
 
   /* OpenACC. */
   struct gfc_expr *async_expr;
diff --git a/gcc/fortran/openmp.c b/gcc/fortran/openmp.c
index 357a1e1..8cf59af 100644
--- a/gcc/fortran/openmp.c
+++ b/gcc/fortran/openmp.c
@@ -1714,8 +1714,56 @@ gfc_match_omp_clauses (gfc_omp_clauses **cp, const omp_mask mask,
 	  if ((mask & OMP_CLAUSE_DEVICE)
 	      && !openacc
 	      && c->device == NULL
-	      && gfc_match ("device ( %e )", &c->device) == MATCH_YES)
-	    continue;
+	      && gfc_match ("device ( ") == MATCH_YES)
+	    {
+	      c->ancestor = false;
+	      if (gfc_match ("device_num : ") == MATCH_YES)
+		{
+		  if (gfc_match ("%e )", &c->device) != MATCH_YES)
+		    {
+		      gfc_error ("Expected integer expression at %C");
+		      break;
+		    }
+		}
+	      else if (gfc_match ("ancestor : ") == MATCH_YES)
+		{
+		  c->ancestor = true;
+		  if (!(gfc_current_ns->omp_requires & OMP_REQ_REVERSE_OFFLOAD))
+		    {
+		      gfc_error ("a %<requires%> directive with the "
+				 "%<reverse_offload%> clause must be "
+				 "specified at %C");
+		      break;
+		    }
+		  locus old_loc2 = gfc_current_locus;
+		  if (gfc_match ("%e )", &c->device) == MATCH_YES)
+		    {
+		      int device = 0;
+		      if (!gfc_extract_int (c->device, &device) && device != 1)
+		      {
+			gfc_current_locus = old_loc2;
+			gfc_error ("the %<device%> clause expression must "
+				   "evaluate to %<1%> at %C");
+			break;
+		      }
+		    }
+		  else
+		    {
+		      gfc_error ("Expected integer expression at %C");
+		      break;
+		    }
+		}
+	      else if (gfc_match ("%e )", &c->device) == MATCH_YES)
+		{
+		}
+	      else
+		{
+		  gfc_error ("Expected integer expression or a single device-"
+			      "modifier %<device_num%> or %<ancestor%> at %C");
+		  break;
+		}
+	      continue;
+	    }
 	  if ((mask & OMP_CLAUSE_DEVICE)
 	      && openacc
 	      && gfc_match ("device ( ") == MATCH_YES
diff --git a/gcc/fortran/trans-openmp.c b/gcc/fortran/trans-openmp.c
index ace4faf..321e7d3 100644
--- a/gcc/fortran/trans-openmp.c
+++ b/gcc/fortran/trans-openmp.c
@@ -3947,6 +3947,10 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses,
 
       c = build_omp_clause (gfc_get_location (&where), OMP_CLAUSE_DEVICE);
       OMP_CLAUSE_DEVICE_ID (c) = device;
+
+      if (clauses->ancestor)
+	OMP_CLAUSE_DEVICE_ANCESTOR (c) = 1;
+
       omp_clauses = gfc_trans_add_clause (c, omp_clauses);
     }
 
diff --git a/gcc/gimplify.c b/gcc/gimplify.c
index 75a4a9d..c6d20cd 100644
--- a/gcc/gimplify.c
+++ b/gcc/gimplify.c
@@ -10088,6 +10088,36 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
 	case OMP_CLAUSE_THREAD_LIMIT:
 	case OMP_CLAUSE_DIST_SCHEDULE:
 	case OMP_CLAUSE_DEVICE:
+	  if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_DEVICE
+	      && OMP_CLAUSE_DEVICE_ANCESTOR (c))
+	    {
+	      if (code != OMP_TARGET)
+		{
+		    error_at (OMP_CLAUSE_LOCATION (c),
+			      "%<device%> clause with %<ancestor%> is only "
+			      "allowed on %<target%> construct");
+		    remove = true;
+		}
+
+	      tree clauses = *orig_list_p;
+	      for (; clauses ; clauses = OMP_CLAUSE_CHAIN (clauses))
+		if (OMP_CLAUSE_CODE (clauses) != OMP_CLAUSE_DEVICE
+		    && OMP_CLAUSE_CODE (clauses) != OMP_CLAUSE_FIRSTPRIVATE
+		    && OMP_CLAUSE_CODE (clauses) != OMP_CLAUSE_PRIVATE
+		    && OMP_CLAUSE_CODE (clauses) != OMP_CLAUSE_DEFAULTMAP
+		    && OMP_CLAUSE_CODE (clauses) != OMP_CLAUSE_MAP
+		   )
+		  {
+		    error_at (OMP_CLAUSE_LOCATION (c),
+			      "with %<ancestor%>, only the %<device%>, "
+			      "%<firstprivate%>, %<private%>, %<defaultmap%>, "
+			      "and %<map%> clauses may appear on the "
+			      "construct");
+		    remove = true;
+		  }
+	    }
+	  /* Fall through.  */
+
 	case OMP_CLAUSE_PRIORITY:
 	case OMP_CLAUSE_GRAINSIZE:
 	case OMP_CLAUSE_NUM_TASKS:
diff --git a/gcc/omp-expand.c b/gcc/omp-expand.c
index 9fd1c65..a9096a1 100644
--- a/gcc/omp-expand.c
+++ b/gcc/omp-expand.c
@@ -9605,6 +9605,8 @@ expand_omp_target (struct omp_region *region)
 	{
 	  device = OMP_CLAUSE_DEVICE_ID (c);
 	  device_loc = OMP_CLAUSE_LOCATION (c);
+	  if (OMP_CLAUSE_DEVICE_ANCESTOR (c))
+	    sorry_at (device_loc, "%<ancestor%> not yet supported");
 	}
       else
 	{
diff --git a/gcc/omp-low.c b/gcc/omp-low.c
index e7049c8..5e2f9d2 100644
--- a/gcc/omp-low.c
+++ b/gcc/omp-low.c
@@ -3101,6 +3101,16 @@ check_omp_nesting_restrictions (gimple *stmt, omp_context *ctx)
       if (gimple_code (ctx->stmt) == GIMPLE_OMP_TARGET
 	  && gimple_omp_target_kind (ctx->stmt) == GF_OMP_TARGET_KIND_REGION)
 	{
+	  c = omp_find_clause (gimple_omp_target_clauses (ctx->stmt),
+			       OMP_CLAUSE_DEVICE);
+	  if (c && OMP_CLAUSE_DEVICE_ANCESTOR (c))
+	    {
+	      error_at (gimple_location (stmt),
+			"OpenMP constructs are not allowed in target region "
+			"with %<ancestor%>");
+	      return false;
+	    }
+
 	  if (gimple_code (stmt) == GIMPLE_OMP_TEAMS && !ctx->teams_nested_p)
 	    ctx->teams_nested_p = true;
 	  else
@@ -4001,6 +4011,20 @@ scan_omp_1_stmt (gimple_stmt_iterator *gsi, bool *handled_ops_p,
 			    "OpenMP runtime API call %qD in a region with "
 			    "%<order(concurrent)%> clause", fndecl);
 		}
+	      if (gimple_code (ctx->stmt) == GIMPLE_OMP_TARGET
+		  && gimple_omp_target_kind (ctx->stmt) ==
+		  GF_OMP_TARGET_KIND_REGION)
+		{
+		  tree c =
+		    omp_find_clause (gimple_omp_target_clauses (ctx->stmt),
+				     OMP_CLAUSE_DEVICE);
+		  if (c && OMP_CLAUSE_DEVICE_ANCESTOR (c))
+		    {
+		      error_at (gimple_location (stmt),
+				"OpenMP runtime API call %qD in a region with "
+				"%<device(ancestor)%> clause", fndecl);
+		    }
+		}
 	    }
 	}
     }
diff --git a/gcc/testsuite/c-c++-common/gomp/target-device-1.c b/gcc/testsuite/c-c++-common/gomp/target-device-1.c
new file mode 100644
index 0000000..dafa643
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/gomp/target-device-1.c
@@ -0,0 +1,34 @@
+/* { dg-do compile } */
+
+void
+foo (void)
+{
+  /* Test to ensure that 'device_num' is parsed correctly in device clauses. */
+
+  int n;
+
+  #pragma omp target device (1)
+  ;
+
+  #pragma omp target device (n)
+  ;
+
+  #pragma omp target device (n + 1)
+  ;
+
+  #pragma omp target device (device_num : 1)
+  ;
+
+  #pragma omp target device (device_num : n)
+  ;
+
+  #pragma omp target device (device_num : n + 1)
+  ;
+
+  #pragma omp target device (invalid : 1) /* { dg-error "expected 'ancestor' or 'device_num'" "" { target *-*-* } } */
+  /* { dg-error "expected '\\)' before 'invalid'" "" { target c } .-1 } */
+  ;
+
+  #pragma omp target device (device_num : n, n) /* { dg-error "expected '\\)' before ','" } */
+  ;
+}
diff --git a/gcc/testsuite/c-c++-common/gomp/target-device-2.c b/gcc/testsuite/c-c++-common/gomp/target-device-2.c
new file mode 100644
index 0000000..b711ea1
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/gomp/target-device-2.c
@@ -0,0 +1,14 @@
+/* { dg-do compile } */
+/* { dg-additional-options "-fdump-tree-original" } */
+
+  /* Test to ensure that device-modifier 'device_num' is parsed correctly in
+     device clauses. */
+
+void
+foo (void)
+{
+  #pragma omp target device (device_num : 42)
+  ;
+}
+
+/* { dg-final { scan-tree-dump "pragma omp target \[^\n\r)]*device\\(42\\)" "original" } } */
diff --git a/gcc/testsuite/c-c++-common/gomp/target-device-ancestor-1.c b/gcc/testsuite/c-c++-common/gomp/target-device-ancestor-1.c
new file mode 100644
index 0000000..11d54f2
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/gomp/target-device-ancestor-1.c
@@ -0,0 +1,11 @@
+/* { dg-do compile } */
+
+void
+foo (void)
+{
+  /* Ensure that a 'requires' directive with the 'reverse_offload' clause was
+     specified.  */
+
+  #pragma omp target device (ancestor : 1) /* { dg-error "a 'requires' directive with the 'reverse_offload' clause must be specified" } */
+  ;
+}
diff --git a/gcc/testsuite/c-c++-common/gomp/target-device-ancestor-2.c b/gcc/testsuite/c-c++-common/gomp/target-device-ancestor-2.c
new file mode 100644
index 0000000..b2067e3
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/gomp/target-device-ancestor-2.c
@@ -0,0 +1,84 @@
+/* { dg-do compile } */
+
+#pragma omp requires reverse_offload /* { dg-message "sorry, unimplemented: 'reverse_offload' clause on 'requires' directive not supported yet" } */
+
+void
+foo (void)
+{
+  /* The following test is marked with 'xfail' because a previous 'sorry' from
+     'reverse_offload' suppresses the 'sorry' for 'ancestor'.  */
+  #pragma omp target device (ancestor: 1) /* { dg-message "" "sorry, unimplemented: 'ancestor' not yet supported" { xfail *-*-* } } */
+  ;
+
+
+  /* Ensure that the integer expression in the 'device' clause for
+     device-modifier 'ancestor' evaluates to '1' in case of a constant.  */
+
+  #pragma omp target device (ancestor : 1)
+  ;
+  #pragma omp target device (ancestor : 42) /* { dg-error "the 'device' clause expression must evaluate to '1'" } */
+  ;
+
+  int n;
+  #pragma omp target device (ancestor : n) /* { dg-message "" "sorry, unimplemented: 'ancestor' not yet supported" { xfail *-*-* } } */
+  ;
+  #pragma omp target device (ancestor : n + 1) /* { dg-message "" "sorry, unimplemented: 'ancestor' not yet supported" { xfail *-*-* } } */
+  ;
+
+
+  /* Ensure that only one 'device' clause appears on the construct.  */
+
+  #pragma omp target device (17) device (42) /* { dg-error "too many 'device' clauses" } */
+  ;
+
+
+  /* Ensure that with 'ancestor' only the 'device', 'firstprivate', 'private',
+     'defaultmap', and 'map' clauses appear on the construct.  */
+
+  #pragma omp target nowait device (ancestor: 1) /* { dg-error "with 'ancestor', only the 'device', 'firstprivate', 'private', 'defaultmap', and 'map' clauses may appear on the construct" } */
+  ;
+  #pragma omp target device (ancestor: 1) nowait /* { dg-error "with 'ancestor', only the 'device', 'firstprivate', 'private', 'defaultmap', and 'map' clauses may appear on the construct" } */
+  ;
+  #pragma omp target nowait device (42)
+  ;
+  #pragma omp target nowait device (device_num: 42)
+  ;
+
+  int a, b, c;
+  #pragma omp target device (ancestor: 1) firstprivate (a) private (b) defaultmap (none) map (c)
+  ;
+
+
+  /* Ensure that 'ancestor' is only used with 'target' constructs (not with
+     'target data', 'target update' etc.).  */
+
+  #pragma omp target data map (a) device (ancestor: 1) /* { dg-error "'device' clause with 'ancestor' is only allowed on 'target' construct" } */
+  ;
+  #pragma omp target enter data map (to: a) device (ancestor: 1) /* { dg-error "'device' clause with 'ancestor' is only allowed on 'target' construct" } */
+  #pragma omp target exit data map (from: a) device (ancestor: 1) /* { dg-error "'device' clause with 'ancestor' is only allowed on 'target' construct" } */
+  #pragma omp target update to (a) device (ancestor: 1) /* { dg-error "'device' clause with 'ancestor' is only allowed on 'target' construct" "" { target *-*-* } } */
+  /* { dg-error "with 'ancestor', only the 'device', 'firstprivate', 'private', 'defaultmap', and 'map' clauses may appear on the construct" "" { target *-*-* } .-1 } */
+
+
+  /* Ensure that no OpenMP constructs appear inside target regions with 
+     'ancestor'.  */
+
+  #pragma omp target device (ancestor: 1)
+    {
+      #pragma omp teams /* { dg-error "OpenMP constructs are not allowed in target region with 'ancestor'" } */
+      ;
+    }
+
+  #pragma omp target device (device_num: 1) 
+    {
+      #pragma omp teams
+      ;
+    }
+
+  #pragma omp target device (1) 
+    {
+      #pragma omp teams
+      ;
+    }
+
+}
diff --git a/gcc/testsuite/c-c++-common/gomp/target-device-ancestor-3.c b/gcc/testsuite/c-c++-common/gomp/target-device-ancestor-3.c
new file mode 100644
index 0000000..5e3a478
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/gomp/target-device-ancestor-3.c
@@ -0,0 +1,37 @@
+#ifdef __cplusplus
+extern "C" {
+#endif
+
+int omp_get_num_teams (void);
+
+#ifdef __cplusplus
+}
+#endif
+
+/* { dg-do compile } */
+
+#pragma omp requires reverse_offload /* { dg-message "sorry, unimplemented: 'reverse_offload' clause on 'requires' directive not supported yet" } */
+
+void
+foo (void)
+{
+  /* Ensure that no calls to OpenMP API runtime routines are allowed inside the
+     corresponding target region.  */
+
+  int a;
+
+  #pragma omp target device (ancestor: 1)
+    {
+      a = omp_get_num_teams (); /* { dg-error "OpenMP runtime API call '\[^\n\r]*omp_get_num_teams\[^\n\r]*' in a region with 'device\\(ancestor\\)' clause" }  */
+    }
+
+  #pragma omp target device (device_num: 1)
+    {
+      a = omp_get_num_teams ();
+    }
+
+  #pragma omp target device (1)
+    {
+      a = omp_get_num_teams ();
+    }
+}
diff --git a/gcc/testsuite/c-c++-common/gomp/target-device-ancestor-4.c b/gcc/testsuite/c-c++-common/gomp/target-device-ancestor-4.c
new file mode 100644
index 0000000..b4b5620
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/gomp/target-device-ancestor-4.c
@@ -0,0 +1,17 @@
+/* { dg-do compile } */
+/* { dg-additional-options "-fdump-tree-original" } */
+
+  /* Test to ensure that device-modifier 'ancestor' is parsed correctly in
+     device clauses. */
+
+#pragma omp requires reverse_offload /* { dg-message "sorry, unimplemented: 'reverse_offload' clause on 'requires' directive not supported yet" } */
+
+void
+foo (void)
+{
+  #pragma omp target device (ancestor: 1) /* { dg-message "" "sorry, unimplemented: 'ancestor' not yet supported" { xfail *-*-* } } */
+  ;
+
+}
+
+/* { dg-final { scan-tree-dump "pragma omp target \[^\n\r)]*device\\(ancestor:1\\)" "original" } } */
diff --git a/gcc/testsuite/gfortran.dg/gomp/target-device-1.f90 b/gcc/testsuite/gfortran.dg/gomp/target-device-1.f90
new file mode 100644
index 0000000..20b9755
--- /dev/null
+++ b/gcc/testsuite/gfortran.dg/gomp/target-device-1.f90
@@ -0,0 +1,67 @@
+! { dg-do compile }
+
+implicit none
+
+integer :: n
+
+!$omp target device (1)
+!$omp end target
+
+!$omp target device (n)
+!$omp end target
+
+!$omp target device (n + 1)
+!$omp end target
+
+!$omp target device (device_num : 1)
+!$omp end target
+
+!$omp target device (device_num : n)
+!$omp end target
+
+!$omp target device (device_num : n + 1)
+!$omp end target
+
+!$omp target device (invalid : 1)  ! { dg-error "Expected integer expression or a single device-modifier 'device_num' or 'ancestor' at" }
+! !$omp end target
+
+!$omp target device ( : 1)  ! { dg-error "Expected integer expression or a single device-modifier 'device_num' or 'ancestor' at" }
+! !$omp end target
+
+!$omp target device ( , : 1)  ! { dg-error "Expected integer expression or a single device-modifier 'device_num' or 'ancestor' at" }
+! !$omp end target
+
+!$omp target device (ancestor, device_num : 1)  ! { dg-error "Expected integer expression or a single device-modifier 'device_num' or 'ancestor' at" }
+! !$omp end target
+
+!$omp target device (ancestor, device_num, ancestor : 1)  ! { dg-error "Expected integer expression or a single device-modifier 'device_num' or 'ancestor' at" }
+! !$omp end target
+
+!$omp target device (device_num device_num : 1)  ! { dg-error "Expected integer expression or a single device-modifier 'device_num' or 'ancestor' at" }
+! !$omp end target
+
+!$omp target device (ancestor device_num : 1)  ! { dg-error "Expected integer expression or a single device-modifier 'device_num' or 'ancestor' at" }
+! !$omp end target
+
+!$omp target device (device_num, invalid : 1)  ! { dg-error "Expected integer expression or a single device-modifier 'device_num' or 'ancestor' at" }
+! !$omp end target
+
+!$omp target device (ancestor, invalid : 1)  ! { dg-error "Expected integer expression or a single device-modifier 'device_num' or 'ancestor' at" }
+! !$omp end target
+
+!$omp target device (ancestor, , , : 1)  ! { dg-error "Expected integer expression or a single device-modifier 'device_num' or 'ancestor' at" }
+! !$omp end target
+
+!$omp target device (invalid, ancestor : 1)  ! { dg-error "xpected integer expression or a single device-modifier 'device_num' or 'ancestor' at" }
+! !$omp end target
+
+!$omp target device (invalid, invalid, ancestor : 1)  ! { dg-error "xpected integer expression or a single device-modifier 'device_num' or 'ancestor' at" }
+! !$omp end target
+
+!$omp target device (device_num invalid : 1)  ! { dg-error "Expected integer expression or a single device-modifier 'device_num' or 'ancestor' at" }
+! !$omp end target
+
+!$omp target device (device_num : n, n)  ! { dg-error "Expected integer expression" }
+! !$omp end target
+
+end
diff --git a/gcc/testsuite/gfortran.dg/gomp/target-device-2.f90 b/gcc/testsuite/gfortran.dg/gomp/target-device-2.f90
new file mode 100644
index 0000000..133b805
--- /dev/null
+++ b/gcc/testsuite/gfortran.dg/gomp/target-device-2.f90
@@ -0,0 +1,12 @@
+! { dg-do compile }
+! { dg-additional-options "-fdump-tree-original" }
+
+! Test to ensure that device-modifier 'device_num' is parsed correctly in
+! device clauses.
+
+!$omp target device (device_num : 42)
+!$omp end target
+
+end
+
+! { dg-final { scan-tree-dump "pragma omp target \[^\n\r)]*device\\(42\\)" "original" } }
diff --git a/gcc/testsuite/gfortran.dg/gomp/target-device-ancestor-1.f90 b/gcc/testsuite/gfortran.dg/gomp/target-device-ancestor-1.f90
new file mode 100644
index 0000000..72a4054
--- /dev/null
+++ b/gcc/testsuite/gfortran.dg/gomp/target-device-ancestor-1.f90
@@ -0,0 +1,9 @@
+! { dg-do compile }
+
+! Ensure that a 'requires' directive with the 'reverse_offload' clause was
+! specified.
+
+!$omp target device (ancestor:1)  ! { dg-error "a 'requires' directive with the 'reverse_offload' clause must be specified" }
+! !$omp end target
+
+end
\ No newline at end of file
diff --git a/gcc/testsuite/gfortran.dg/gomp/target-device-ancestor-2.f90 b/gcc/testsuite/gfortran.dg/gomp/target-device-ancestor-2.f90
new file mode 100644
index 0000000..117a1d0
--- /dev/null
+++ b/gcc/testsuite/gfortran.dg/gomp/target-device-ancestor-2.f90
@@ -0,0 +1,92 @@
+! { dg-do compile }
+
+implicit none
+
+integer :: a, b, c
+
+!$omp requires reverse_offload  ! { dg-error "Sorry, 'reverse_offload' clause at \\(1\\) on REQUIRES directive is not yet supported" }
+
+
+! The following test case is marked with 'xfail' because a previous 'sorry' from
+! 'reverse_offload' suppresses the 'sorry' for 'ancestor'.
+
+!$omp target device (ancestor: 1)  ! { dg-message "" "sorry, unimplemented: 'ancestor' not yet supported" { xfail *-*-* } }
+!$omp end target
+
+!$omp target device (ancestor : a)  ! { dg-message "" "sorry, unimplemented: 'ancestor' not yet supported" { xfail *-*-* } }
+!$omp end target
+
+!$omp target device (ancestor : a + 1)  ! { dg-message "" "sorry, unimplemented: 'ancestor' not yet supported" { xfail *-*-* } }
+!$omp end target
+
+
+! Ensure that the integer expression in the 'device' clause for
+! device-modifier 'ancestor' evaluates to '1' in case of a constant.
+
+!$omp target device (ancestor: 42)  ! { dg-error "the 'device' clause expression must evaluate to '1'" }
+! !$omp end target
+
+!$omp target device (device_num:42)
+!$omp end target
+
+!$omp target device (42)
+!$omp end target
+
+
+! Ensure that no OpenMP constructs appear inside target regions with 'ancestor'.
+! The following test case is marked with 'xfail' because a previous 'sorry' from
+! 'reverse_offload' suppresses the 'sorry' for 'ancestor'.
+
+!$omp target device (ancestor: 1)
+  !$omp teams  ! { dg-error "" "OpenMP constructs are not allowed in target region with 'ancestor'" { xfail *-*-* } }
+  !$omp end teams
+!$omp end target
+
+!$omp target device (device_num: 1)
+  !$omp teams
+  !$omp end teams
+!$omp end target
+
+!$omp target device (1)
+  !$omp teams
+  !$omp end teams
+!$omp end target
+
+
+! Ensure that with 'ancestor' only the 'device', 'firstprivate', 'private',
+! 'defaultmap', and 'map' clauses appear on the construct.
+! The following test case is marked with 'xfail' because a previous 'sorry' from
+! 'reverse_offload' suppresses the 'sorry' for 'ancestor'.
+
+!$omp target nowait device (ancestor: 1)  ! { dg-error "" "with 'ancestor', only the 'device', 'firstprivate', 'private', 'defaultmap', and 'map' clauses may appear on the construct" { xfail *-*-* } }
+!$omp end target
+
+!$omp target device (ancestor: 1) nowait  ! { dg-error "" "with 'ancestor', only the 'device', 'firstprivate', 'private', 'defaultmap', and 'map' clauses may appear on the construct" { xfail *-*-* } }
+!$omp end target
+
+!$omp target nowait device (device_num: 1)
+!$omp end target
+
+!$omp target nowait device (1)
+!$omp end target
+
+!$omp target device (ancestor: 1) firstprivate (a) private (b) defaultmap (none) map (c)
+!$omp end target
+
+
+! Ensure that 'ancestor' is only used with 'target' constructs (not with
+! 'target data', 'target update' etc.).
+! The following test case is marked with 'xfail' because a previous 'sorry' from
+! 'reverse_offload' suppresses the 'sorry' for 'ancestor'.
+
+!$omp target data map (a) device (ancestor: 1)  ! { dg-error "" "'device' clause with 'ancestor' is only allowed on 'target' construct" { xfail *-*-* } }
+!$omp end target data
+
+!$omp target enter data map (to: a) device (ancestor: 1)  ! { dg-error "" "'device' clause with 'ancestor' is only allowed on 'target' construct" { xfail *-*-* } }
+!$omp target exit data map (from: a) device (ancestor: 1)  ! { dg-error "" "'device' clause with 'ancestor' is only allowed on 'target' construct" { xfail *-*-* } }
+
+!$omp target update to (a) device (ancestor: 1)  ! { dg-error "'device' clause with 'ancestor' is only allowed on 'target' construct" "" { xfail *-*-* } }
+! { dg-error "with 'ancestor', only the 'device', 'firstprivate', 'private', 'defaultmap', and 'map' clauses may appear on the construct" "" { xfail *-*-* } .-1 }
+
+
+end
\ No newline at end of file
diff --git a/gcc/testsuite/gfortran.dg/gomp/target-device-ancestor-3.f90 b/gcc/testsuite/gfortran.dg/gomp/target-device-ancestor-3.f90
new file mode 100644
index 0000000..f1145bd
--- /dev/null
+++ b/gcc/testsuite/gfortran.dg/gomp/target-device-ancestor-3.f90
@@ -0,0 +1,33 @@
+! { dg-do compile }
+
+! This testcase ensure that no calls to OpenMP API runtime routines are allowed
+! inside the corresponding target region.
+
+module my_omp_mod
+ use iso_c_binding
+ interface
+   integer function omp_get_thread_num ()
+   end
+ end interface
+end
+
+subroutine f1 ()
+  use my_omp_mod
+  implicit none
+  integer :: n
+
+  !$omp requires reverse_offload  ! { dg-error "Sorry, 'reverse_offload' clause at \\(1\\) on REQUIRES directive is not yet supported" }
+
+  !$omp target device (ancestor : 1)
+    n = omp_get_thread_num ()  ! { dg-error "" "OpenMP runtime API call 'omp_get_thread_num' in a region with 'device\\(ancestor\\)' clause" { xfail *-*-* } }
+  !$omp end target
+
+  !$omp target device (device_num : 1)
+    n = omp_get_thread_num ()
+  !$omp end target
+
+  !$omp target device (1)
+    n = omp_get_thread_num ()
+  !$omp end target
+
+end
\ No newline at end of file
diff --git a/gcc/testsuite/gfortran.dg/gomp/target-device-ancestor-4.f90 b/gcc/testsuite/gfortran.dg/gomp/target-device-ancestor-4.f90
new file mode 100644
index 0000000..540b3d0
--- /dev/null
+++ b/gcc/testsuite/gfortran.dg/gomp/target-device-ancestor-4.f90
@@ -0,0 +1,14 @@
+! { dg-do compile }
+! { dg-additional-options "-fdump-tree-original" }
+
+! Test to ensure that device-modifier 'ancestor' is parsed correctly in
+! device clauses.
+
+!$omp requires reverse_offload  ! { dg-error "Sorry, 'reverse_offload' clause at \\(1\\) on REQUIRES directive is not yet supported" }
+
+!$omp target device (ancestor : 1)  ! { dg-message "" "sorry, unimplemented: 'ancestor' not yet supported" { xfail *-*-* } }
+!$omp end target
+
+end
+
+! { dg-final { scan-tree-dump "pragma omp target \[^\n\r)]*device\\(ancestor:1\\)" "original" } }
diff --git a/gcc/tree-pretty-print.c b/gcc/tree-pretty-print.c
index fde07df..042b44a 100644
--- a/gcc/tree-pretty-print.c
+++ b/gcc/tree-pretty-print.c
@@ -986,6 +986,8 @@ dump_omp_clause (pretty_printer *pp, tree clause, int spc, dump_flags_t flags)
 
     case OMP_CLAUSE_DEVICE:
       pp_string (pp, "device(");
+      if (OMP_CLAUSE_DEVICE_ANCESTOR (clause))
+	pp_string (pp, "ancestor:");
       dump_generic_node (pp, OMP_CLAUSE_DEVICE_ID (clause),
 			 spc, flags, false);
       pp_right_paren (pp);
diff --git a/gcc/tree.h b/gcc/tree.h
index 8bdf16d..1988a11 100644
--- a/gcc/tree.h
+++ b/gcc/tree.h
@@ -1673,6 +1673,10 @@ class auto_suppress_location_wrappers
 #define OMP_CLAUSE_DEVICE_TYPE_KIND(NODE) \
   (OMP_CLAUSE_SUBCODE_CHECK (NODE, OMP_CLAUSE_DEVICE_TYPE)->omp_clause.subcode.device_type_kind)
 
+/* True if there is a device clause with a device-modifier 'ancestor'.  */
+#define OMP_CLAUSE_DEVICE_ANCESTOR(NODE) \
+  (OMP_CLAUSE_SUBCODE_CHECK (NODE, OMP_CLAUSE_DEVICE)->base.public_flag)
+
 #define OMP_CLAUSE_COLLAPSE_EXPR(NODE) \
   OMP_CLAUSE_OPERAND (OMP_CLAUSE_SUBCODE_CHECK (NODE, OMP_CLAUSE_COLLAPSE), 0)
 #define OMP_CLAUSE_COLLAPSE_ITERVAR(NODE) \


More information about the Gcc-patches mailing list