Follow up to the just committed support for 'device_type(nohost)'
on 'target' - adding now 'device_type(host)' and competing this
feature. ('any' is also supported.)
Comment before I commit it?
Tobias
OpenMP: Handle 'device_type(host)' on 'target'
When 'device_type(host)' is specified on a target region, no offload code
is generated for the outlined region. Additionally, it is honored for
implicit 'declare target'.
The code is executed as if 'device(omp_initial_device)' were present
(i.e. also no issus with OMP_TARGET_OFFLOAD=mandatory).
gcc/ChangeLog:
* omp-expand.cc (expand_omp_target): For 'device_type(host)',
invoke with device(omp_initial_device)'.
* omp-low.cc (create_omp_child_function): Take host_only
argument to prevent generating offload code.
scan_omp_parallel, scan_omp_task, scan_omp_target, scan_omp_teams,
lower_omp_target): Update call.
* omp-offload.cc (omp_discover_declare_target_tgt_fn_r,
omp_discover_declare_target_fn_r): Do not add implicit 'declare
target' to target regions with 'device_type(host)'.
libgomp/ChangeLog:
* libgomp.texi (Implementations status): Mark 'device_type'
clause on target as implemented.
* testsuite/libgomp.c/target-device-type-3.c: New test.
gcc/testsuite/ChangeLog:
* c-c++-common/gomp/target-device-type-1.c: Remove dg-sorry.
* gfortran.dg/gomp/target-device-type-1.f90: Likewise.
gcc/omp-expand.cc | 12 +++-
gcc/omp-low.cc | 26 ++++-----
gcc/omp-offload.cc | 10 +++-
.../c-c++-common/gomp/target-device-type-1.c | 2 +-
.../gfortran.dg/gomp/target-device-type-1.f90 | 2 +-
libgomp/libgomp.texi | 4 +-
libgomp/testsuite/libgomp.c/target-device-type-3.c | 65 ++++++++++++++++++++++
7 files changed, 97 insertions(+), 24 deletions(-)
diff --git a/gcc/omp-expand.cc b/gcc/omp-expand.cc
index 388f84eddff..ca4d5d0db9e 100644
--- a/gcc/omp-expand.cc
+++ b/gcc/omp-expand.cc
@@ -10014,6 +10014,7 @@ expand_omp_target (struct omp_region *region)
tree clauses = gimple_omp_target_clauses (entry_stmt);
bool is_ancestor = false;
+ bool is_host_only = false;
child_fn = child_fn2 = NULL_TREE;
child_cfun = NULL;
if (offloaded)
@@ -10021,6 +10022,9 @@ expand_omp_target (struct omp_region *region)
c = omp_find_clause (clauses, OMP_CLAUSE_DEVICE);
if (ENABLE_OFFLOADING && c)
is_ancestor = OMP_CLAUSE_DEVICE_ANCESTOR (c);
+ c = omp_find_clause (clauses, OMP_CLAUSE_DEVICE_TYPE);
+ if (c && OMP_CLAUSE_DEVICE_TYPE_KIND (c) == OMP_CLAUSE_DEVICE_TYPE_HOST)
+ is_host_only = true;
child_fn = gimple_omp_target_child_fn (entry_stmt);
child_cfun = DECL_STRUCT_FUNCTION (child_fn);
}
@@ -10211,7 +10215,7 @@ expand_omp_target (struct omp_region *region)
{
if (in_lto_p)
DECL_PRESERVE_P (child_fn) = 1;
- if (!is_ancestor)
+ if (!is_ancestor && !is_host_only)
vec_safe_push (offload_funcs, child_fn);
}
@@ -10393,8 +10397,10 @@ expand_omp_target (struct omp_region *region)
}
else
{
- c = omp_find_clause (clauses, OMP_CLAUSE_DEVICE);
- if (c)
+ if (is_host_only)
+ device = build_int_cst (integer_type_node,
+ GOMP_DEVICE_HOST_FALLBACK - 1);
+ else if ((c = omp_find_clause (clauses, OMP_CLAUSE_DEVICE)) != NULL_TREE)
{
device = OMP_CLAUSE_DEVICE_ID (c);
/* Ensure 'device' is of the correct type. */
diff --git a/gcc/omp-low.cc b/gcc/omp-low.cc
index 99f73e501c0..ec421657c68 100644
--- a/gcc/omp-low.cc
+++ b/gcc/omp-low.cc
@@ -2131,10 +2131,10 @@ omp_maybe_offloaded_ctx (omp_context *ctx)
}
/* Build a decl for the omp child function. It'll not contain a body
- yet, just the bare decl. */
+ yet, just the bare decl. If HOST_ONLY, do not create a device version. */
static void
-create_omp_child_function (omp_context *ctx, bool task_copy)
+create_omp_child_function (omp_context *ctx, bool task_copy, bool host_only)
{
tree decl, type, name, t;
@@ -2189,7 +2189,7 @@ create_omp_child_function (omp_context *ctx, bool task_copy)
DECL_FUNCTION_VERSIONED (decl)
= DECL_FUNCTION_VERSIONED (current_function_decl);
- if (omp_maybe_offloaded_ctx (ctx))
+ if (omp_maybe_offloaded_ctx (ctx) && !host_only)
{
cgraph_node::get_create (decl)->offloadable = 1;
if (ENABLE_OFFLOADING)
@@ -2435,7 +2435,7 @@ scan_omp_parallel (gimple_stmt_iterator *gsi, omp_context *outer_ctx)
DECL_NAMELESS (name) = 1;
TYPE_NAME (ctx->record_type) = name;
TYPE_ARTIFICIAL (ctx->record_type) = 1;
- create_omp_child_function (ctx, false);
+ create_omp_child_function (ctx, false, false);
gimple_omp_parallel_set_child_fn (stmt, ctx->cb.dst_fn);
scan_sharing_clauses (gimple_omp_parallel_clauses (stmt), ctx);
@@ -2488,7 +2488,7 @@ scan_omp_task (gimple_stmt_iterator *gsi, omp_context *outer_ctx)
DECL_NAMELESS (name) = 1;
TYPE_NAME (ctx->record_type) = name;
TYPE_ARTIFICIAL (ctx->record_type) = 1;
- create_omp_child_function (ctx, false);
+ create_omp_child_function (ctx, false, false);
gimple_omp_task_set_child_fn (stmt, ctx->cb.dst_fn);
scan_sharing_clauses (gimple_omp_task_clauses (stmt), ctx);
@@ -2502,7 +2502,7 @@ scan_omp_task (gimple_stmt_iterator *gsi, omp_context *outer_ctx)
DECL_NAMELESS (name) = 1;
TYPE_NAME (ctx->srecord_type) = name;
TYPE_ARTIFICIAL (ctx->srecord_type) = 1;
- create_omp_child_function (ctx, true);
+ create_omp_child_function (ctx, true, false);
}
scan_omp (gimple_omp_body_ptr (stmt), ctx);
@@ -3209,7 +3209,10 @@ scan_omp_target (gomp_target *stmt, omp_context *outer_ctx)
if (offloaded)
{
- create_omp_child_function (ctx, false);
+ tree c = omp_find_clause (clauses, OMP_CLAUSE_DEVICE_TYPE);
+ bool host_only
+ = c && OMP_CLAUSE_DEVICE_TYPE_KIND (c) == OMP_CLAUSE_DEVICE_TYPE_HOST;
+ create_omp_child_function (ctx, false, host_only);
gimple_omp_target_set_child_fn (stmt, ctx->cb.dst_fn);
}
@@ -3268,7 +3271,7 @@ scan_omp_teams (gomp_teams *stmt, omp_context *outer_ctx)
DECL_NAMELESS (name) = 1;
TYPE_NAME (ctx->record_type) = name;
TYPE_ARTIFICIAL (ctx->record_type) = 1;
- create_omp_child_function (ctx, false);
+ create_omp_child_function (ctx, false, false);
gimple_omp_teams_set_child_fn (stmt, ctx->cb.dst_fn);
scan_sharing_clauses (gimple_omp_teams_clauses (stmt), ctx);
@@ -13284,13 +13287,6 @@ 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_HOST)
- sorry_at (OMP_CLAUSE_LOCATION (c),
- "the %<device_type(host)%> is not supported");
- break;
case OMP_CLAUSE_USES_ALLOCATORS:
allocator = OMP_CLAUSE_USES_ALLOCATORS_ALLOCATOR (c);
tree new_allocator = lookup_decl (allocator, ctx);
diff --git a/gcc/omp-offload.cc b/gcc/omp-offload.cc
index 7cd2a572b7c..bbe13154d8f 100644
--- a/gcc/omp-offload.cc
+++ b/gcc/omp-offload.cc
@@ -272,7 +272,10 @@ omp_discover_declare_target_tgt_fn_r (tree *tp, int *walk_subtrees, void *data)
else if (TREE_CODE (*tp) == OMP_TARGET)
{
tree c = omp_find_clause (OMP_CLAUSES (*tp), OMP_CLAUSE_DEVICE);
- if (c && OMP_CLAUSE_DEVICE_ANCESTOR (c))
+ tree c2 = omp_find_clause (OMP_CLAUSES (*tp), OMP_CLAUSE_DEVICE_TYPE);
+ if ((!c || !OMP_CLAUSE_DEVICE_ANCESTOR (c))
+ && (!c2 || (OMP_CLAUSE_DEVICE_TYPE_KIND (c2)
+ != OMP_CLAUSE_DEVICE_TYPE_HOST)))
*walk_subtrees = 0;
}
return NULL_TREE;
@@ -286,7 +289,10 @@ omp_discover_declare_target_fn_r (tree *tp, int *walk_subtrees, void *data)
if (TREE_CODE (*tp) == OMP_TARGET)
{
tree c = omp_find_clause (OMP_CLAUSES (*tp), OMP_CLAUSE_DEVICE);
- if (!c || !OMP_CLAUSE_DEVICE_ANCESTOR (c))
+ tree c2 = omp_find_clause (OMP_CLAUSES (*tp), OMP_CLAUSE_DEVICE_TYPE);
+ if ((!c || !OMP_CLAUSE_DEVICE_ANCESTOR (c))
+ && (!c2 || (OMP_CLAUSE_DEVICE_TYPE_KIND (c2)
+ != OMP_CLAUSE_DEVICE_TYPE_HOST)))
walk_tree_without_duplicates (&OMP_TARGET_BODY (*tp),
omp_discover_declare_target_tgt_fn_r,
data);
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 758330436cd..15817717daa 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
@@ -13,7 +13,7 @@ void f ()
#pragma omp target device_type ( nohost )
;
-#pragma omp target device_type ( host ) // { dg-message "sorry, unimplemented: the 'device_type\\(host\\)' is not supported" }
+#pragma omp target device_type ( host )
;
}
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 164fc4f1fd6..fa1265f5ed2 100644
--- a/gcc/testsuite/gfortran.dg/gomp/target-device-type-1.f90
+++ b/gcc/testsuite/gfortran.dg/gomp/target-device-type-1.f90
@@ -10,7 +10,7 @@
!$omp target device_type ( nohost )
!$omp end target
-!$omp target device_type ( host ) ! { dg-message "sorry, unimplemented: the 'device_type\\(host\\)' is not supported" }
+!$omp target device_type ( host )
!$omp end target
end
diff --git a/libgomp/libgomp.texi b/libgomp/libgomp.texi
index 63c5cf475b7..4e2bdc280e9 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 P @tab @code{device_type(host)} unsupported
+ @tab Y @tab
@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-3.c b/libgomp/testsuite/libgomp.c/target-device-type-3.c
new file mode 100644
index 00000000000..faedcabdf22
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c/target-device-type-3.c
@@ -0,0 +1,65 @@
+/* { dg-do run } */
+/* { dg-additional-options -fdump-tree-ompexp } */
+/* { dg-additional-options -foffload-options=-fdump-tree-optimized-raw-asmname } */
+
+/* This code is supposed to run on the host independent of the current
+ default-device-var ICV - and it shall not create any device code for
+ neither the target region nor the called functions. */
+
+// NOTE: With offloading,
+// UNRESOLVED:
+// is expected for the "optimized" tests as those dumps are only enabled
+// for the no-host side.
+
+
+// Due to 'device_type(nohost)', the outlined region shall not be marked
+// as 'target entrypoint' - nor should 'ggg' become implicitly
+// 'declare target' -- check for the absence of those __attributes__:
+
+/* { dg-final { scan-tree-dump-times "omp declare target" 1 "ompexp" } } */
+/* { dg-final { scan-tree-dump-not "omp target entrypoint" "ompexp" } } */
+
+
+// Expect that GOMP_target explicitly invokes the host device, i.e.
+// -3 == (shifted) omp_initial_device (aka host fallback)
+
+/* { dg-final { scan-tree-dump "__builtin_GOMP_target_ext \\(-3," "ompexp" } } */
+
+
+// On the device side, expect only hhh:
+
+/* { dg-final { scan-tree-dump "hhh" "optimized" { target offload_device } } } */
+/* { dg-final { scan-tree-dump-not "ggg" "optimized" { target offload_device } } } */
+/* { dg-final { scan-tree-dump-not "_omp_fn" "optimized" { target offload_device } } } */
+
+
+#include <omp.h>
+
+// Use 'declare target' such that -fdump-tree-optimized creates a file.
+// Otherwise, there is dump file (only a stub assembly file)
+// That's the only purpose of this function.
+int hhh() { return 5; }
+#pragma omp declare target enter(hhh)
+
+
+int ggg() {
+ return 42;
+}
+
+int f(int *x) {
+ bool initial_dev = false;
+ #pragma omp target map(tofrom: initial_dev) device_type(host)
+ {
+ initial_dev = omp_is_initial_device ();
+ (*x) = (*x) + 1 + ggg();
+ }
+ if (!initial_dev)
+ __builtin_abort ();
+ return *x;
+}
+
+int main() {
+ int i = 5;
+ if (f(&i) != 5 + 1 + 42)
+ __builtin_abort();
+}
