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(-)
@@ -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);
@@ -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)
@@ -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
@@ -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.
@@ -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
@@ -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:
@@ -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:
new file mode 100644
@@ -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 } */
@@ -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:
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