https://gcc.gnu.org/g:97a43fc4dbc5cfb1174a13313dda7637c838d110
commit r17-2401-g97a43fc4dbc5cfb1174a13313dda7637c838d110 Author: Tobias Burnus <[email protected]> Date: Tue Jul 14 21:44:48 2026 +0200 OpenMP: Handle 'device_type(nohost)' on 'target' The parsing support was already there, this just adds code to handle 'nohost' in the middle end. That's handled by adding if-else code to the target's structure block using the omp_is_initial_device builtin in the condition that is replaced by a constant (late) during compilation. Thus, with optimization turned on, only the error message or the target body remains for the host and the nonhost region, respectively. Note that 'device_type(host)' is not yet implemented. Note also there is no change to implicit handling, i.e. functions called in the region will implicitly 'declare target' with 'device_type(any)', i.e. there is no attempt made to add those with 'nohost', which is also the most sensible. gcc/ChangeLog: * gimplify.cc (gimplify_omp_workshare): Handle 'device_type(nohost)' on 'target'. * omp-low.cc (lower_omp_target): Remove sorry for this case. libgomp/ChangeLog: * error.c (GOMP_error): When passing msglen == (size_t) -2, do not show 'error directive encountered:' in the message. * libgomp.texi (Implementation Status): Mark target + device_type as partially implemented. * testsuite/libgomp.c/target-device-type-1.c: New test. * testsuite/libgomp.c/target-device-type-2.c: New test. gcc/testsuite/ChangeLog: * c-c++-common/gomp/target-device-type-1.c: Update dg-sorry as now only device_type(host) remains unimplemented. * gfortran.dg/gomp/target-device-type-1.f90: Likewise. Diff: --- gcc/gimplify.cc | 40 ++++++++++++++++++++++ gcc/omp-low.cc | 4 +-- .../c-c++-common/gomp/target-device-type-1.c | 4 +-- .../gfortran.dg/gomp/target-device-type-1.f90 | 4 +-- libgomp/error.c | 6 ++++ libgomp/libgomp.texi | 2 +- libgomp/testsuite/libgomp.c/target-device-type-1.c | 21 ++++++++++++ libgomp/testsuite/libgomp.c/target-device-type-2.c | 27 +++++++++++++++ 8 files changed, 101 insertions(+), 7 deletions(-) diff --git a/gcc/gimplify.cc b/gcc/gimplify.cc index 6c5d182ffab2..223b55b7e93e 100644 --- a/gcc/gimplify.cc +++ b/gcc/gimplify.cc @@ -18990,6 +18990,46 @@ gimplify_omp_workshare (tree *expr_p, gimple_seq *pre_p) stmt = gimple_build_omp_scope (body, OMP_CLAUSES (expr)); break; case OMP_TARGET: + /* Handle 'device_type(nohost)'. */ + if (tree c = omp_find_clause (OMP_CLAUSES (expr), OMP_CLAUSE_DEVICE_TYPE)) + if (OMP_CLAUSE_DEVICE_TYPE_KIND (c) == OMP_CLAUSE_DEVICE_TYPE_NOHOST) + { + location_t loc = gimple_location (body); + tree fn = builtin_decl_explicit (BUILT_IN_OMP_IS_INITIAL_DEVICE); + fn = build_call_expr_loc (loc, fn, 0); + tree err = builtin_decl_explicit (BUILT_IN_GOMP_ERROR); + const char *str = "Executing device-type 'nohost' target region " + "on the host"; + const int len = strlen (str) + 1; + tree msg = build_string (len, str); + TREE_TYPE (msg) = build_array_type_nelts (char_type_node, len); + TREE_READONLY (msg) = 1; + TREE_STATIC (msg) = 1; + msg = build_fold_addr_expr (msg); + tree msglen = fold_build2 (MINUS_EXPR, size_type_node, + build_all_ones_cst (size_type_node), + size_one_node); + err = build_call_expr_loc (loc, err, 2, msg, msglen); + tree l1 = create_artificial_label (UNKNOWN_LOCATION); + tree l2 = create_artificial_label (UNKNOWN_LOCATION); + tree l3 = create_artificial_label (UNKNOWN_LOCATION); + + gimple_seq new_body = NULL; + + gimplify_expr (&fn, &new_body, NULL, is_gimple_val, fb_rvalue); + gcond *cond = gimple_build_cond (NE_EXPR, fn, + build_int_cst (TREE_TYPE (fn), 0), + l1, l2); + gimplify_seq_add_stmt (&new_body, cond); + gimplify_seq_add_stmt (&new_body, gimple_build_label (l1)); + gimplify_and_add (err, &new_body); + gimplify_seq_add_stmt (&new_body, gimple_build_goto (l3)); + gimplify_seq_add_stmt (&new_body, gimple_build_label (l2)); + gimple_seq_add_seq (&new_body, body); + gimplify_seq_add_stmt (&new_body, gimple_build_label (l3)); + body = gimple_build_bind (NULL_TREE, new_body, make_node (BLOCK)); + gimple_set_location (body, loc); + } stmt = gimple_build_omp_target (body, GF_OMP_TARGET_KIND_REGION, OMP_CLAUSES (expr), iterator_loops_seq); break; diff --git a/gcc/omp-low.cc b/gcc/omp-low.cc index c578e8ece49a..99f73e501c02 100644 --- a/gcc/omp-low.cc +++ b/gcc/omp-low.cc @@ -13287,9 +13287,9 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx) 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) + if (OMP_CLAUSE_DEVICE_TYPE_KIND (c) == OMP_CLAUSE_DEVICE_TYPE_HOST) sorry_at (OMP_CLAUSE_LOCATION (c), - "only the %<device_type(any)%> is supported"); + "the %<device_type(host)%> is not supported"); break; case OMP_CLAUSE_USES_ALLOCATORS: allocator = OMP_CLAUSE_USES_ALLOCATORS_ALLOCATOR (c); 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 index e64349baae3e..758330436cd1 100644 --- a/gcc/testsuite/c-c++-common/gomp/target-device-type-1.c +++ b/gcc/testsuite/c-c++-common/gomp/target-device-type-1.c @@ -10,10 +10,10 @@ void f () #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 ( nohost ) ; -#pragma omp target device_type ( host ) +#pragma omp target device_type ( host ) // { dg-message "sorry, unimplemented: the 'device_type\\(host\\)' is not supported" } ; } diff --git a/gcc/testsuite/gfortran.dg/gomp/target-device-type-1.f90 b/gcc/testsuite/gfortran.dg/gomp/target-device-type-1.f90 index be33bb6a19d7..164fc4f1fd60 100644 --- a/gcc/testsuite/gfortran.dg/gomp/target-device-type-1.f90 +++ b/gcc/testsuite/gfortran.dg/gomp/target-device-type-1.f90 @@ -7,10 +7,10 @@ !$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 target device_type ( nohost ) !$omp end target -!$omp target device_type ( host ) +!$omp target device_type ( host ) ! { dg-message "sorry, unimplemented: the 'device_type\\(host\\)' is not supported" } !$omp end target end diff --git a/libgomp/error.c b/libgomp/error.c index cd7c98cf25ec..0fa311faeca1 100644 --- a/libgomp/error.c +++ b/libgomp/error.c @@ -105,11 +105,17 @@ GOMP_warning (const char *msg, size_t msglen) gomp_error ("error directive encountered"); } +/* Handle 'omp error' directive; for msglen == -2, it is used internally + by the compiler-generated code such that 'error directive encountered' + is not printed. */ + void GOMP_error (const char *msg, size_t msglen) { if (msg && msglen == (size_t) -1) gomp_fatal ("fatal error: error directive encountered: %s", msg); + else if (msg && msglen == (size_t) -2) + gomp_fatal ("fatal error: %s", msg); else if (msg) { fputs ("\nlibgomp: fatal error: error directive encountered: ", stderr); diff --git a/libgomp/libgomp.texi b/libgomp/libgomp.texi index 5bbc856544ec..63c5cf475b7b 100644 --- a/libgomp/libgomp.texi +++ b/libgomp/libgomp.texi @@ -580,7 +580,7 @@ to address of matching mapped list item per 5.1, Sect. 2.21.7.2 @tab N @tab @code{target_data}, @code{target_exit_data} and @code{target_update} @tab N @tab @item New @code{device_type} clause to the @code{target} directive - @tab N @tab + @tab P @tab @code{device_type(host)} unsupported @item @code{target_data} as composite construct @tab N @tab @item @code{nowait} clause with reverse-offload @code{target} directives @tab N @tab diff --git a/libgomp/testsuite/libgomp.c/target-device-type-1.c b/libgomp/testsuite/libgomp.c/target-device-type-1.c new file mode 100644 index 000000000000..a1e1af6979ad --- /dev/null +++ b/libgomp/testsuite/libgomp.c/target-device-type-1.c @@ -0,0 +1,21 @@ +/* { dg-do run { target offload_device } } */ +/* { dg-additional-options "-fdump-tree-gimple" } */ + +int +f (int *x) +{ + #pragma omp target device_type(nohost) + (*x)++; + return *x; +} + +int +main () +{ + int i = 3; + if (f(&i) != 4) + __builtin_abort (); +} + +/* { dg-final { scan-tree-dump "= __builtin_omp_is_initial_device \\(\\);" "gimple" } } */ +/* { dg-final { scan-tree-dump "__builtin_GOMP_error \\(\"Executing device-type ..nohost.. target region on the host\"," "gimple" } } */ diff --git a/libgomp/testsuite/libgomp.c/target-device-type-2.c b/libgomp/testsuite/libgomp.c/target-device-type-2.c new file mode 100644 index 000000000000..37fa62216c93 --- /dev/null +++ b/libgomp/testsuite/libgomp.c/target-device-type-2.c @@ -0,0 +1,27 @@ +/* { dg-do run } */ +/* { dg-additional-options "-fdump-tree-gimple" } */ + +#include <omp.h> + +int +f (int *x) +{ + #pragma omp target device_type(nohost) + (*x)++; + return *x; +} + +int +main () +{ + omp_set_default_device (omp_initial_device); + int i = 3; + if (f(&i) != 4) + __builtin_abort (); +} + +/* { dg-final { scan-tree-dump "= __builtin_omp_is_initial_device \\(\\);" "gimple" } } */ +/* { dg-final { scan-tree-dump "__builtin_GOMP_error \\(\"Executing device-type ..nohost.. target region on the host\"," "gimple" } } */ + +/* { dg-output "libgomp: fatal error: Executing device-type 'nohost' target region on the host" } */ +/* { dg-shouldfail "'nohost' target region on the host" } */
