[gcc r17-2401] OpenMP: Handle 'device_type(nohost)' on 'target'

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" } */