[gcc r17-2404] OpenMP: Handle 'device_type(host)' on 'target'
Tobias Burnus via Gcc-cvs <[email protected]>
| Newsgroups | gmane.comp.gcc.cvs |
|---|---|
| Message-ID | <[email protected]> |
https://gcc.gnu.org/g:2468c79741f97c2aa0cbb4baaccb86fd4a150910 commit r17-2404-g2468c79741f97c2aa0cbb4baaccb86fd4a150910 Author: Tobias Burnus <[email protected]> Date: Wed Jul 15 08:43:26 2026 +0200 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. Diff: --- 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 | 2 +- libgomp/testsuite/libgomp.c/target-device-type-3.c | 65 ++++++++++++++++++++++ 7 files changed, 96 insertions(+), 23 deletions(-) diff --git a/gcc/omp-expand.cc b/gcc/omp-expand.cc index 388f84eddff0..ca4d5d0db9e5 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 99f73e501c02..ec421657c68b 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 7cd2a572b7c0..bbe13154d8f0 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 758330436cd1..15817717daa9 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 164fc4f1fd60..fa1265f5ed21 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 63c5cf475b7b..d090e2e86d63 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 000000000000..faedcabdf227 --- /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(); +}