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

@@ -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 ();

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)

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 { } {

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

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

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

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

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

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

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

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

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

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
 

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

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

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

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};

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};

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