Another commit in the series to fix issues for 'declare target', including issues with requires self_maps, and improve/fix issues related to the 'device_type' clause.
This one adds 'device_type' clause to the 'target' construct, albeit errors out for anything but 'any'. Committed as r16-6245-g1ed9526113c970. * * * Parser wise, it is the last Fortran patch in this series (except for some trans*.cc code gen changes + removing 'sorry'), for C/C++, the 'local' clause parsing to 'declare target' is still missing. In terms of the middle end: - 'device_type(host)' shouldn't enable offloading which also affects implicit declare target. - For 'self_maps' declare target's 'enter' (aka 'to') clause becomes 'link', except for 'local' - which also needs to be implemented. (incl. implicit declare target) - Finally, the nohost handling needs to be improved/implemented. (It exists in some way for OpenACC and for reverse offload, but needs to be improved; also declare-target variables are unimplemented.) At least the second bullet is a correctness issue and not a quality of implementation one. Tobias
commit 1ed9526113c970126ca710b516e5260a25f1def9 Author: Tobias Burnus <[email protected]> Date: Thu Dec 18 12:20:36 2025 +0100 OpenMP: Add parser support for target's device_type clause Enables the existing 'device_type(nohost|host|any)' for the 'target' construct; for now, it will fail with a 'sorry, unimplemented' for all but 'any'. gcc/c-family/ChangeLog: * c-omp.cc (c_omp_split_clauses): Handle target's device_type clause. gcc/c/ChangeLog: * c-parser.cc (OMP_TARGET_CLAUSE_MASK): Add device_type clause. gcc/cp/ChangeLog: * parser.cc (OMP_TARGET_CLAUSE_MASK): Add device_type clause. gcc/fortran/ChangeLog: * dump-parse-tree.cc (show_omp_clauses): Handle device_type clause. * openmp.cc (gfc_match_omp_clauses): Reorder to match 'device' after 'device_...' to avoid parse errors. (OMP_TARGET_CLAUSES): Add device_type clause. * trans-openmp.cc (gfc_trans_omp_clauses, gfc_split_omp_clauses): Handle device_type clause. gcc/ChangeLog: * gimplify.cc (gimplify_scan_omp_clauses): Handle OpenMP device_type clause. * omp-low.cc (scan_sharing_clauses): Likewise. (lower_omp_target): Print 'sorry, unimplemented' for device_type clause value other than 'any'. gcc/testsuite/ChangeLog: * c-c++-common/gomp/target-device-type-1.c: New test. * gfortran.dg/gomp/target-device-type-1.f90: New test. --- gcc/c-family/c-omp.cc | 1 + gcc/c/c-parser.cc | 1 + gcc/cp/parser.cc | 1 + gcc/fortran/dump-parse-tree.cc | 14 ++++ gcc/fortran/openmp.cc | 75 +++++++++++----------- gcc/fortran/trans-openmp.cc | 24 +++++++ gcc/gimplify.cc | 2 + gcc/omp-low.cc | 9 +++ .../c-c++-common/gomp/target-device-type-1.c | 24 +++++++ .../gfortran.dg/gomp/target-device-type-1.f90 | 21 ++++++ 10 files changed, 135 insertions(+), 37 deletions(-) diff --git a/gcc/c-family/c-omp.cc b/gcc/c-family/c-omp.cc index 9dd0d1450dd..69b01253c0c 100644 --- a/gcc/c-family/c-omp.cc +++ b/gcc/c-family/c-omp.cc @@ -2176,6 +2176,7 @@ c_omp_split_clauses (location_t loc, enum tree_code code, { /* First the clauses that are unique to some constructs. */ case OMP_CLAUSE_DEVICE: + case OMP_CLAUSE_DEVICE_TYPE: case OMP_CLAUSE_DEFAULTMAP: case OMP_CLAUSE_DEPEND: case OMP_CLAUSE_DYN_GROUPPRIVATE: diff --git a/gcc/c/c-parser.cc b/gcc/c/c-parser.cc index 6e3c7bea3bc..d6d0b0ed415 100644 --- a/gcc/c/c-parser.cc +++ b/gcc/c/c-parser.cc @@ -27217,6 +27217,7 @@ c_parser_omp_target_exit_data (location_t loc, c_parser *parser, | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_DEFAULTMAP) \ | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_DEPEND) \ | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_DEVICE) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_DEVICE_TYPE) \ | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_DYN_GROUPPRIVATE) \ | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_FIRSTPRIVATE) \ | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_HAS_DEVICE_ADDR) \ diff --git a/gcc/cp/parser.cc b/gcc/cp/parser.cc index d9c8cfada93..18df3b1014e 100644 --- a/gcc/cp/parser.cc +++ b/gcc/cp/parser.cc @@ -50663,6 +50663,7 @@ cp_parser_omp_target_update (cp_parser *parser, cp_token *pragma_tok, | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_DEFAULTMAP) \ | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_DEPEND) \ | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_DEVICE) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_DEVICE_TYPE) \ | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_DYN_GROUPPRIVATE) \ | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_FIRSTPRIVATE) \ | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_HAS_DEVICE_ADDR) \ diff --git a/gcc/fortran/dump-parse-tree.cc b/gcc/fortran/dump-parse-tree.cc index 2a4ebb0fa0f..db6d54f5fc7 100644 --- a/gcc/fortran/dump-parse-tree.cc +++ b/gcc/fortran/dump-parse-tree.cc @@ -1984,6 +1984,20 @@ show_omp_clauses (gfc_omp_clauses *omp_clauses) fputs (" NOWAIT", dumpfile); if (omp_clauses->collapse) fprintf (dumpfile, " COLLAPSE(%d)", omp_clauses->collapse); + if (omp_clauses->device_type != OMP_DEVICE_TYPE_UNSET) + { + const char *s; + switch (omp_clauses->device_type) + { + case OMP_DEVICE_TYPE_HOST: s = "host"; break; + case OMP_DEVICE_TYPE_NOHOST: s = "nohost"; break; + case OMP_DEVICE_TYPE_ANY: s = "any"; break; + case OMP_DEVICE_TYPE_UNSET: gcc_unreachable (); + } + fputs (" DEVICE_TYPE(", dumpfile); + fputs (s, dumpfile); + fputc (')', dumpfile); + } for (list_type = 0; list_type < OMP_LIST_NUM; list_type++) if (omp_clauses->lists[list_type] != NULL) { diff --git a/gcc/fortran/openmp.cc b/gcc/fortran/openmp.cc index af89c87b0ab..4527068f974 100644 --- a/gcc/fortran/openmp.cc +++ b/gcc/fortran/openmp.cc @@ -3013,6 +3013,43 @@ gfc_match_omp_clauses (gfc_omp_clauses **cp, const omp_mask mask, OMP_MAP_DETACH, false, allow_derived)) continue; + if ((mask & OMP_CLAUSE_DEVICEPTR) + && gfc_match ("deviceptr ( ") == MATCH_YES + && gfc_match_omp_map_clause (&c->lists[OMP_LIST_MAP], + OMP_MAP_FORCE_DEVICEPTR, false, + allow_derived)) + continue; + if ((mask & OMP_CLAUSE_DEVICE_TYPE) + && gfc_match_dupl_check (c->device_type == OMP_DEVICE_TYPE_UNSET, + "device_type", true) == MATCH_YES) + { + if (gfc_match ("host") == MATCH_YES) + c->device_type = OMP_DEVICE_TYPE_HOST; + else if (gfc_match ("nohost") == MATCH_YES) + c->device_type = OMP_DEVICE_TYPE_NOHOST; + else if (gfc_match ("any") == MATCH_YES) + c->device_type = OMP_DEVICE_TYPE_ANY; + else + { + gfc_error ("Expected HOST, NOHOST or ANY at %C"); + break; + } + if (gfc_match (" )") != MATCH_YES) + break; + continue; + } + if ((mask & OMP_CLAUSE_DEVICE_RESIDENT) + && gfc_match_omp_variable_list + ("device_resident (", + &c->lists[OMP_LIST_DEVICE_RESIDENT], true) == MATCH_YES) + continue; + if ((mask & OMP_CLAUSE_DEVICE) + && openacc + && gfc_match ("device ( ") == MATCH_YES + && gfc_match_omp_map_clause (&c->lists[OMP_LIST_MAP], + OMP_MAP_FORCE_TO, true, + /* allow_derived = */ true)) + continue; if ((mask & OMP_CLAUSE_DEVICE) && !openacc && ((m = gfc_match_dupl_check (!c->device, "device", true)) @@ -3072,42 +3109,6 @@ gfc_match_omp_clauses (gfc_omp_clauses **cp, const omp_mask mask, } continue; } - if ((mask & OMP_CLAUSE_DEVICE) - && openacc - && gfc_match ("device ( ") == MATCH_YES - && gfc_match_omp_map_clause (&c->lists[OMP_LIST_MAP], - OMP_MAP_FORCE_TO, true, - /* allow_derived = */ true)) - continue; - if ((mask & OMP_CLAUSE_DEVICEPTR) - && gfc_match ("deviceptr ( ") == MATCH_YES - && gfc_match_omp_map_clause (&c->lists[OMP_LIST_MAP], - OMP_MAP_FORCE_DEVICEPTR, false, - allow_derived)) - continue; - if ((mask & OMP_CLAUSE_DEVICE_TYPE) - && gfc_match ("device_type ( ") == MATCH_YES) - { - if (gfc_match ("host") == MATCH_YES) - c->device_type = OMP_DEVICE_TYPE_HOST; - else if (gfc_match ("nohost") == MATCH_YES) - c->device_type = OMP_DEVICE_TYPE_NOHOST; - else if (gfc_match ("any") == MATCH_YES) - c->device_type = OMP_DEVICE_TYPE_ANY; - else - { - gfc_error ("Expected HOST, NOHOST or ANY at %C"); - break; - } - if (gfc_match (" )") != MATCH_YES) - break; - continue; - } - if ((mask & OMP_CLAUSE_DEVICE_RESIDENT) - && gfc_match_omp_variable_list - ("device_resident (", - &c->lists[OMP_LIST_DEVICE_RESIDENT], true) == MATCH_YES) - continue; if ((mask & OMP_CLAUSE_DIST_SCHEDULE) && c->dist_sched_kind == OMP_SCHED_NONE && gfc_match ("dist_schedule ( static") == MATCH_YES) @@ -5136,7 +5137,7 @@ cleanup: | OMP_CLAUSE_IS_DEVICE_PTR | OMP_CLAUSE_IN_REDUCTION \ | OMP_CLAUSE_THREAD_LIMIT | OMP_CLAUSE_ALLOCATE \ | OMP_CLAUSE_HAS_DEVICE_ADDR | OMP_CLAUSE_USES_ALLOCATORS \ - | OMP_CLAUSE_DYN_GROUPPRIVATE) + | OMP_CLAUSE_DYN_GROUPPRIVATE | OMP_CLAUSE_DEVICE_TYPE) #define OMP_TARGET_DATA_CLAUSES \ (omp_mask (OMP_CLAUSE_DEVICE) | OMP_CLAUSE_MAP | OMP_CLAUSE_IF \ | OMP_CLAUSE_USE_DEVICE_PTR | OMP_CLAUSE_USE_DEVICE_ADDR) diff --git a/gcc/fortran/trans-openmp.cc b/gcc/fortran/trans-openmp.cc index a07dc2ec0e9..254fc934af1 100644 --- a/gcc/fortran/trans-openmp.cc +++ b/gcc/fortran/trans-openmp.cc @@ -5266,6 +5266,28 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses, omp_clauses = gfc_trans_add_clause (c, omp_clauses); } + if (clauses->device_type != OMP_DEVICE_TYPE_UNSET) + { + enum omp_clause_device_type_kind type; + switch (clauses->device_type) + { + case OMP_DEVICE_TYPE_HOST: + type = OMP_CLAUSE_DEVICE_TYPE_HOST; + break; + case OMP_DEVICE_TYPE_NOHOST: + type = OMP_CLAUSE_DEVICE_TYPE_NOHOST; + break; + case OMP_DEVICE_TYPE_ANY: + type = OMP_CLAUSE_DEVICE_TYPE_ANY; + break; + case OMP_DEVICE_TYPE_UNSET: + gcc_unreachable (); + } + c = build_omp_clause (gfc_get_location (&where), OMP_CLAUSE_DEVICE_TYPE); + OMP_CLAUSE_DEVICE_TYPE_KIND (c) = type; + omp_clauses = gfc_trans_add_clause (c, omp_clauses); + } + if (clauses->dyn_groupprivate) { gfc_init_se (&se, NULL); @@ -8051,6 +8073,8 @@ gfc_split_omp_clauses (gfc_code *code, = code->ext.omp_clauses->if_expr; clausesa[GFC_OMP_SPLIT_TARGET].nowait = code->ext.omp_clauses->nowait; + clausesa[GFC_OMP_SPLIT_TARGET].device_type + = code->ext.omp_clauses->device_type; } if (mask & GFC_OMP_MASK_TEAMS) { diff --git a/gcc/gimplify.cc b/gcc/gimplify.cc index ee565336b1d..751b4697271 100644 --- a/gcc/gimplify.cc +++ b/gcc/gimplify.cc @@ -14885,6 +14885,7 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p, case OMP_CLAUSE_INIT: case OMP_CLAUSE_USE: case OMP_CLAUSE_DESTROY: + case OMP_CLAUSE_DEVICE_TYPE: break; case OMP_CLAUSE_DYN_GROUPPRIVATE: @@ -16373,6 +16374,7 @@ end_adjust_omp_map_clause: case OMP_CLAUSE_INCLUSIVE: case OMP_CLAUSE_EXCLUSIVE: case OMP_CLAUSE_USES_ALLOCATORS: + case OMP_CLAUSE_DEVICE_TYPE: break; case OMP_CLAUSE_NOHOST: diff --git a/gcc/omp-low.cc b/gcc/omp-low.cc index 6fd685cdecd..4540a25d1ad 100644 --- a/gcc/omp-low.cc +++ b/gcc/omp-low.cc @@ -1768,6 +1768,7 @@ scan_sharing_clauses (tree clauses, omp_context *ctx) case OMP_CLAUSE_FINALIZE: case OMP_CLAUSE_TASK_REDUCTION: case OMP_CLAUSE_ALLOCATE: + case OMP_CLAUSE_DEVICE_TYPE: break; case OMP_CLAUSE_ALIGNED: @@ -1994,6 +1995,7 @@ scan_sharing_clauses (tree clauses, omp_context *ctx) case OMP_CLAUSE_INIT: case OMP_CLAUSE_USE: case OMP_CLAUSE_DESTROY: + case OMP_CLAUSE_DEVICE_TYPE: break; case OMP_CLAUSE__CACHE_: @@ -13098,6 +13100,13 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx) DECL_HAS_VALUE_EXPR_P (new_var) = 1; } break; + case OMP_CLAUSE_DEVICE_TYPE: + /* FIXME: Ensure that 'nohost' also has not implied before that + 'g->have_offload = true' or an implicit declare target. */ + if (OMP_CLAUSE_DEVICE_TYPE_KIND (c) != OMP_CLAUSE_DEVICE_TYPE_ANY) + sorry_at (OMP_CLAUSE_LOCATION (c), + "only the %<device_type(any)%> is supported"); + break; } if (offloaded) diff --git a/gcc/testsuite/c-c++-common/gomp/target-device-type-1.c b/gcc/testsuite/c-c++-common/gomp/target-device-type-1.c new file mode 100644 index 00000000000..e64349baae3 --- /dev/null +++ b/gcc/testsuite/c-c++-common/gomp/target-device-type-1.c @@ -0,0 +1,24 @@ +// { dg-do compile } +// { dg-additional-options "-fdump-tree-gimple" } + +void f () +{ + +#pragma omp target + ; + +#pragma omp target device_type ( any ) + ; + +#pragma omp target device_type ( nohost ) // { dg-message "sorry, unimplemented: only the 'device_type\\(any\\)' is supported" } + ; + +#pragma omp target device_type ( host ) + ; + +} + +// { dg-final { scan-tree-dump-times "#pragma omp target num_teams\\(-2\\) thread_limit\\(0\\)\[\\r\\n\]" 1 "gimple" } } +// { dg-final { scan-tree-dump-times "#pragma omp target num_teams\\(-2\\) thread_limit\\(0\\) device_type\\(any\\)" 1 "gimple" } } +// { dg-final { scan-tree-dump-times "#pragma omp target num_teams\\(-2\\) thread_limit\\(0\\) device_type\\(nohost\\)" 1 "gimple" } } +// { dg-final { scan-tree-dump-times "#pragma omp target num_teams\\(-2\\) thread_limit\\(0\\) device_type\\(host\\)" 1 "gimple" } } diff --git a/gcc/testsuite/gfortran.dg/gomp/target-device-type-1.f90 b/gcc/testsuite/gfortran.dg/gomp/target-device-type-1.f90 new file mode 100644 index 00000000000..be33bb6a19d --- /dev/null +++ b/gcc/testsuite/gfortran.dg/gomp/target-device-type-1.f90 @@ -0,0 +1,21 @@ +! { dg-do compile } +! { dg-additional-options "-fdump-tree-gimple" } + +!$omp target +!$omp end target + +!$omp target device_type ( any ) +!$omp end target + +!$omp target device_type ( nohost ) ! { dg-message "sorry, unimplemented: only the 'device_type\\(any\\)' is supported" } +!$omp end target + +!$omp target device_type ( host ) +!$omp end target + +end + +! { dg-final { scan-tree-dump-times "#pragma omp target num_teams\\(-2\\) thread_limit\\(0\\)\[\\r\\n\]" 1 "gimple" } } +! { dg-final { scan-tree-dump-times "#pragma omp target num_teams\\(-2\\) thread_limit\\(0\\) device_type\\(any\\)" 1 "gimple" } } +! { dg-final { scan-tree-dump-times "#pragma omp target num_teams\\(-2\\) thread_limit\\(0\\) device_type\\(nohost\\)" 1 "gimple" } } +! { dg-final { scan-tree-dump-times "#pragma omp target num_teams\\(-2\\) thread_limit\\(0\\) device_type\\(host\\)" 1 "gimple" } }
