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, - "% requires at least " - "% for %<-misa=%>"); + { + warning_at (input_location, 0, + "% requires at " + "least % 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: