From 1cf9fda4936de54198858b8f54cd9707a3725f4e Mon Sep 17 00:00:00 2001 From: Andrew Stubbs Date: Mon, 8 Dec 2025 16:18:59 +0000 Subject: [PATCH] amdgcn: Adjust failure mode for gfx908 USM Unified Shared Memory does not appear to work well on gfx908, which is why we disabled xnack by default. For this reason it makes sense to inform the user as compile time, but this is causing trouble in the testsuite which assumes that USM only fails at runtime. This patch changes the gfx908 compile time message to a warning only (in case some other target does this differently), and prevents the tests from attempting to run in host-fallback mode (given that that is not what they are trying to test). It also changes the existing warning to only fire once. The patch assumes that effective target "omp_usm" also implies self-maps. gcc/ChangeLog: * config/gcn/gcn.cc (gcn_init_cumulative_args): Only warn once. Use "required" instead of "enabled" in the warning. * config/gcn/mkoffload.cc (process_asm): Warn, don't error. Use "required" instead of "on" in the warning. libgomp/ChangeLog: * testsuite/lib/libgomp.exp (check_effective_target_omp_usm): New. * testsuite/libgomp.c++/target-std__array-concurrent-usm.C: Require working Unified Shared Memory to run the test. * testsuite/libgomp.c++/target-std__bitset-concurrent-usm.C: Likewise. * testsuite/libgomp.c++/target-std__deque-concurrent-usm.C: Likewise. * testsuite/libgomp.c++/target-std__forward_list-concurrent-usm.C: Likewise. * testsuite/libgomp.c++/target-std__list-concurrent-usm.C: Likewise. * testsuite/libgomp.c++/target-std__map-concurrent-usm.C: Likewise. * testsuite/libgomp.c++/target-std__multimap-concurrent-usm.C: Likewise. * testsuite/libgomp.c++/target-std__multiset-concurrent-usm.C: Likewise. * testsuite/libgomp.c++/target-std__set-concurrent-usm.C: Likewise. * testsuite/libgomp.c++/target-std__span-concurrent-usm.C: Likewise. * testsuite/libgomp.c++/target-std__valarray-concurrent-usm.C: Likewise. * testsuite/libgomp.c++/target-std__vector-concurrent-usm.C: Likewise. * testsuite/libgomp.c-c++-common/target-implicit-map-4.c: Likewise. * testsuite/libgomp.c-c++-common/target-link-3.c: Likewise. * testsuite/libgomp.c-c++-common/target-link-4.c: Likewise. * testsuite/libgomp.fortran/self_maps.f90: Likewise. --- gcc/config/gcn/gcn.cc | 9 +++++--- gcc/config/gcn/mkoffload.cc | 9 +++++--- libgomp/testsuite/lib/libgomp.exp | 23 +++++++++++++++++++ .../target-std__array-concurrent-usm.C | 1 + .../target-std__bitset-concurrent-usm.C | 1 + .../target-std__deque-concurrent-usm.C | 1 + .../target-std__forward_list-concurrent-usm.C | 1 + .../target-std__list-concurrent-usm.C | 1 + .../target-std__map-concurrent-usm.C | 1 + .../target-std__multimap-concurrent-usm.C | 1 + .../target-std__multiset-concurrent-usm.C | 1 + .../target-std__set-concurrent-usm.C | 1 + .../target-std__span-concurrent-usm.C | 1 + .../target-std__valarray-concurrent-usm.C | 1 + .../target-std__vector-concurrent-usm.C | 1 + .../target-implicit-map-4.c | 1 + .../libgomp.c-c++-common/target-link-3.c | 1 + .../libgomp.c-c++-common/target-link-4.c | 1 + .../testsuite/libgomp.fortran/self_maps.f90 | 1 + 19 files changed, 51 insertions(+), 6 deletions(-) diff --git a/gcc/config/gcn/gcn.cc b/gcc/config/gcn/gcn.cc index a729ea4de36b..54abf8c1a749 100644 --- a/gcc/config/gcn/gcn.cc +++ b/gcc/config/gcn/gcn.cc @@ -2940,14 +2940,17 @@ gcn_init_cumulative_args (CUMULATIVE_ARGS *cum /* Argument info to init */ , if (!caller && cfun->machine->normal_function) gcn_detect_incoming_pointer_arg (fndecl); - if ((omp_requires_mask & (OMP_REQUIRES_UNIFIED_SHARED_MEMORY - | OMP_REQUIRES_SELF_MAPS)) + static bool warned_xnack = 0; + if (!warned_xnack + && (omp_requires_mask & (OMP_REQUIRES_UNIFIED_SHARED_MEMORY + | OMP_REQUIRES_SELF_MAPS)) && gcn_devices[gcn_arch].xnack_default != HSACO_ATTR_UNSUPPORTED && flag_xnack == HSACO_ATTR_OFF) { warning_at (UNKNOWN_LOCATION, 0, - "Unified Shared Memory is enabled, but XNACK is disabled"); + "Unified Shared Memory is required, but XNACK is disabled"); inform (UNKNOWN_LOCATION, "Try -foffload-options=-mxnack=any"); + warned_xnack = 1; } reinit_regs (); diff --git a/gcc/config/gcn/mkoffload.cc b/gcc/config/gcn/mkoffload.cc index d9d89c64f950..ac6aae52adb3 100644 --- a/gcc/config/gcn/mkoffload.cc +++ b/gcc/config/gcn/mkoffload.cc @@ -627,9 +627,12 @@ process_asm (FILE *in, FILE *out, FILE *cfile, uint32_t omp_requires) || TEST_XNACK_ON (elf_flags) || xnack_required); if (TEST_XNACK_OFF (elf_flags) && xnack_required) - fatal_error (input_location, - "conflicting settings; XNACK is forced off but Unified " - "Shared Memory is on"); + { + warning (input_location, + "conflicting settings; XNACK is forced off but Unified " + "Shared Memory is required"); + xnack_required = 0; + } /* Start generating the C code. */ if (gcn_stack_size) diff --git a/libgomp/testsuite/lib/libgomp.exp b/libgomp/testsuite/lib/libgomp.exp index 076b775560fa..cce2e93f8577 100644 --- a/libgomp/testsuite/lib/libgomp.exp +++ b/libgomp/testsuite/lib/libgomp.exp @@ -725,6 +725,29 @@ int main() { } } "-lhipblas" ] } +# return 1 if OpenMP Unified Shared Memory is supported by offload devices + +proc check_effective_target_omp_usm { } { + if { [check_effective_target_offload_device_nvptx] + || [check_effective_target_offload_target_amdgcn] } { + if [check_runtime usm_available_ { + #include + #pragma omp requires unified_shared_memory + int main () + { + int a; + #pragma omp target map(from: a) + a = omp_is_initial_device (); + return a; + } + } ] { + return 1 + } + } + + return 0 +} + # return 1 if OpenMP Device Managed Memory is supported proc check_effective_target_omp_managedmem { } { diff --git a/libgomp/testsuite/libgomp.c++/target-std__array-concurrent-usm.C b/libgomp/testsuite/libgomp.c++/target-std__array-concurrent-usm.C index 9923783bcb12..aa36f7109e93 100644 --- a/libgomp/testsuite/libgomp.c++/target-std__array-concurrent-usm.C +++ b/libgomp/testsuite/libgomp.c++/target-std__array-concurrent-usm.C @@ -1,3 +1,4 @@ +/* { dg-require-effective-target omp_usm } */ #pragma omp requires unified_shared_memory self_maps #define MEM_SHARED diff --git a/libgomp/testsuite/libgomp.c++/target-std__bitset-concurrent-usm.C b/libgomp/testsuite/libgomp.c++/target-std__bitset-concurrent-usm.C index 9023ef85c55b..d08ea7109714 100644 --- a/libgomp/testsuite/libgomp.c++/target-std__bitset-concurrent-usm.C +++ b/libgomp/testsuite/libgomp.c++/target-std__bitset-concurrent-usm.C @@ -1,3 +1,4 @@ +/* { dg-require-effective-target omp_usm } */ #pragma omp requires unified_shared_memory self_maps #define MEM_SHARED diff --git a/libgomp/testsuite/libgomp.c++/target-std__deque-concurrent-usm.C b/libgomp/testsuite/libgomp.c++/target-std__deque-concurrent-usm.C index 863a1de76874..b30ade4f0865 100644 --- a/libgomp/testsuite/libgomp.c++/target-std__deque-concurrent-usm.C +++ b/libgomp/testsuite/libgomp.c++/target-std__deque-concurrent-usm.C @@ -1,3 +1,4 @@ +/* { dg-require-effective-target omp_usm } */ #pragma omp requires unified_shared_memory self_maps #define MEM_SHARED diff --git a/libgomp/testsuite/libgomp.c++/target-std__forward_list-concurrent-usm.C b/libgomp/testsuite/libgomp.c++/target-std__forward_list-concurrent-usm.C index 60d5cee5ef36..65004b25510b 100644 --- a/libgomp/testsuite/libgomp.c++/target-std__forward_list-concurrent-usm.C +++ b/libgomp/testsuite/libgomp.c++/target-std__forward_list-concurrent-usm.C @@ -1,3 +1,4 @@ +/* { dg-require-effective-target omp_usm } */ #pragma omp requires unified_shared_memory self_maps #define MEM_SHARED diff --git a/libgomp/testsuite/libgomp.c++/target-std__list-concurrent-usm.C b/libgomp/testsuite/libgomp.c++/target-std__list-concurrent-usm.C index 5057bf96fef6..3cdd44db4277 100644 --- a/libgomp/testsuite/libgomp.c++/target-std__list-concurrent-usm.C +++ b/libgomp/testsuite/libgomp.c++/target-std__list-concurrent-usm.C @@ -1,3 +1,4 @@ +/* { dg-require-effective-target omp_usm } */ #pragma omp requires unified_shared_memory self_maps #define MEM_SHARED diff --git a/libgomp/testsuite/libgomp.c++/target-std__map-concurrent-usm.C b/libgomp/testsuite/libgomp.c++/target-std__map-concurrent-usm.C index fe37426b8c80..b7d3dd822a77 100644 --- a/libgomp/testsuite/libgomp.c++/target-std__map-concurrent-usm.C +++ b/libgomp/testsuite/libgomp.c++/target-std__map-concurrent-usm.C @@ -1,3 +1,4 @@ +/* { dg-require-effective-target omp_usm } */ #pragma omp requires unified_shared_memory self_maps #define MEM_SHARED diff --git a/libgomp/testsuite/libgomp.c++/target-std__multimap-concurrent-usm.C b/libgomp/testsuite/libgomp.c++/target-std__multimap-concurrent-usm.C index 79f9245117f4..f243790a6384 100644 --- a/libgomp/testsuite/libgomp.c++/target-std__multimap-concurrent-usm.C +++ b/libgomp/testsuite/libgomp.c++/target-std__multimap-concurrent-usm.C @@ -1,3 +1,4 @@ +/* { dg-require-effective-target omp_usm } */ #pragma omp requires unified_shared_memory self_maps #define MEM_SHARED diff --git a/libgomp/testsuite/libgomp.c++/target-std__multiset-concurrent-usm.C b/libgomp/testsuite/libgomp.c++/target-std__multiset-concurrent-usm.C index 2d8075663bd3..d869e8937ffd 100644 --- a/libgomp/testsuite/libgomp.c++/target-std__multiset-concurrent-usm.C +++ b/libgomp/testsuite/libgomp.c++/target-std__multiset-concurrent-usm.C @@ -1,3 +1,4 @@ +/* { dg-require-effective-target omp_usm } */ #pragma omp requires unified_shared_memory self_maps #define MEM_SHARED diff --git a/libgomp/testsuite/libgomp.c++/target-std__set-concurrent-usm.C b/libgomp/testsuite/libgomp.c++/target-std__set-concurrent-usm.C index 54f62e3e84ba..5fbf91b2e077 100644 --- a/libgomp/testsuite/libgomp.c++/target-std__set-concurrent-usm.C +++ b/libgomp/testsuite/libgomp.c++/target-std__set-concurrent-usm.C @@ -1,3 +1,4 @@ +/* { dg-require-effective-target omp_usm } */ #pragma omp requires unified_shared_memory self_maps #define MEM_SHARED diff --git a/libgomp/testsuite/libgomp.c++/target-std__span-concurrent-usm.C b/libgomp/testsuite/libgomp.c++/target-std__span-concurrent-usm.C index 7ef16bfb574c..09f98790b3cb 100644 --- a/libgomp/testsuite/libgomp.c++/target-std__span-concurrent-usm.C +++ b/libgomp/testsuite/libgomp.c++/target-std__span-concurrent-usm.C @@ -1,4 +1,5 @@ // { dg-additional-options "-std=c++20" } +/* { dg-require-effective-target omp_usm } */ #pragma omp requires unified_shared_memory self_maps diff --git a/libgomp/testsuite/libgomp.c++/target-std__valarray-concurrent-usm.C b/libgomp/testsuite/libgomp.c++/target-std__valarray-concurrent-usm.C index 41ec80ee900d..828b67c39306 100644 --- a/libgomp/testsuite/libgomp.c++/target-std__valarray-concurrent-usm.C +++ b/libgomp/testsuite/libgomp.c++/target-std__valarray-concurrent-usm.C @@ -1,3 +1,4 @@ +/* { dg-require-effective-target omp_usm } */ #pragma omp requires unified_shared_memory self_maps #define MEM_SHARED diff --git a/libgomp/testsuite/libgomp.c++/target-std__vector-concurrent-usm.C b/libgomp/testsuite/libgomp.c++/target-std__vector-concurrent-usm.C index 967bff3b81ad..835f6d5287eb 100644 --- a/libgomp/testsuite/libgomp.c++/target-std__vector-concurrent-usm.C +++ b/libgomp/testsuite/libgomp.c++/target-std__vector-concurrent-usm.C @@ -1,3 +1,4 @@ +/* { dg-require-effective-target omp_usm } */ #pragma omp requires unified_shared_memory self_maps #define MEM_SHARED diff --git a/libgomp/testsuite/libgomp.c-c++-common/target-implicit-map-4.c b/libgomp/testsuite/libgomp.c-c++-common/target-implicit-map-4.c index d0b0cd178c00..97bb97abccf9 100644 --- a/libgomp/testsuite/libgomp.c-c++-common/target-implicit-map-4.c +++ b/libgomp/testsuite/libgomp.c-c++-common/target-implicit-map-4.c @@ -4,6 +4,7 @@ and for not mapping the stack variables 'A' and 'B' (not mapped but accessible -> USM makes this tested feature even more important.) */ +/* { dg-require-effective-target omp_usm } */ #pragma omp requires unified_shared_memory /* Ensure that defaultmap(default : pointer) uses correct OpenMP 5.2 diff --git a/libgomp/testsuite/libgomp.c-c++-common/target-link-3.c b/libgomp/testsuite/libgomp.c-c++-common/target-link-3.c index c707b38b7d46..96642353d4ab 100644 --- a/libgomp/testsuite/libgomp.c-c++-common/target-link-3.c +++ b/libgomp/testsuite/libgomp.c-c++-common/target-link-3.c @@ -3,6 +3,7 @@ #include #include +/* { dg-require-effective-target omp_usm } */ #pragma omp requires unified_shared_memory int A[3] = {-3,-4,-5}; diff --git a/libgomp/testsuite/libgomp.c-c++-common/target-link-4.c b/libgomp/testsuite/libgomp.c-c++-common/target-link-4.c index 785055e216d7..009c521a9966 100644 --- a/libgomp/testsuite/libgomp.c-c++-common/target-link-4.c +++ b/libgomp/testsuite/libgomp.c-c++-common/target-link-4.c @@ -3,6 +3,7 @@ #include #include +/* { dg-require-effective-target omp_usm } */ #pragma omp requires self_maps int A[3] = {-3,-4,-5}; diff --git a/libgomp/testsuite/libgomp.fortran/self_maps.f90 b/libgomp/testsuite/libgomp.fortran/self_maps.f90 index 208fd1c71d5c..60889687c2cd 100644 --- a/libgomp/testsuite/libgomp.fortran/self_maps.f90 +++ b/libgomp/testsuite/libgomp.fortran/self_maps.f90 @@ -1,4 +1,5 @@ ! Basic test whether self_maps work +! { dg-require-effective-target omp_usm } module m !$omp requires self_maps -- 2.47.3