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.

diff --git a/gcc/config/gcn/gcn.cc b/gcc/config/gcn/gcn.cc
index a729ea4de36b00aa05d4134eda56a8b614024903..54abf8c1a749a188704ed71489649ebddf94f774 100644 (file)
--- 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 d9d89c64f950fb1b2a31077bb8c30efd4f8915ab..ac6aae52adb378c02552db4ee329f2b50d2b7db1 100644 (file)
--- 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 076b775560fa1c5aa4fe3dedab0974ead3d46437..cce2e93f8577cc1d62b63646ec7db72538a18d3b 100644 (file)
--- 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 <omp.h>
+           #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 9923783bcb12dd5ee8803ea2facca8ba9e0759c1..aa36f7109e93a282a4a159ddb79a1d5bb383d67d 100644 (file)
--- 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 9023ef85c55b3dce84750c6d5abe5509292924cf..d08ea710971444502da523e3ddf076b3b8a660ef 100644 (file)
--- 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 863a1de76874417b743f3a3e062937c6f3fbe5ed..b30ade4f08652199530fb9a2a53f555bfa446fc8 100644 (file)
--- 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 60d5cee5ef36f7abeee903abd3ef40b1fdee0b53..65004b25510b4c8d7711f28e639f0647b5841e21 100644 (file)
--- 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 5057bf96fef6c3dcc2dfaff04bc4b85d2cc2e168..3cdd44db427725180e5017a0df6e1485a65d3564 100644 (file)
--- 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 fe37426b8c8026ae284f3d71bacff7c321522ab4..b7d3dd822a77d5e4a9434146a219351a516b6cc7 100644 (file)
--- 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 79f9245117f41e37f0d06e3ec20608298347511c..f243790a6384e438efa1a15ecf64ce7cc6473e7b 100644 (file)
--- 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 2d8075663bd3a98e65937c11ba634771004372b6..d869e8937ffd146ca93742966ba0101be3df4b56 100644 (file)
--- 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 54f62e3e84ba258934d59b01e35832c0c425dee1..5fbf91b2e077d8e3338f6bb704ffb4ddbdfd38c7 100644 (file)
--- 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 7ef16bfb574c75d95afbed562cd943e11c48f080..09f98790b3cba39068f43e89d3e62492ba7f9521 100644 (file)
--- 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 41ec80ee900d2ca73288224ea48fac943ce7a3d2..828b67c39306dafe4a48bd2c9b94efaddfbe3e0a 100644 (file)
--- 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 967bff3b81ad909f4ec1309d8d4d167ee3253bb5..835f6d5287eb09fe5af68faf8ab73ab04358ca7b 100644 (file)
--- 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 d0b0cd178c0068d85a5d904520d2795795a5de00..97bb97abccf9aa63ef10bd7a2c71dfa04d33b745 100644 (file)
--- 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 c707b38b7d46a8ef22499342e08cb7d9e251f7fa..96642353d4ab7797d6b56e4c423a7f86c267a10e 100644 (file)
--- 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 <stdint.h>
 #include <omp.h>
 
+/* { 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 785055e216d714d9b5e92242308a321ee86589d9..009c521a9966b7f8dd71d8bbe3484dc8a160dc0f 100644 (file)
--- 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 <stdint.h>
 #include <omp.h>
 
+/* { 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 208fd1c71d5c64b7abc4ce0baceb074e0649417c..60889687c2cd030f97818cc9bea1296d02d86009 100644 (file)
--- 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