https://gcc.gnu.org/g:1f0018091064c26334f4dd67b3472f93c62b0b80
commit r17-3917-g1f0018091064c26334f4dd67b3472f93c62b0b80 Author: Tobias Burnus <[email protected]> Date: Fri Sep 4 09:20:40 2026 +0200 libgomp: GOMP_RUNTIME_USM - fix pointer attachment + update .texi [PR127186] If GOMP_RUNTIME_USM enables self mapping: Handle GOMP_MAP_ATTACH pointer attachment to global static (pointer) variables in map claues on 'target enter/exit data'. Update the documentation to GOMP_RUNTIME_USM to explicitly state what is not supported, removing the reference to 'requires self_maps/USM'. Because: The two features have diverged by now: * For GOMP_RUNTIME_USM, only 'target'/'target data' with 'always' modifier is mishandled (being tracked in the PR); there is an error abort in that case. * For 'requires self_maps/unified_shard_memory', currently more goes wrong but there is a pending compiler patch that turns all 'enter' into 'link' clauses when the requirement is specified, which should solve the issue completely without requiring the changes to libgomp. libgomp/ChangeLog: PR libgomp/127186 * libgomp.texi (GOMP_RUNTIME_USM): Update caveat section. * target.c (gomp_target_enter_exit_data_usm): Handle GOMP_MAP_ATTACH. * testsuite/libgomp.c/usm_env_handling-3.c: New test. Diff: --- libgomp/libgomp.texi | 19 +++-- libgomp/target.c | 32 +++++-- libgomp/testsuite/libgomp.c/usm_env_handling-3.c | 103 +++++++++++++++++++++++ 3 files changed, 141 insertions(+), 13 deletions(-) diff --git a/libgomp/libgomp.texi b/libgomp/libgomp.texi index 82f0a65bd503..92d82f2f66c4 100644 --- a/libgomp/libgomp.texi +++ b/libgomp/libgomp.texi @@ -5246,12 +5246,12 @@ This is currently not specified in more detail, and subject to change. @cindex Environment Variable @table @asis @item @emph{Description}: -Control the behavior of devices which can access host moemory, specifically +Control the behavior of devices which can access host memory, specifically whether data mapping copies data or uses self-mapping, such that the device directly accesses the corresponding host memory. Setting the environment variable to @code{DISABLED} does not make use of Unified Shared Memory by default. Setting it to @code{AUTO} conditionally enables Unified Shared -Memory on specific integrated GPUs (APUs, in the verbiage of AMD). In most +Memory (USM) on specific integrated GPUs (APUs, in the verbiage of AMD). In most workloads on these types of GPUs, self-mapping is expected to reduce runtime overhead. Lastly, setting it to @code{ENABLED} unconditionally attempts to make use of Unified Shared Memory if the device claims to support it. Of @@ -5261,11 +5261,16 @@ behavior to take precedence over the behavior set by this environment variable. For caveats around Unified Shared Memory, please consult the relevant sections in @ref{Offload-Target Specifics}. -When Unified Shared Memory is enabled, whether by a @code{requires} directive -or by this environment variable being set to @code{AUTO} or @code{ENABLED}, -global static variables that appear in a @code{declare target} directive are -not updated between device and host. That is, data copying under this model -for global static variables is ignored. +When unified-shared memory is enabled by this environment variable, the +following is currently @emph{unsupported} for only global static variables +that are specified in an @code{enter} clause on a @code{declare target} +directive: specifying those variables in a @code{map} clause with @code{always} +modifier on a @code{target} or @code{target data} directive. However, both +@code{map} clauses without @code{always} modifier and using the +@code{target update}, @code{target enter data}, or @code{target exit data} +directives is fully supported. When the unsupported case is +encountered, the program aborts with an error message. +@c FIXME: This is tracked in PR libgomp/127186 OpenACC is unaffected by the value of this environment variable, however, hybrid usage of OpenMP and OpenACC offloading results in the first used diff --git a/libgomp/target.c b/libgomp/target.c index 081701ca6342..f508040e0cad 100644 --- a/libgomp/target.c +++ b/libgomp/target.c @@ -4672,7 +4672,8 @@ gomp_exit_data (struct gomp_device_descr *devicep, size_t mapnum, /* Copy data for global static non-link variable, only; to be use with GOMP_RUNTIME_USM. Global variables with 'link' already point to the host and can be ignored. Only mapping with ALWAYS needs to be handled - as static has infinite ref count. */ + as static has infinite ref count - additionally handle pointer attachment + to global static pointer variables. */ static void gomp_target_enter_exit_data_usm (bool map_entering_p, @@ -4694,13 +4695,15 @@ gomp_target_enter_exit_data_usm (bool map_entering_p, for (size_t i = 0; i < mapnum; i++) { int kind = get_kind (short_mapkind, kinds, i); - if (!sizes[i]) - continue; - if (!GOMP_MAP_ALWAYS_P (typemask & kind)) + int map_kind = typemask & kind; + if ((!sizes[i] || !GOMP_MAP_ALWAYS_P (map_kind)) + && map_kind != GOMP_MAP_ATTACH) continue; struct splay_tree_key_s cur_node; cur_node.host_start = (uintptr_t) hostaddrs[i]; - cur_node.host_end = cur_node.host_start + sizes[i]; + cur_node.host_end = cur_node.host_start + + (map_kind == GOMP_MAP_ATTACH + ? sizeof (void *) : sizes[i]); splay_tree_key n = splay_tree_lookup (&devicep->mem_map, &cur_node); if (!n || n->refcount != REFCOUNT_INFINITY) continue; @@ -4719,7 +4722,24 @@ gomp_target_enter_exit_data_usm (bool map_entering_p, void *devaddr = (void *) (n->tgt->tgt_start + n->tgt_offset + cur_node.host_start - n->host_start); size_t size = cur_node.host_end - cur_node.host_start; - + uintptr_t addr; + if (map_kind == GOMP_MAP_ATTACH) + { + /* Check whether the pointee is a device variable; if so, we need + to use the device address. + NOTE: Assumes that that variable is 'always' mapped such that an + update is actually required. */ + cur_node.host_start = (uintptr_t) *(void **) hostaddr; + cur_node.host_end = cur_node.host_start + 1; + splay_tree_key n2 = splay_tree_lookup (&devicep->mem_map, &cur_node); + if (n2 && n2->refcount == REFCOUNT_INFINITY) + { + addr = (n2->tgt->tgt_start + n2->tgt_offset + + cur_node.host_start - n2->host_start); + hostaddr = &addr; + } + devaddr += sizes[i]; + } if (map_entering_p) gomp_copy_host2dev (devicep, NULL, devaddr, hostaddr, size, false, NULL); diff --git a/libgomp/testsuite/libgomp.c/usm_env_handling-3.c b/libgomp/testsuite/libgomp.c/usm_env_handling-3.c new file mode 100644 index 000000000000..45c5226e6ad8 --- /dev/null +++ b/libgomp/testsuite/libgomp.c/usm_env_handling-3.c @@ -0,0 +1,103 @@ +/* { dg-do run } */ +/* { dg-set-target-env-var GOMP_RUNTIME_USM "enabled" } */ + +// PR libgomp/127186 + +/* Check that attaching to a static global variable works with + GOMP_RUNTIME_USM="enabled", which otherwise implies self mapping. + + This testcase also works with 'requires self_maps', normal offloading, + and host fallback. */ + +#include <stdint.h> +#include <omp.h> + +#ifdef USE_SELF_MAPS + #pragma omp requires self_maps +#endif + + +// 'declare target' global variables: +int *p = nullptr; +int *pa[10] = {}; +int *p2 = nullptr; + +struct S { + int *p, *pa[10]; +} s; + +int Edev = 0; + +#pragma omp declare target enter(p, pa, p2, s, Edev) + + +// This function is executed either on the host (fallback) or on the device. +void +check_addr (intptr_t s_p_addr, intptr_t p_addr, intptr_t pa_addr, intptr_t s_pa_addr) +{ +#if 0 + __builtin_printf ("%p / %p / %p / %p / %p (%p) - device %s\n", + p, pa[1], s.p, s.pa[1], p2, (void*)&Edev, + __builtin_omp_is_initial_device() ? "host" : "nohost"); +#endif + + if (p == nullptr + || pa[1] == nullptr + || s.p != (void*) s_p_addr + || s.pa[1] == nullptr + || p2 != &Edev) + __builtin_abort (); + + if (p_addr != 0 && (void*) p_addr != p) + __builtin_abort (); + if (pa_addr != 0 && (void*) pa_addr != pa[1]) + __builtin_abort (); + if (s_pa_addr != 0 && (void*) s_pa_addr != s.pa[1]) + __builtin_abort (); +} + + +int +main () +{ + int A,B,C,D; + p = &A; + pa[1] = &B; + s.p = &C; + s.pa[1] = &D; + p2 = &Edev; + + // Check whether self mapping is active + bool self_mapping_p = false; + #pragma omp target map(to: self_mapping_p) + self_mapping_p = true; + __builtin_printf ("DEBUG: self_mapping_p = %d, num_devs = %d\n", + self_mapping_p, omp_get_num_devices ()); + +#if 0 + __builtin_printf ("%p / %p / %p / %p / %p (%p) - host\n", + p, pa[1], s.p, s.pa[1], p2, (void*)&Edev); +#endif + + // Map A, B, C, D (via the pointer) - and attach the pointee + // to the global variable: + + #pragma omp target enter data map(to: p, p[:1]) + #pragma omp target enter data map(to: pa[1][0]) + #pragma omp target enter data map(to: s.pa[1][0]) + #pragma omp target enter data map(to: p2[:0]) + + #pragma omp target enter data map(to: C) + // This one works for the wrong reason (i.e. always is the host address) + #pragma omp target update to(s.p) + + // Check whether the pointer attachment was successful on the default device + intptr_t s_p_addr = (intptr_t) s.p; + + intptr_t p_addr = self_mapping_p ? (intptr_t) &A : 0; + intptr_t pa_addr = self_mapping_p ? (intptr_t) &B : 0; + intptr_t s_pa_addr = self_mapping_p ? (intptr_t) &D : 0; + + #pragma omp target firstprivate(s_p_addr, p_addr, pa_addr, s_pa_addr) + check_addr (s_p_addr, p_addr, pa_addr, s_pa_addr); +}
