* Re: [Patch v2] C, C++, Fortran, OpenMP: Add support for device-modifiers for 'omp target device'
[not found] ` <20210823174733.GY2380545@tucnak>
@ 2021-08-25 10:14 ` Marcel Vollweiler
2021-08-30 6:26 ` Jakub Jelinek
0 siblings, 1 reply; 8+ messages in thread
From: Marcel Vollweiler @ 2021-08-25 10:14 UTC (permalink / raw)
To: Jakub Jelinek; +Cc: gcc-patches, fortran, Tobias Burnus
[-- Attachment #1: Type: text/plain, Size: 9603 bytes --]
Hi Jakub,
I applied all your suggested changes and checked for no test regressions
on x86_64-linux with nvptx offloading. The revised patch is attached.
Do you think that it's ok to commit the code?
Thanks,
Marcel
Am 23.08.2021 um 19:47 schrieb Jakub Jelinek:
> On Fri, Aug 20, 2021 at 09:18:32PM +0200, Marcel Vollweiler wrote:
>
>> --- 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 )
>
> Please remove all the >>>>>s.
>> +
>> + OpenMP 5.0:
>> + device ( [device-modifier :] integer-expression )
>> +
>> + device-modifier:
>> + ancestor | device_num */
>>
>
>> + /* 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");
>
> [BI think this diagnostics is confusing, it tells the user that it has to
> do something but doesn't tell why. It is also not a parser error.
> So I think it should be instead
> error_at (tok->location, "%<ancestor%> device modifier not "
> "preceded by %<requires%> directive "
> "with %<reverse_offload%> clause");
>
>> + parens.skip_until_found_close (parser);
>> + return list;
>> + }
>> + ancestor = true;
>> + }
>
>> + 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 */
>
> For the C FE, I'd suggest to move this to the c_parser_omp_clause_device
> routine like other similar checking is done there too. And you can use
> if (TREE_CODE (t) == INTEGER_CST && !integer_onep (t))
>> + error_at (tok->location, "a %<requires%> directive with the "
>
>> + "%<reverse_offload%> clause must be "
>> + "specified");
>
> See above.
>
>> @@ -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;
>
> But in C++ the INTEGER_CST checking shouldn't be done here, because
> the argument could be type or value dependent.
>
>> --- 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)
>
> !integer_onep (t)
>
>> + 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");
>
> See above.
>
>> + else if (gfc_match ("%e )", &c->device) == MATCH_YES)
>> + {
>> + }
>> + else
>
> Better != MATCH_YES and drop the {} 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
>
>> + 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;
>> + }
>
> Formatting, {/} are correctly indented, but error_at and remove should be
> indented 2 columns to the right from that, not 4 columns.
> Also, it should have break; there too, so that it doesn't fallthrough
> to the next one:
>> +
>> + 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. */
>> +
>
>> @@ -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)
>
> Formatting. Neither == nor = should be at the end of lines. So,
> 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);
>
> And probably use a tree tgt_clauses = gimple_omp_target_clauses (ctx->stmt);
> temporary to make tree c = omp_find_clause (tgt_clauses, OMP_CLAUSE_DEVICE);
> fit nicely.
>
>> + 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);
>> + }
>
> Single statement in if body shouldn't be wrapped with {}s.
>> + }
>> }
>> }
>> }
>> 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;
>
> Better use
> foo (int n)
> such that it isn't an uninitialized use. Or initialize n to something.
>
>> +
>> + #pragma omp target device (1)
>> + ;
>> +
>> + #pragma omp target device (n)
>> + ;
>> +
>> + #pragma omp target device (n + 1)
>> + ;
>> +
>> +end
>> \ No newline at end of file
>
> Please avoid these in all the tests.
>
> Otherwise LGTM.
>
> Jakub
>
-----------------
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
[-- Attachment #2: device-modifiers-for-device-clause-v3.diff --]
[-- Type: text/plain, Size: 33573 bytes --]
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.
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..efbf759 100644
--- a/gcc/c/c-parser.c
+++ b/gcc/c/c-parser.c
@@ -15864,37 +15864,87 @@ 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)
+ {
+ error_at (tok->location, "%<ancestor%> device modifier not "
+ "preceded by %<requires%> directive "
+ "with %<reverse_offload%> clause");
+ 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;
}
+ if (ancestor && TREE_CODE (t) == INTEGER_CST && !integer_onep (t))
+ {
+ error_at (expr_loc, "the %<device%> clause expression must evaluate to "
+ "%<1%>");
+ 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/cp/parser.c b/gcc/cp/parser.c
index 93698aa..2c1e202 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, "%<ancestor%> device modifier not "
+ "preceded by %<requires%> directive "
+ "with %<reverse_offload%> clause");
+ 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..7352f90 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
+ && !integer_onep (t))
+ {
+ 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..16fab73 100644
--- a/gcc/fortran/openmp.c
+++ b/gcc/fortran/openmp.c
@@ -1714,8 +1714,53 @@ 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 ("%<ancestor%> device modifier not "
+ "preceded by %<requires%> directive "
+ "with %<reverse_offload%> clause 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)
+ {
+ 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..653d4cf 100644
--- a/gcc/gimplify.c
+++ b/gcc/gimplify.c
@@ -10088,6 +10088,38 @@ 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;
+ break;
+ }
+
+ 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;
+ break;
+ }
+ }
+ /* 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..65252d6 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,17 @@ 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 tgt_clauses = gimple_omp_target_clauses (ctx->stmt);
+ tree c = omp_find_clause (tgt_clauses, 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..9822862
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/gomp/target-device-1.c
@@ -0,0 +1,32 @@
+/* { dg-do compile } */
+
+void
+foo (int n)
+{
+ /* Test to ensure that 'device_num' is parsed correctly in device clauses. */
+
+ #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..b3c1ce8
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/gomp/target-device-ancestor-1.c
@@ -0,0 +1,13 @@
+/* { 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 "'ancestor' device modifier not preceded by 'requires' directive with 'reverse_offload' clause" } */
+ /* { dg-error "expected '\\)' before 'ancestor'" "" { target c } .-1 } */
+
+ ;
+}
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..cf05c50
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/gomp/target-device-ancestor-2.c
@@ -0,0 +1,82 @@
+/* { dg-do compile } */
+
+#pragma omp requires reverse_offload /* { dg-message "sorry, unimplemented: 'reverse_offload' clause on 'requires' directive not supported yet" } */
+
+void
+foo (int n)
+{
+ /* 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'" } */
+ ;
+
+ #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 = 0, b = 0, c = 0;
+ #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 *-*-* } } */
+
+
+ /* 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..9a170db
--- /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 "'ancestor' device modifier not preceded by 'requires' directive with 'reverse_offload' clause" }
+! !$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) \
^ permalink raw reply [flat|nested] 8+ messages in thread