nvptx/mkoffload.cc: Warn instead of error when reverse offload is not possible

Reverse offload requests at least -misa=sm_35; with this patch, a warning
instead of an error is shown, still permitting reverse offload for all
other configured device types. This is achieved by not calling
GOMP_offload_register_ver (and stopping generating pointless 'static const char'
variables, once known.)

The tool_name as progname changes adds "nvptx " and "gcn " to the
"mkoffload: warning/error:" diagnostic.

gcc/ChangeLog:

	* config/nvptx/mkoffload.cc (process): Replace a fatal_error by
	a warning + not enabling offloading if -misa=sm_30 prevents
	reverse offload.
	(main): Use tool_name as progname for diagnostic.
	* config/gcn/mkoffload.cc (main): Likewise.

libgomp/ChangeLog:

	* libgomp.texi (Offload-Target Specifics: nvptx): Document
	that reverse offload requires >= -march=sm_35.
	* testsuite/libgomp.c-c++-common/requires-4.c: Build for nvptx
	with -misa=sm_35.
	* testsuite/libgomp.c-c++-common/requires-5.c: Likewise.
	* testsuite/libgomp.c-c++-common/requires-6.c: Likewise.
	* testsuite/libgomp.c-c++-common/reverse-offload-1.c: Likewise.
	* testsuite/libgomp.fortran/reverse-offload-1.f90: Likewise.
	* testsuite/libgomp.c/reverse-offload-sm30.c: New test.

gcc/config/gcn/mkoffload.cc

@@ -805,7 +805,7 @@ main (int argc, char **argv)
   FILE *cfile = stdout;
   const char *outname = 0;
 
-  progname = "mkoffload";
+  progname = tool_name;
   diagnostic_initialize (global_dc, 0);
 
   obstack_init (&files_to_cleanup);

gcc/config/nvptx/mkoffload.cc

@@ -324,9 +324,19 @@ process (FILE *in, FILE *out, uint32_t omp_requires)
 	{
 	  if (sm_ver && sm_ver[0] == '3' && sm_ver[1] == '0'
 	      && sm_ver[2] == '\n')
-	    fatal_error (input_location,
-			 "%<omp requires reverse_offload%> requires at least "
-			 "%<sm_35%> for %<-misa=%>");
+	    {
+	      warning_at (input_location, 0,
+			  "%<omp requires reverse_offload%> requires at "
+			  "least %<sm_35%> for "
+			  "%<-foffload-options=nvptx-none=-march=%> - disabling"
+			  " offload-code generation for this device type");
+	      /* As now an empty file is compiled and there is no call to
+		 GOMP_offload_register_ver, this device type is effectively
+		 disabled.  */
+	      fflush (out);
+	      ftruncate (fileno (out), 0);
+	      return;
+	    }
 	  sm_ver2 = sm_ver;
 	  version2 = version;
 	}
@@ -526,7 +536,7 @@ main (int argc, char **argv)
   FILE *out = stdout;
   const char *outname = 0;
 
-  progname = "mkoffload";
+  progname = tool_name;
   diagnostic_initialize (global_dc, 0);
 
   if (atexit (mkoffload_cleanup) != 0)

libgomp/libgomp.texi

@@ -4386,6 +4386,9 @@ The implementation remark:
 @item I/O within OpenMP target regions and OpenACC parallel/kernels is supported
       using the C library @code{printf} functions and the Fortran
       @code{print}/@code{write} statements.
+@item Compilation OpenMP code that contains @code{requires reverse_offload}
+      requires at least @code{-march=sm_35}, compiling for @code{-march=sm_30}
+      is not supported.
 @end itemize
 
 

libgomp/testsuite/libgomp.c-c++-common/requires-4.c

@@ -1,4 +1,5 @@
 /* { dg-additional-options "-flto" } */
+/* { dg-additional-options "-foffload-options=nvptx-none=-misa=sm_35" { target { offload_target_nvptx } } } */
 /* { dg-additional-sources requires-4-aux.c } */
 
 /* Check no diagnostic by device-compiler's or host compiler's lto1.

libgomp/testsuite/libgomp.c-c++-common/requires-5.c

@@ -1,3 +1,4 @@
+/* { dg-additional-options "-foffload-options=nvptx-none=-misa=sm_35" { target { offload_target_nvptx } } } */
 /* { dg-additional-sources requires-5-aux.c } */
 
 /* Depending on offload device capabilities, it may print something like the

libgomp/testsuite/libgomp.c-c++-common/requires-6.c

@@ -1,3 +1,5 @@
+/* { dg-additional-options "-foffload-options=nvptx-none=-misa=sm_35" { target { offload_target_nvptx } } } */
+
 #pragma omp requires unified_shared_memory, unified_address, reverse_offload
 
 /* The requires line is not active as there is none of:

libgomp/testsuite/libgomp.c-c++-common/reverse-offload-1.c

@@ -1,4 +1,5 @@
 /* { dg-do run }  */
+/* { dg-additional-options "-foffload-options=nvptx-none=-misa=sm_35" { target { offload_target_nvptx } } } */
 /* { dg-additional-sources reverse-offload-1-aux.c } */
 
 /* Check that reverse offload works in particular:

libgomp/testsuite/libgomp.c/reverse-offload-sm30.c

@@ -0,0 +1,15 @@
+/* { dg-do link { target { offload_target_nvptx } } } */
+/* { dg-additional-options "-foffload-options=nvptx-none=-march=sm_30 -foffload=-mptx=_" } */
+
+#pragma omp requires reverse_offload
+
+int
+main ()
+{
+  #pragma omp target
+    {
+    }
+  return 0;
+}
+
+/* { dg-warning "'omp requires reverse_offload' requires at least 'sm_35' for '-foffload-options=nvptx-none=-march=' - disabling offload-code generation for this device type" "" { target *-*-* } 0 } */

libgomp/testsuite/libgomp.fortran/reverse-offload-1.f90

@@ -1,4 +1,5 @@
 ! { dg-do run }
+! { dg-additional-options "-foffload-options=nvptx-none=-misa=sm_35" { target { offload_target_nvptx } } }
 ! { dg-additional-sources reverse-offload-1-aux.f90 }
 
 ! Check that reverse offload works in particular: