Hi all,
attached patch closes (hopefully!) the only loop hole in the
GOMP_RUNTIME_USM={enabled,auto} implementation that could cause
conforming programs to give invalid results; all other should either
work or fail with a runtime error message ('sorry, unsupported').
Namely, the patch handles pointer attachment to a global static
variables specified in a 'declare target' directive - and it does
so for both a pointee being only on the host and being another
global static 'declare target' variable.
It also updates the documentation to make clear what exactly is
(currently) unsupported. [The fixed issue + still missing bits
are tracked in PR127186.]
Comments, remarks, concerns before I commit it?
@Sandra: if you have time, esp. documentation comments are welcome;
the feature itself is documented at
https://gcc.gnu.org/onlinedocs/libgomp/GOMP_005fRUNTIME_005fUSM.html
and comment to the unmodified part as fine as well.
Tobias
PS: As mentioned, the caveat does not really apply anymore to in
the same way to the env var based self mapping and the 'omp requires'
one. The mentioned fix for the latter is provided by the patches
https://gcc.gnu.org/pipermail/gcc-patches/2026-August/728747.html and
https://gcc.gnu.org/pipermail/gcc-patches/2026-August/728748.html
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.
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 82f0a65bd50..92d82f2f66c 100644
--- a/libgomp/libgomp.texi
+++ b/libgomp/libgomp.texi
@@ -5238,42 +5238,47 @@ Enable debugging output. The variable should be set to @code{0}
If enabled, some debugging output is printed during execution.
This is currently not specified in more detail, and subject to change.
@end table
@node GOMP_RUNTIME_USM
@section @env{GOMP_RUNTIME_USM} -- Control runtime Unified Shared Memory application
@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
additional note, specifying the @code{unified_shared_memory} or
@code{self_maps} clauses with the @code{requires} directive causes that
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
framework taking precedent. As an example, if one uses OpenMP offloading and
then later OpenACC offloading, the capabilities the device possesses reflect
the original OpenMP model.
@item @emph{See also}:
@ref{Offload-Target Specifics}, @ref{AMD Radeon}
@end table
diff --git a/libgomp/target.c b/libgomp/target.c
index 081701ca634..f508040e0ca 100644
--- a/libgomp/target.c
+++ b/libgomp/target.c
@@ -4664,70 +4664,90 @@ gomp_exit_data (struct gomp_device_descr *devicep, size_t mapnum,
{
free (hostaddrs);
free (sizes);
free (kinds);
free (iterator_count);
}
}
/* 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,
struct gomp_device_descr *devicep,
size_t mapnum, void **hostaddrs,
size_t *sizes, unsigned short *kinds)
{
constexpr bool short_mapkind = true; /* OpenMP */
constexpr int typemask = 0xff;
size_t *iterator_count = NULL;
bool iterators_p = gomp_merge_iterator_maps (&mapnum, &hostaddrs, &sizes,
(void **)&kinds, &iterator_count);
gomp_mutex_lock (&devicep->lock);
if (devicep->state == GOMP_DEVICE_FINALIZED)
{
gomp_mutex_unlock (&devicep->lock);
return;
}
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;
if (n->host_start > cur_node.host_start
|| n->host_end < cur_node.host_end)
{
gomp_mutex_unlock (&devicep->lock);
gomp_fatal ("Trying to update [%p..%p) object when only [%p..%p) is "
"mapped",
(void *) cur_node.host_start,
(void *) cur_node.host_end,
(void *) n->host_start,
(void *) n->host_end);
}
void *hostaddr = (void *) cur_node.host_start;
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);
else
gomp_copy_dev2host (devicep, NULL, hostaddr, devaddr, size);
}
gomp_mutex_unlock (&devicep->lock);
if (iterators_p)
{
free (hostaddrs);
free (sizes);
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 00000000000..45c5226e6ad
--- /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);
+}
