Hi Thomas,
coming back to the patch itself – having sent comments in previous email.
As mentioned, I think you either want to have:
* Check whether the runtime knows that USM is supported for the default
device (that is a non-host device), i.e. a no-host device is available
with 'required unified_shared_memory/self_maps')
* Check the above and check in addition that self-mapping happens with
'required unified_shared_memory'.
I am not sure whether we need the second one, given that using 'requires
self_maps' or a runtime check in the program works. (IMHO, it does not
harm but is not really needed.) But the first one is surely useful.
In any case, you add the second one, only – while
'map-alloc-comp-9-usm.f90' requires the first one.
* * *
Thomas Schwinge wrote:
From 46fc59b5cdaa42c4dc9edaee7d52194c1f45b6b3 Mon Sep 17 00:00:00 2001
From: Thomas Schwinge<tschwi...@baylibre.com>
Date: Fri, 9 May 2025 15:05:57 +0200
Subject: [PATCH] Add effective-target 'offload_device_usm',
'libgomp.c-c++-common/target-usm-1.c'
Also use the new effective-target 'offload_device_usm' for restricting
'libgomp.fortran/map-alloc-comp-9-usm.f90' testing; the latter being a USM
variant of 'libgomp.fortran/map-alloc-comp-9.f90'.
The wording for effective-target should be expanded as it is not clear
what is tested for. (Cf. comments above).
diff --git a/libgomp/testsuite/lib/libgomp.exp
b/libgomp/testsuite/lib/libgomp.exp
index a620f8c2a09..cd32be1ca68 100644
--- a/libgomp/testsuite/lib/libgomp.exp
+++ b/libgomp/testsuite/lib/libgomp.exp
@@ -449,6 +449,8 @@ proc check_effective_target_offload_device_nonshared_as { }
{
}
# Return 1 if offload device is available and it has shared address space.
+# This doesn't consider whether '#pragma omp requires unified_shared_memory'
+# may be used to switch into shared-memory mode.
This one actually does not check whether USM works but
whether self-mapping happens by default.
For instance, when the following code does not fail,
int a = 6;
int *p = &a;
#pragma omp target firstprivate(p)
if (*p != 6) __builtin_abort ();
... there is support for USM (for at least some host memory).
However, this does not imply hat 'libgomp' knows about this
(i.e. 'requires unified_shared_memory' might exclude that device)
and even if libgomp knows it, it might still default to copy
mapping.
Example for the former:
* Two AMD GPUs, only one supporting USM. The current check
is not per device but for all AMD GPUs.
* Example for the second: All GPUs currently. The plan is
to default to self mapping for APUs (MI300a, Grace-Hopper,
...) and to permit to switch to it by an env var, but
that still has to be implemented.
BTW: In principle, we could also add such a check. Although, it
might have false positive in case some but not all host memory
is accessible.
That's actually what I want to have formap-alloc-comp-9.f90 Like: { dg-additional-options "-DUSM_SUPPORTED" {
target ... } } while the requires USM (map-alloc-comp-9.f90)is a partial work around.
* * *
+# Return 1 if, with '#pragma omp requires unified_shared_memory' in effect, an
+# USM-capable offload device is available (not considering host-fallback
+# execution).
+proc check_effective_target_offload_device_usm { } {
Probably simpler to use 'self_maps' and state so in the name?
+ return [check_runtime_nocache offload_device_usm {
+ #pragma omp requires unified_shared_memory
+ #include <omp.h>
+ int main ()
+ {
+ int a;
This variable is not initialized! You surely want to use "= 0" here!
+ #pragma omp target map(from: a)
This checks currently whether self-maps is the default with USM.
With 'requires self_maps' or with 'map(tofrom:' it would check whether
USM is (known to libgomp to be) supported by the device.
At least for map-alloc-comp-9-usm.f90, self map is not required (and
actually not wanted) – but host-memory access is required.
* * *
+++ b/libgomp/testsuite/libgomp.c-c++-common/target-usm-1.c
@@ -0,0 +1,46 @@
+/* If we have an offload device that is capable of USM...
+ { dg-do run { target offload_device_usm } } */
+
+/* ..., and we request USM... */
+#pragma omp requires unified_shared_memory
+/* (..., which in the GCC implementation equals 'self_maps'...) */
Can you please avoid mixing USM with self_maps? If I had more time
during in GCC 15, that assumption wouldn't be true anymore.
Thus, if you just care about USM, use 'requires unified_shared_memory'.
But if you assume self mapping, use 'requires self_maps'.
[Adding APU detection and an env var is planned for GCC 16]
* * *
However, this is a rather safe assumption (at least in GCC): USM
supported → self_maps supported – and vice versa, even though OpenMP
does not guarantee it and some system might not imply it. self_maps
supported → USM supported is surely always the case.
At least as I intent to get this working in GCC, both will work.
+++ b/libgomp/testsuite/libgomp.fortran/map-alloc-comp-9-usm.f90
@@ -1,3 +1,4 @@
+! { dg-skip-if {} { ! offload_device_usm } }
! { dg-additional-options "-cpp -DUSE_USM_REQUIREMENT=1 -Wno-openmp" }
As mentioned, this patch strictly does not depend on self maps – and if
we had an env var to disable self mapping, it would use it.
(As mentioned above, I intent to add an env var for this.)
Tobias
