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.

diff --git a/gcc/config/gcn/mkoffload.cc b/gcc/config/gcn/mkoffload.cc
index 24d3273..6403780 100644 (file)
--- a/gcc/config/gcn/mkoffload.cc
+++ b/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);

diff --git a/gcc/config/nvptx/mkoffload.cc b/gcc/config/nvptx/mkoffload.cc
index 834b205..854cd72 100644 (file)
--- a/gcc/config/nvptx/mkoffload.cc
+++ b/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)

diff --git a/libgomp/libgomp.texi b/libgomp/libgomp.texi
index 4eaad43..1f402d6 100644 (file)
--- a/libgomp/libgomp.texi
+++ b/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
 
 

diff --git a/libgomp/testsuite/libgomp.c-c++-common/requires-4.c b/libgomp/testsuite/libgomp.c-c++-common/requires-4.c
index 6ed5a5f..5883eff 100644 (file)
--- a/libgomp/testsuite/libgomp.c-c++-common/requires-4.c
+++ b/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.

diff --git a/libgomp/testsuite/libgomp.c-c++-common/requires-5.c b/libgomp/testsuite/libgomp.c-c++-common/requires-5.c
index 7fe0c73..d43d78d 100644 (file)
--- a/libgomp/testsuite/libgomp.c-c++-common/requires-5.c
+++ b/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

diff --git a/libgomp/testsuite/libgomp.c-c++-common/requires-6.c b/libgomp/testsuite/libgomp.c-c++-common/requires-6.c
index b00c745..a25b4d2 100644 (file)
--- a/libgomp/testsuite/libgomp.c-c++-common/requires-6.c
+++ b/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:

diff --git a/libgomp/testsuite/libgomp.c-c++-common/reverse-offload-1.c b/libgomp/testsuite/libgomp.c-c++-common/reverse-offload-1.c
index 976e129..52d828c 100644 (file)
--- a/libgomp/testsuite/libgomp.c-c++-common/reverse-offload-1.c
+++ b/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:

diff --git a/libgomp/testsuite/libgomp.c/reverse-offload-sm30.c b/libgomp/testsuite/libgomp.c/reverse-offload-sm30.c
new file mode 100644 (file)
index 0000000..fbfeae1
--- /dev/null
+++ b/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 } */

diff --git a/libgomp/testsuite/libgomp.fortran/reverse-offload-1.f90 b/libgomp/testsuite/libgomp.fortran/reverse-offload-1.f90
index 7cfb8b6..de68011 100644 (file)
--- a/libgomp/testsuite/libgomp.fortran/reverse-offload-1.f90
+++ b/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: