[gcc r17-3807] libgomp: OpenMP - improve GOMP_RUNTIME_USM handling
Tobias Burnus
burnus@gcc.gnu.org
Tue Sep 1 08:39:34 GMT 2026
https://gcc.gnu.org/g:82104cde7ac0ba2a80194048ca309e26c655ca6d
commit r17-3807-g82104cde7ac0ba2a80194048ca309e26c655ca6d
Author: Tobias Burnus <tburnus@baylibre.com>
Date: Tue Sep 1 10:35:47 2026 +0200
libgomp: OpenMP - improve GOMP_RUNTIME_USM handling
Follow up to the commit 'libgomp: Add Controls For Runtime USM'
r17-3803-g486b381b7ae6c1 that adds the GOMP_RUNTIME_USM environemnt
variable with values 'disabled', 'enabled', and 'auto'.
For devices with OFFLOAD_CAP_SHARED_MEM capability, do not add
'link' variables to the splay tree as those are initialized to point
to the host variable and will never be touched afterwards.
The rest of the code: OFFLOAD_CAP_SHARED_MEM implies that self mapping
is possible; however, unless all declare-target global variables are
turned into 'link' variables (as effectively mandated for 'requires
self_maps'), data mapping still needs to be honored for those.
For 'requires self_maps' and 'requires unified_shared_memory', a pending
patch takes care of this in the compiler. However, for
GOMP_RUNTIME_USM={enabled,auto} it needs to be handled in the runtime.
This commit handles it in 'target update' and 'target enter/exit data',
but not in 'target data' and 'target'. For the latter two, a fatal error
(sorry) is printed if encountering such a variable with REFCOUNT_INFINITY,
but only if the 'always' map-type modifier is used for it.
NOTE: There is a loop hole in the diagnostic as 'target map(to: A)' will
use the host variable directly not the device variable and without 'always'
no diagnostic is shown. However, the lacking diagnostic is only an issue
if the variable is both using in the target region and directly accessing
the global variable in a called function, otherwise its value is either
inconsistent before or the code should work.
libgomp/ChangeLog:
* target.c (gomp_load_image_to_device): With shared memory, do not
add 'link' variables to the splay tree.
(GOMP_target_update_ext): For OFFLOAD_CAP_SHARED_MEM, still call
gomp_update unless requires self_maps/USM is set.
(gomp_update): For OFFLOAD_CAP_SHARED_MEM, only handle global
static variables (other than 'link').
(GOMP_target_ext, GOMP_target_data_ext): For OFFLOAD_CAP_SHARED_MEM
without 'requires self_maps/USM', abort with a 'sorry' for
global static variables other than 'link' as the MAP clause would
otherwise be ignored.
(gomp_target_enter_exit_data_usm): New; handle global static variables
other than 'link' with OFFLOAD_CAP_SHARED_MEM.
(GOMP_target_enter_exit_data): Call it in that case, unless there
is a 'requires self_maps/USM'.
* testsuite/libgomp.c/usm_env_handling-1.c: New test.
Diff:
---
libgomp/target.c | 170 +++++++++++++++++++++--
libgomp/testsuite/libgomp.c/usm_env_handling-1.c | 132 ++++++++++++++++++
2 files changed, 293 insertions(+), 9 deletions(-)
diff --git a/libgomp/target.c b/libgomp/target.c
index cd41c6552a2d..f688f33378d6 100644
--- a/libgomp/target.c
+++ b/libgomp/target.c
@@ -2540,6 +2540,12 @@ gomp_update (struct gomp_device_descr *devicep, size_t mapnum, void **hostaddrs,
splay_tree_key n = splay_tree_lookup (&devicep->mem_map, &cur_node);
if (n)
{
+ /* With shared memory, static global variables still need to be
+ handled, except for 'link', pointing already to the host
+ variable. */
+ if ((devicep->capabilities & GOMP_OFFLOAD_CAP_SHARED_MEM)
+ && n->refcount != REFCOUNT_INFINITY)
+ continue;
int kind = get_kind (short_mapkind, kinds, i);
if (n->host_start > cur_node.host_start
|| n->host_end < cur_node.host_end)
@@ -2820,17 +2826,22 @@ gomp_load_image_to_device (struct gomp_device_descr *devicep, unsigned version,
k->refcount = is_link_var ? REFCOUNT_LINK : REFCOUNT_INFINITY;
k->dynamic_refcount = 0;
k->aux = NULL;
+
+ /* If GOMP_OFFLOAD_CAP_SHARED_MEM: For link variable, set the device
+ variable to point to the host variable (as required for self_maps);
+ as no data transfer is needed for 'link' variables with shared memory,
+ skip adding it to the splay tree. */
+ if (is_link_var && devicep->capabilities & GOMP_OFFLOAD_CAP_SHARED_MEM)
+ {
+ gomp_copy_host2dev (devicep, NULL, (void *) target_var->start,
+ &k->host_start, sizeof (void *), false, NULL);
+ continue;
+ }
+
array->left = NULL;
array->right = NULL;
splay_tree_insert (&devicep->mem_map, array);
array++;
-
- if (is_link_var
- && ((omp_requires_mask
- & (GOMP_REQUIRES_UNIFIED_SHARED_MEMORY | GOMP_REQUIRES_SELF_MAPS))
- || (devicep->capabilities & GOMP_OFFLOAD_CAP_SHARED_MEM)))
- gomp_copy_host2dev (devicep, NULL, (void *) target_var->start,
- &k->host_start, sizeof (void *), false, NULL);
}
/* Last entry is for the ICV struct variable; if absent, start = end = 0. */
@@ -3699,6 +3710,34 @@ GOMP_target_ext (int device, void (*fn) (void *), size_t mapnum,
struct gomp_offload_session *session
= gomp_offload_session_new (devicep, gomp_alloca);
+ if ((devicep->capabilities & GOMP_OFFLOAD_CAP_SHARED_MEM)
+ && !(omp_requires_mask
+ & (GOMP_REQUIRES_UNIFIED_SHARED_MEMORY | GOMP_REQUIRES_SELF_MAPS)))
+ {
+ /* When USM is anabled (cf. GOMP_RUNTIME_USM) but enter clauses are not
+ automatically turned into 'link' clauses in the compiler, they still
+ need to be handled in map clauses.
+ TODO: Implement this feature; for now fail with an error if those
+ are encountered.
+ FIXME: Not only 'always' has to be handled as variables in the target
+ region and an (indirect) access to the global variable in a function
+ call must refer to the same variable. */
+ constexpr bool short_mapkind = true; /* OpenMP */
+ constexpr int typemask = 0xff;
+ for (size_t i = 0; i < mapnum; i++)
+ if (GOMP_MAP_ALWAYS_P (typemask & get_kind (short_mapkind, kinds, i)))
+ {
+ 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];
+ splay_tree_key n = splay_tree_lookup (&devicep->mem_map, &cur_node);
+ if (n && n->refcount == REFCOUNT_INFINITY)
+ gomp_fatal ("Sorry, unimplemented: TARGET with ALWAYS "
+ "modifier for device static variables when "
+ "GOMP_RUNTIME_USM is enabled");
+ }
+ }
+
if (devicep->capabilities & GOMP_OFFLOAD_CAP_SHARED_MEM)
{
if (!fpc_done)
@@ -4352,6 +4391,33 @@ GOMP_target_data_ext (int device, size_t mapnum, void **hostaddrs,
{
struct gomp_device_descr *devicep = resolve_device (device, true);
+ if (devicep
+ && ((devicep->capabilities & GOMP_OFFLOAD_CAP_SHARED_MEM)
+ && !(omp_requires_mask
+ & (GOMP_REQUIRES_UNIFIED_SHARED_MEMORY
+ | GOMP_REQUIRES_SELF_MAPS))))
+ {
+ /* When USM is anabled (cf. GOMP_RUNTIME_USM) but enter clauses are not
+ automatically turned into 'link' clauses in the compiler, they still
+ need to be handled in map clauses.
+ TODO: Implement this feature; for now fail with an error if those
+ are encountered. */
+ constexpr bool short_mapkind = true; /* OpenMP */
+ const int typemask = 0xff;
+ for (size_t i = 0; i < mapnum; i++)
+ if (GOMP_MAP_ALWAYS_P (typemask & get_kind (short_mapkind, kinds, i)))
+ {
+ 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];
+ splay_tree_key n = splay_tree_lookup (&devicep->mem_map, &cur_node);
+ if (n && n->refcount == REFCOUNT_INFINITY)
+ gomp_fatal ("Sorry, unimplemented: TARGET DATA with ALWAYS "
+ "modifier for device static variables when "
+ "GOMP_RUNTIME_USM is enabled");
+ }
+ }
+
if (devicep == NULL
|| !(devicep->capabilities & GOMP_OFFLOAD_CAP_OPENMP_400)
|| devicep->capabilities & GOMP_OFFLOAD_CAP_SHARED_MEM)
@@ -4444,9 +4510,16 @@ GOMP_target_update_ext (int device, size_t mapnum, void **hostaddrs,
}
}
+ /* For GOMP_OFFLOAD_CAP_SHARED_MEM, still static global varibles
+ that lack the link clause need to be handled, unless the
+ USM/self_maps requirement is set (as then all such variables
+ have the link clause); cf. also GOMP_RUNTIME_USM_DISABLED).
+ Note that USM/self_maps implies that all available devices
+ support CAP_SHARED_MEM. */
if (devicep == NULL
|| !(devicep->capabilities & GOMP_OFFLOAD_CAP_OPENMP_400)
- || devicep->capabilities & GOMP_OFFLOAD_CAP_SHARED_MEM)
+ || (omp_requires_mask
+ & (GOMP_REQUIRES_UNIFIED_SHARED_MEMORY | GOMP_REQUIRES_SELF_MAPS)))
return;
struct gomp_thread *thr = gomp_thread ();
@@ -4596,6 +4669,73 @@ 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. */
+
+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))
+ 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];
+ 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;
+
+ 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);
+ free (kinds);
+ free (iterator_count);
+ }
+}
+
void
GOMP_target_enter_exit_data (int device, size_t mapnum, void **hostaddrs,
size_t *sizes, unsigned short *kinds,
@@ -4651,7 +4791,8 @@ GOMP_target_enter_exit_data (int device, size_t mapnum, void **hostaddrs,
if (devicep == NULL
|| !(devicep->capabilities & GOMP_OFFLOAD_CAP_OPENMP_400)
- || devicep->capabilities & GOMP_OFFLOAD_CAP_SHARED_MEM)
+ || (omp_requires_mask
+ & (GOMP_REQUIRES_UNIFIED_SHARED_MEMORY | GOMP_REQUIRES_SELF_MAPS)))
return;
struct gomp_thread *thr = gomp_thread ();
@@ -4672,6 +4813,17 @@ GOMP_target_enter_exit_data (int device, size_t mapnum, void **hostaddrs,
}
}
+ /* When USM is anabled (cf. GOMP_RUNTIME_USM) but enter clauses are not
+ automatically turned into 'link' clauses in the compiler, they still
+ need to be honored. */
+ if (devicep->capabilities & GOMP_OFFLOAD_CAP_SHARED_MEM)
+ {
+ bool map_entering_p = (flags & GOMP_TARGET_FLAG_EXIT_DATA) == 0;
+ gomp_target_enter_exit_data_usm (map_entering_p, devicep, mapnum,
+ hostaddrs, sizes, kinds);
+ return;
+ }
+
htab_t refcount_set = htab_create (mapnum);
/* The variables are mapped separately such that they can be released
diff --git a/libgomp/testsuite/libgomp.c/usm_env_handling-1.c b/libgomp/testsuite/libgomp.c/usm_env_handling-1.c
new file mode 100644
index 000000000000..f909851ff6fa
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c/usm_env_handling-1.c
@@ -0,0 +1,132 @@
+/* { dg-do run } */
+/* { dg-set-target-env-var GOMP_RUNTIME_USM "enabled" } */
+
+/* Check that a global static variable is consistently updated on the device.
+
+ With GOMP_RUNTIME_USM=enabled, for devices supporting host access
+ (shared memory), self maps are used - however, for 'declare target'
+ variables, actual data copying still needs to happen.
+
+ Check that this works. */
+
+// FIXME: There are known issues with TARGET and TARGET DATA
+// The code still makes the wrong cases pass (exit code 0) but still prints an ERROR.
+
+#include <omp.h>
+#include <stdio.h>
+#include <stdlib.h>
+#include <string.h>
+
+
+int A[1] = {1};
+#pragma omp declare target enter(A)
+
+
+int
+getA ()
+{
+ return A[0];
+}
+
+void
+checkA (int expected, int also_ok, const char *dir)
+{
+ if (expected != A[0])
+ printf ("ERROR: %d (expected: %d) %s running on the %s\n", A[0], expected,
+ dir, omp_is_initial_device () ? "HOST" : "DEVICE");
+ if (expected != A[0] && also_ok != A[0])
+ __builtin_abort ();
+}
+
+void
+incA ()
+{
+ ++A[0];
+}
+
+
+int
+main ()
+{
+ // Debug output
+ const char *env_var = getenv ("GOMP_RUNTIME_USM");
+ bool is_shared_mem = false;
+ #pragma omp target map(to: is_shared_mem)
+ is_shared_mem = true;
+ bool is_usm = (is_shared_mem && env_var
+ && strcasecmp ("enabled", env_var) == 0
+ && omp_get_num_devices () > 0);
+ printf ("DEBUG: GOMP_RUNTIME_USM = %s, is_shared_mem = %s -> usm = %s\n",
+ env_var ? env_var : "<unset>", is_shared_mem ? "true" : "false",
+ is_usm ? "true" : "false");
+
+
+ // Check that 'TARGET UPDATE' works
+ A[0] = 2;
+ #pragma omp target update to(A)
+ #pragma omp target
+ checkA (2, 2, "after TARGET");
+
+
+ // Check that 'TARGET ENTER DATA' works
+ A[0] = 3;
+ #pragma omp target enter data map(always,to: A)
+
+ #pragma omp target
+ {
+ checkA (3, 3, "after TARGET ENTER DATA");
+ incA ();
+ }
+
+ // Check that 'TARGET EXIT DATA' works
+ #pragma omp target exit data map(always,from: A)
+ checkA (4, 4, "after TARGET EXIT DATA");
+
+ // FIXME: Early return because of the issues below.
+ if (is_usm)
+ {
+ printf ("SKIPPED: RUNTIME USM does not support TARGET and TARGET DATA and "
+ "fails with a runtime error - see testcase\n");
+ return 0;
+ }
+
+ // The following FIXME expect that the code just works without
+ // Aborting - if you comment the gomp_fatal error, it will pass as follows
+
+ // Check that 'TARGET DATA' works
+ A[0] = 5;
+ #pragma omp target data map(always, tofrom: A)
+ {
+ #pragma omp target
+ {
+ // FIXME: Should print 5 - with USM it would print the old device value 4, but currently aborts with:
+ // libgomp: Sorry, unimplemented: TARGET DATA with ALWAYS modifier for device static variables when GOMP_RUNTIME_USM is enabled
+ checkA (5, !is_usm ? 5 : 4 /* FIXME: old device value */,
+ "in TARGET DATA");
+ incA ();
+ }
+ }
+ // FIXME: Should print 6 - with USM prints old host value 5
+ checkA (6, !is_usm ? 6 : 5 /* FIXME: old host value */,
+ "after TARGET DATA");
+
+ // Check that 'TARGET' works
+ A[0] = 7;
+ #pragma omp target map(always,tofrom: A)
+ {
+ // FIXME: Should print 7 - with USM it would print the old device value 5, but currently aborts with:
+ // libgomp: Sorry, unimplemented: TARGET with ALWAYS modifier for device static variables when GOMP_RUNTIME_USM is enabled
+
+ // The following accesses the global static variable
+ checkA (7, !is_usm ? 7 : 5 /* FIXME: old device value */,
+ "in TARGET");
+ incA ();
+
+ // The following accesses the variable as translated by the TARGET construct
+ if ((!is_usm && A[0] != 8) || (is_usm && A[0] != 7 /* FIXME: uses host var, not device var */))
+ __builtin_abort ();
+ ++A[0];
+ }
+ checkA (9, !is_usm ? 9 : 8 /* FIXME: old host value, incremented once on the device */,
+ "after TARGET");
+}
More information about the Gcc-cvs
mailing list