Hi Tobias!
On 2026-07-14T22:38:40+0200, Tobias Burnus <[email protected]> wrote:
> 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?
..., or after. My question is: how did you test this? I'm seeing:
@@ -7893,7 +7893,7 @@ PASS: c-c++-common/gomp/reverse-offload-1.c (test for
excess errors)
PASS: c-c++-common/gomp/reverse-offload-1.c scan-tree-dump-times omplower
"__attribute__\\(\\(omp declare target\\)\\)[\n\r]*int called_in_target1" 1
PASS: c-c++-common/gomp/reverse-offload-1.c scan-tree-dump-times omplower
"__attribute__\\(\\(omp declare target\\)\\)[\n\r]*int called_in_target2" 1
PASS: c-c++-common/gomp/reverse-offload-1.c scan-tree-dump-times omplower
"__attribute__\\(\\(omp declare target\\)\\)[\n\r]*void tg_fn" 1
[-PASS:-]{+FAIL:+} c-c++-common/gomp/reverse-offload-1.c
scan-tree-dump-times omplower "omp declare target[^ ]" 3
Same for C++.
..., and for AMD as well as NVIDIA GPU offloading configurations:
UNSUPPORTED: libgomp.c/../libgomp.c-c++-common/reverse-offload-1-aux.c
[-PASS:-]{+FAIL:+} libgomp.c/../libgomp.c-c++-common/reverse-offload-1.c
(test for excess errors)
[-PASS:-]{+UNRESOLVED:+}
libgomp.c/../libgomp.c-c++-common/reverse-offload-1.c [-execution
test-]{+compilation failed to produce executable+}
[...]/source-gcc/libgomp/testsuite/libgomp.c/../libgomp.c-c++-common/reverse-offload-1.c:19:12:
error: variable 'global_var' has been referenced in offloaded code but hasn't
been marked to be included in the offloaded code
lto1: fatal error: errors during merging of translation units
compilation terminated.
gcn mkoffload: fatal error:
[...]/build-gcc/gcc/x86_64-pc-linux-gnu-accel-amdgcn-amdhsa-gcc returned 1 exit
status
@@ -2549,8 +2549,8 @@ PASS: libgomp.c/../libgomp.c-c++-common/target-42.c
(test for warnings, line 10
PASS: libgomp.c/../libgomp.c-c++-common/target-42.c (test for warnings,
line 12)
PASS: libgomp.c/../libgomp.c-c++-common/target-42.c (test for warnings,
line 14)
PASS: libgomp.c/../libgomp.c-c++-common/target-42.c (test for warnings,
line 15)
[-PASS:-]{+FAIL:+} libgomp.c/../libgomp.c-c++-common/target-42.c (test for
excess errors)
[-PASS:-]{+UNRESOLVED:+} libgomp.c/../libgomp.c-c++-common/target-42.c
[-execution test-]{+compilation failed to produce executable+}
[...]
[...]/source-gcc/libgomp/testsuite/libgomp.c/../libgomp.c-c++-common/target-42.c:4:1:
error: function 'foo' has been referenced in offloaded code but hasn't been
marked to be included in the offloaded code
lto1: fatal error: errors during merging of translation units
compilation terminated.
gcn mkoffload: fatal error:
[...]/build-gcc/gcc/x86_64-pc-linux-gnu-accel-amdgcn-amdhsa-gcc returned 1 exit
status
[-PASS:-]{+FAIL:+} libgomp.fortran/reverse-offload-1.f90 -O0 (test for
excess errors)
[-PASS:-]{+UNRESOLVED:+} libgomp.fortran/reverse-offload-1.f90 -O0
[-execution test-]{+compilation failed to produce executable+}
[Etc.]
[...]/source-gcc/libgomp/testsuite/libgomp.fortran/reverse-offload-1.f90:53:14:
error: variable 'prog_var' has been referenced in offloaded code but hasn't
been marked to be included in the offloaded code
[...]/source-gcc/libgomp/testsuite/libgomp.fortran/reverse-offload-1.f90:22:14:
error: variable 'global_var' has been referenced in offloaded code but hasn't
been marked to be included in the offloaded code
lto1: fatal error: errors during merging of translation units
compilation terminated.
gcn mkoffload: fatal error:
[...]/build-gcc/gcc/x86_64-pc-linux-gnu-accel-amdgcn-amdhsa-gcc returned 1 exit
status
[-PASS:-]{+FAIL:+} libgomp.fortran/reverse-offload-6.f90 -O (test for
excess errors)
[-PASS:-]{+UNRESOLVED:+} libgomp.fortran/reverse-offload-6.f90 -O
[-execution test-]{+compilation failed to produce executable+}
[...]/source-gcc/libgomp/testsuite/libgomp.fortran/reverse-offload-6.f90:9:14:
error: variable 'x' has been referenced in offloaded code but hasn't been
marked to be included in the offloaded code
[...]/source-gcc/libgomp/testsuite/libgomp.fortran/reverse-offload-6.f90:10:14:
error: variable 'dev_num2' has been referenced in offloaded code but hasn't
been marked to be included in the offloaded code
lto1: fatal error: errors during merging of translation units
compilation terminated.
gcn mkoffload: fatal error:
[...]/build-gcc/gcc/x86_64-pc-linux-gnu-accel-amdgcn-amdhsa-gcc returned 1 exit
status
Or, is something wrong on my end?
Grüße
Thomas
> 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();
> +}
