From mboxrd@z Thu Jan 1 00:00:00 1970 Return-Path: Received: by sourceware.org (Postfix, from userid 1534) id 373CF385702A; Mon, 12 Sep 2022 13:28:52 +0000 (GMT) DKIM-Filter: OpenDKIM Filter v2.11.0 sourceware.org 373CF385702A DKIM-Signature: v=1; a=rsa-sha256; c=relaxed/relaxed; d=gcc.gnu.org; s=default; t=1662989332; bh=QggGY640TB9iWj4Jw+WEKkM001Z+rNywtLdWgi5l9Q0=; h=From:To:Subject:Date:From; b=iGO3YIdmjIWZjQwAj+rQyX6Gq6QNndniw41tTbLVTZhes+h80VwmiOV/0rYdhusa9 7IFity7v7G8Lu/TTCvtvxkMicUTneBGu6uUBQIbPQXUmD5wDEcnbCMqpzeJb1Hmf8e Fy7QEPxn71RgDnd6Xdyo+/2uiHuDBi5bYu3cz6fk= MIME-Version: 1.0 Content-Transfer-Encoding: 7bit Content-Type: text/plain; charset="utf-8" From: Tobias Burnus To: gcc-cvs@gcc.gnu.org Subject: [gcc r13-2625] nvptx/mkoffload.cc: Warn instead of error when reverse offload is not possible X-Act-Checkin: gcc X-Git-Author: Tobias Burnus X-Git-Refname: refs/heads/master X-Git-Oldrev: 06b30eecdd9822842a0ff21548ab92f01adb2795 X-Git-Newrev: 6b43f556f392a7165582aca36a19fe7389d995b2 Message-Id: <20220912132852.373CF385702A@sourceware.org> Date: Mon, 12 Sep 2022 13:28:52 +0000 (GMT) List-Id: https://gcc.gnu.org/g:6b43f556f392a7165582aca36a19fe7389d995b2 commit r13-2625-g6b43f556f392a7165582aca36a19fe7389d995b2 Author: Tobias Burnus Date: Mon Sep 12 15:25:13 2022 +0200 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: --- gcc/config/gcn/mkoffload.cc | 2 +- gcc/config/nvptx/mkoffload.cc | 18 ++++++++++++++---- 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 +++++++++++++++ .../testsuite/libgomp.fortran/reverse-offload-1.f90 | 1 + 9 files changed, 39 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..854cd72f3c7 100644 --- 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, - "% requires at least " - "% for %<-misa=%>"); + { + warning_at (input_location, 0, + "% requires at " + "least % 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 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..fbfeae1fd41 --- /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 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: