nvptx/mkoffload.cc: Warn instead of error when reverse offload is not possible (was: Re: [Patch][2/3][v2] nvptx: libgomp+mkoffload.cc: Prepare for reverse offload fn lookup)

Message ID cd13de52-a268-8a3d-ea84-6be9d9599cbb@codesourcery.com
State New, archived
Headers
Series nvptx/mkoffload.cc: Warn instead of error when reverse offload is not possible (was: Re: [Patch][2/3][v2] nvptx: libgomp+mkoffload.cc: Prepare for reverse offload fn lookup) |

Commit Message

Tobias Burnus Sept. 12, 2022, 12:02 p.m. UTC
  On 09.09.22 17:36, Jakub Jelinek wrote:

Wonder if we instead shouldn't arrange for silent request for no PTX
offloading (or one with warning?) if sm_30 and reverse offload is needed.
Error might be too harsh, the program can still offload to GCN or host
just fine...


Attached patch now implements the warning. I think silently failing is
not the proper solution. It is too confusing and without telling the
user, they may not notice this issue.


I also changed the "progname" variable; it seems as if the only user is
the diagnostic machinery and it seems to make sense to have:

nvptx mkoffload: warning: 'omp requires reverse_offload' requires at least
'sm_35' for '-misa=' - disabling offload-code generation for this device type

instead of guessing whether "mkoffload" is for the host, gcn or nvptx. I know
that the common way is to use the binary name ("lto1") or ... but I still
think this prefix – which is the tool_name. Makes sense.

(BTW: gcc/config/i386/intelmic-mkoffload.cc uses "mkoffload-intelmic".)

I also changed "-misa=" in the diagnostic to "-march=" as Tom changed the default
from "-misa=" (now an alias) to "-march=" in GCC 12 (+plus added -march-map= in addition).

I also added a testcase for this – and updated the testsuite for the dg-warning.

OK for mainline?

Tobias

PS: As no code for nvptx is generated, there is no "device present bit not used" warning
with GOMP_DEBUG, but as there is a compile-time warning, I guess that's fine.


-----------------
Siemens Electronic Design Automation GmbH; Anschrift: Arnulfstraße 201, 80634 München; Gesellschaft mit beschränkter Haftung; Geschäftsführer: Thomas Heurung, Frank Thürauf; Sitz der Gesellschaft: München; Registergericht München, HRB 106955
  

Comments

Jakub Jelinek Sept. 12, 2022, 12:10 p.m. UTC | #1
On Mon, Sep 12, 2022 at 02:02:16PM +0200, Tobias Burnus wrote:
> +	    {
> +	      warning_at (input_location, 0,
> +			  "%<omp requires reverse_offload%> requires at "
> +			  "least %<sm_35%> for %<-march=%> - disabling "
> +			  "offload-code generation for this device type");

I wonder whether it shouldn't talk about
-foffload-options=nvptx-none=-march=
instead of just -march=.

Otherwise LGTM.

	Jakub
  

Patch

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                             |  2 +-
 gcc/config/nvptx/mkoffload.cc                           | 17 +++++++++++++----
 libgomp/libgomp.texi                                    |  3 +++
 libgomp/testsuite/libgomp.c-c++-common/requires-4.c     |  1 +
 libgomp/testsuite/libgomp.c-c++-common/requires-5.c     |  1 +
 libgomp/testsuite/libgomp.c-c++-common/requires-6.c     |  2 ++
 .../testsuite/libgomp.c-c++-common/reverse-offload-1.c  |  1 +
 libgomp/testsuite/libgomp.c/reverse-offload-sm30.c      | 15 +++++++++++++++
 libgomp/testsuite/libgomp.fortran/reverse-offload-1.f90 |  1 +
 9 files changed, 38 insertions(+), 5 deletions(-)

diff --git a/gcc/config/gcn/mkoffload.cc b/gcc/config/gcn/mkoffload.cc
index 24d327355e3..64037806acf 100644
--- 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 834b2059aac..5f3e07ad066 100644
--- a/gcc/config/nvptx/mkoffload.cc
+++ b/gcc/config/nvptx/mkoffload.cc
@@ -324,9 +324,18 @@  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 %<-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 +535,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 4eaad4348bb..1f402d6df79 100644
--- 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 6ed5a5f647a..5883eff0d93 100644
--- 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 7fe0c735d27..d43d78db6fa 100644
--- 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 b00c7459bbc..a25b4d2dedd 100644
--- 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 976e129f560..52d828caf1c 100644
--- 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
index 00000000000..14aed0132b7
--- /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 '-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 7cfb8b6552e..de68011f8f7 100644
--- 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: