[gcc r16-6245] OpenMP: Add parser support for target's device_type clause
Tobias Burnus
burnus@gcc.gnu.org
Thu Dec 18 11:20:59 GMT 2025
https://gcc.gnu.org/g:1ed9526113c970126ca710b516e5260a25f1def9
commit r16-6245-g1ed9526113c970126ca710b516e5260a25f1def9
Author: Tobias Burnus <tburnus@baylibre.com>
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.
Diff:
---
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 9dd0d1450dd8..69b01253c0cd 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 6e3c7bea3bc8..d6d0b0ed4151 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 d9c8cfada935..18df3b1014e9 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 2a4ebb0fa0f9..db6d54f5fc74 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 af89c87b0ab0..4527068f9745 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 a07dc2ec0e98..254fc934af19 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 ee565336b1d7..751b46972714 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 6fd685cdecd3..4540a25d1ad6 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 000000000000..e64349baae3e
--- /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 000000000000..be33bb6a19d7
--- /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" } }
More information about the Gcc-cvs
mailing list