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.

19 files changed, 51 insertions, 6 deletions

diff --git a/gcc/config/gcn/gcn.cc b/gcc/config/gcn/gcn.cc
index a729ea4..54abf8c 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 d9d89c6..ac6aae5 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 076b775..cce2e93 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 <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 9923783..aa36f71 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 9023ef8..d08ea71 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 863a1de..b30ade4 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 60d5cee..65004b2 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 5057bf9..3cdd44d 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 fe37426..b7d3dd8 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 79f9245..f243790 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 2d80756..d869e89 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 54f62e3..5fbf91b2 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 7ef16bf..09f9879 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 41ec80e..828b67c 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 967bff3..835f6d5 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 d0b0cd1..97bb97a 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 c707b38..9664235 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 <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 785055e..009c521 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 <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 208fd1c..6088968 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