This is kind of a fallout of the review of the 'requires self_maps'
patch by Chung-Lin as that one also touches 'device_type([no]host)'
in a way (albeit for 'declare target').
[I use the now-implemented feature in an code snipped in
the patch review.]
The attached patch now permits
#pragma omp target device_type(nohost)
f();
which gets turned into
#pragma omp target
if (__builtin_is_initial_device() != 0)
GOMP_error (...);
else
f();
where the builtin eventually turns into a constant such that -Og
(or higher) will remove the respectively unreachable code on the
host and the device side.
Still to be added is support for 'device_type(host)'.
Comments before I push his patch?
Tobias
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.
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 | 13 +++++--
libgomp/testsuite/libgomp.c/target-device-type-1.c | 21 ++++++++++++
libgomp/testsuite/libgomp.c/target-device-type-2.c | 27 +++++++++++++++
8 files changed, 111 insertions(+), 8 deletions(-)
diff --git a/gcc/gimplify.cc b/gcc/gimplify.cc
index 6c5d182ffab..b2b16e3b0a7 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)'. */
+ for (tree c = OMP_CLAUSES (expr); c; c = OMP_CLAUSE_CHAIN (c))
+ if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_DEVICE_TYPE
+ && OMP_CLAUSE_DEVICE_TYPE_KIND (c) == OMP_CLAUSE_DEVICE_TYPE_NOHOST)
+ {
+ tree fn = builtin_decl_explicit (BUILT_IN_OMP_IS_INITIAL_DEVICE);
+ fn = build_call_expr (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 (OMP_CLAUSE_LOCATION (c), 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, NULL_TREE);
+ break;
+ }
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 c578e8ece49..99f73e501c0 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 e64349baae3..758330436cd 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 be33bb6a19d..164fc4f1fd6 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 cd7c98cf25e..0fa311faeca 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 5bbc856544e..a87600eefef 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 00000000000..a1e1af6979a
--- /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 00000000000..37fa62216c9
--- /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" } */
