[gcc r17-3917] libgomp: GOMP_RUNTIME_USM - fix pointer attachment + update .texi [PR127186]
Tobias Burnus
burnus@gcc.gnu.org
Fri Sep 4 07:21:09 GMT 2026
https://gcc.gnu.org/g:1f0018091064c26334f4dd67b3472f93c62b0b80
commit r17-3917-g1f0018091064c26334f4dd67b3472f93c62b0b80
Author: Tobias Burnus <tburnus@baylibre.com>
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);
+}
More information about the Gcc-cvs
mailing list