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
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) \
