Now committed as r15-3836-g4cb20dc043cf70
Contrary to the originally posted patch, it also acts on the newer/newly
added 'omp requires self_maps'.
In the area of (unified-)shared memory/self maps, the next step seems to
be to do still mapping for static variables – before moving to
refinements like how to handle implicit 'declare target' for static
variables, …
For this piece of code, we also want to run it for APUs even when no USM
has been requested, avoid adding those to the mapping table (for self
maps) and do a more efficient mapping (e.g. memcpy or avoid multiple locks).
Tobias
Tobias Burnus wrote:
short version: I think the patch as posted is fine and no action
beyond is needed for this one issue.
See below for the long version.
Possibly modifications (now or as follow up):
- using memcpy + or let the plugin do it
- not adding link variables to the splay tree with 'USM'.
Thomas Schwinge wrote:
Tested on x86-64-gnu-linux and nvptx offloading (that supports USM).
(I yet have to set up such a USM configuration...)
You already used an USM config, e.g., when running gfx90a (likewise:
gfx90c), except that USM on mainline it currently only works if you
explicitly set 'export HSA_XNACK=1'.
For Nvptx, you need a post-Volta GPU with the open-kernels driver,
which is for newer driver versions the default.
* * *
Do I understand correctly that even if
'GOMP_REQUIRES_UNIFIED_SHARED_MEMORY', we cannot just skip all the
'mem_map' setup in 'gomp_load_image_to_device' etc., because we're not
(yet?) setting 'GOMP_OFFLOAD_CAP_SHARED_MEM'?
We actually do set GOMP_OFFLOAD_CAP_SHARED_MEM with 'requires
unified_shared_memory'.
But, indeed, we cannot skip the memory mapping parts – due to the way
we handle static variables.
* * *
+
+ if (is_link_var
+ && (omp_requires_mask & GOMP_REQUIRES_UNIFIED_SHARED_MEMORY))
+ gomp_copy_host2dev (devicep, NULL, (void *) target_var->start,
+ &k->host_start, sizeof (void *), false, NULL);
}
Calling 'gomp_copy_host2dev' looks a bit funny given we've just
determined USM (..., but I'm not asking for plain 'memcpy').
I guess a plain memcpy would do as well. [Assuming that the device's
static variable is host accessible, which it probably is and should be.]
I add it to my to-do list for USM-related tasks to change this;
possibly moving it to the plugin side has some advantages? Possibly
not adding it to the splay tree if not needed. (Cf. below for env var
discussion.)
Regarding the unload: For 'declare target link(A)', we have, e.g.,
'static int *A' on the device side. Thus, we could do 'A = NULL' – and
rather should do 'A = {clobber}', but that's rather pointless in
general and especially when unloading the image.
What's the advantage/rationale of doing this here vs. in
'gomp_map_vars_internal' for 'REFCOUNT_LINK'? (May be worth a source
code comment?)
(A, B, C refers to the following example.)
We don't see 'A' (or 'B') in the GOMP_target_ext call and thus not in
gomp_map_vars_internal.
Besides: We only want to do the initialization once and not every time
gomp_map_vars_internal is called.
I think the following program may help to understand the issue and the
patch better.
Note: While A, B, C are 'int …[3]' on the host, on the device we only
have 'int B[3]' while for A it's 'int *A' and C only exists on the host.
* * *
#pragma requires unified_shared_memory
static int A[3], B[3], C[3];
#pragma omp declare target link(A) enter(B)
#pragma omp begin declare target
void f(int *p)
{
A[2] += B[2] + p[2]; // p points to the host's C variable
}
#pragma omp end declare target
void foo(int dev) {
int *ptr = C;
#pragma omp target firstprivate(ptr) device(dev)
f (ptr);
}
* * *
Here, 'ptr' (and thus 'p') point to the host 'C' variable, both before
the target
region and inside the target region.
'B' points to the device local version of the variable.
And 'A' on a non-host device is likely to be NULL ('static int *A' +
.BSS) before this patch.
Or pointing to the host's 'A' with this patch.
* * *
With A pointing to the host version (and likewise 'p' pointing to the
host C), host fallback
and device version yield identical result for 'A' and for 'C' (via
ptr/p). — However, 'B' on
host and non-host device have nothing in common. While that might be
fine, in general it is not.
Hence, in order to get for a .BSS valued 'B' the same result on host
and device, we need, e.g.
#pragma omp data map(always: B) device(dev)
foo (dev);
to call 'foo' to ensure that the two 'B' are in sync.
* * *
Code wise, this means that with GOMP_OFFLOAD_CAP_SHARED_MEM, we still
have
to apply the map for 'declare target enter(…)' variables, except if host
and device share the same code – but that should only be the case for
host fallback (= initial device) and, possibly,
GOMP_OFFLOAD_CAP_NATIVE_EXEC.
* * *
NOTE: OpenMP still permits to honor explicit 'map' with 'requires
unified_shared_memory',
only with 'self' maps, copying the data in 'map' is explicitly
disallowed.
* * *
This patch + honoring 'map' for static (non-'link'?) variables even with
GOMP_OFFLOAD_CAP_SHARED_MEM where the main items for the USM follow-up
patches,
I meant by "More USM cleanup/fixes/extensions to make it _more_
useful" on slide 16
of
https://gcc.gnu.org/wiki/cauldron2024#cauldron2024talks.openmp_openacc_and_offloading_in_gcc
Plus, to go a bit beyond:
- offering a flag to change 'declare target enter(…)' to 'link(…)'
[RFC: enable it by default for 'requires unified_shared_memory'?]
- switching to GOMP_OFFLOAD_CAP_SHARED_MEM by default for APUs
(= same memory controller) for performance
- Adding a GOMP_ environment variable to toggle between mapping vs. USM
access on systems not detected as being APUs. (That is: systems that
support USM but use an interconnect or page migration for the memory
access. Possibly, also overriding the USM detection for systems which
can access the host memory but due to some own memory are not
recognized
(→ device attributes) as being USM devices.
And possibly also forcing to honor explicit maps with requires
(Example for the latter is Andrew's gfx1103, which reports
HSA_AMD_SYSTEM_INFO_SVM_ACCESSIBLE_BY_DEFAULT = 'false'; I assume it
can still access all host memory, but I might be wrong.)
- Documenting how GCC handles this in libgomp.texi
BTW: See
https://gcc.gnu.org/onlinedocs/libgomp/Offload-Target-Specifics.html
for how USM is detected for nvptx + gcn devices.
* * *
For completeness, I also intent to look at Andrew's pinned
memory/(pseudo)USM
patches; they are useful but address other aspects as those listed above
→ https://gcc.gnu.org/pipermail/gcc-patches/2024-June/654331.html
→ https://gcc.gnu.org/pipermail/gcc-patches/2024-May/652932.html
commit 4cb20dc043cf70b8a1b4846c86599cc1ff9680d9
Author: Tobias Burnus <tbur...@baylibre.com>
Date: Tue Sep 24 17:41:39 2024 +0200
libgomp: with USM, init 'link' variables with host address
If requires unified_shared_memory or self_maps is set, make
'declare target link' variables to point initially to the host pointer.
libgomp/ChangeLog:
* target.c (gomp_load_image_to_device): For requires
unified_shared_memory, update 'link' vars to point to the host var.
* testsuite/libgomp.c-c++-common/target-link-3.c: New test.
* testsuite/libgomp.c-c++-common/target-link-4.c: New test.
---
libgomp/target.c | 6 +++
.../testsuite/libgomp.c-c++-common/target-link-3.c | 52 ++++++++++++++++++++++
.../testsuite/libgomp.c-c++-common/target-link-4.c | 52 ++++++++++++++++++++++
3 files changed, 110 insertions(+)
diff --git a/libgomp/target.c b/libgomp/target.c
index 6918694a843..cf62af61f3b 100644
--- a/libgomp/target.c
+++ b/libgomp/target.c
@@ -2454,6 +2454,12 @@ gomp_load_image_to_device (struct gomp_device_descr *devicep, unsigned version,
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)))
+ 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. */
diff --git a/libgomp/testsuite/libgomp.c-c++-common/target-link-3.c b/libgomp/testsuite/libgomp.c-c++-common/target-link-3.c
new file mode 100644
index 00000000000..c707b38b7d4
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c-c++-common/target-link-3.c
@@ -0,0 +1,52 @@
+/* { dg-do run } */
+
+#include <stdint.h>
+#include <omp.h>
+
+#pragma omp requires unified_shared_memory
+
+int A[3] = {-3,-4,-5};
+static int q = -401;
+#pragma omp declare target link(A, q)
+
+#pragma omp begin declare target
+void
+f (uintptr_t *pA, uintptr_t *pq)
+{
+ if (A[0] != 1 || A[1] != 2 || A[2] != 3 || q != 42)
+ __builtin_abort ();
+ A[0] = 13;
+ A[1] = 14;
+ A[2] = 15;
+ q = 23;
+ *pA = (uintptr_t) &A[0];
+ *pq = (uintptr_t) &q;
+}
+#pragma omp end declare target
+
+int
+main ()
+{
+ uintptr_t hpA = (uintptr_t) &A[0];
+ uintptr_t hpq = (uintptr_t) &q;
+ uintptr_t dpA, dpq;
+
+ A[0] = 1;
+ A[1] = 2;
+ A[2] = 3;
+ q = 42;
+
+ for (int i = 0; i <= omp_get_num_devices (); ++i)
+ {
+ #pragma omp target device(device_num: i) map(dpA, dpq)
+ f (&dpA, &dpq);
+ if (hpA != dpA || hpq != dpq)
+ __builtin_abort ();
+ if (A[0] != 13 || A[1] != 14 || A[2] != 15 || q != 23)
+ __builtin_abort ();
+ A[0] = 1;
+ A[1] = 2;
+ A[2] = 3;
+ q = 42;
+ }
+}
diff --git a/libgomp/testsuite/libgomp.c-c++-common/target-link-4.c b/libgomp/testsuite/libgomp.c-c++-common/target-link-4.c
new file mode 100644
index 00000000000..785055e216d
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c-c++-common/target-link-4.c
@@ -0,0 +1,52 @@
+/* { dg-do run } */
+
+#include <stdint.h>
+#include <omp.h>
+
+#pragma omp requires self_maps
+
+int A[3] = {-3,-4,-5};
+static int q = -401;
+#pragma omp declare target link(A, q)
+
+#pragma omp begin declare target
+void
+f (uintptr_t *pA, uintptr_t *pq)
+{
+ if (A[0] != 1 || A[1] != 2 || A[2] != 3 || q != 42)
+ __builtin_abort ();
+ A[0] = 13;
+ A[1] = 14;
+ A[2] = 15;
+ q = 23;
+ *pA = (uintptr_t) &A[0];
+ *pq = (uintptr_t) &q;
+}
+#pragma omp end declare target
+
+int
+main ()
+{
+ uintptr_t hpA = (uintptr_t) &A[0];
+ uintptr_t hpq = (uintptr_t) &q;
+ uintptr_t dpA, dpq;
+
+ A[0] = 1;
+ A[1] = 2;
+ A[2] = 3;
+ q = 42;
+
+ for (int i = 0; i <= omp_get_num_devices (); ++i)
+ {
+ #pragma omp target device(device_num: i) map(dpA, dpq)
+ f (&dpA, &dpq);
+ if (hpA != dpA || hpq != dpq)
+ __builtin_abort ();
+ if (A[0] != 13 || A[1] != 14 || A[2] != 15 || q != 23)
+ __builtin_abort ();
+ A[0] = 1;
+ A[1] = 2;
+ A[2] = 3;
+ q = 42;
+ }
+}
