Hi Thomas,
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
Tobias
