From mboxrd@z Thu Jan 1 00:00:00 1970 Return-Path: Received: by sourceware.org (Postfix, from userid 1534) id 44623385C317; Mon, 12 Sep 2022 13:50:30 +0000 (GMT) DKIM-Filter: OpenDKIM Filter v2.11.0 sourceware.org 44623385C317 DKIM-Signature: v=1; a=rsa-sha256; c=relaxed/relaxed; d=gcc.gnu.org; s=default; t=1662990630; bh=RW9dSOa93r7VtR8pu+ObztrJvtsw88Q9oUQfkZHLS78=; h=From:To:Subject:Date:From; b=YWwWY34B6GZlE9MFNWwPcGWdF3kQZCWq83V57xn5RX/Mo+9jtVxQQmtFfQ2Y+mUbc hsmxl2yMIb3oouMp72WtZ/TW+KF4gey5bdWKad+4V6sginiznV024SO2pQs4+sGu9Z Az0+NccXq0dH8JeRyKGwJktr8XHvhm6VZ6msKs5c= Content-Type: text/plain; charset="us-ascii" MIME-Version: 1.0 Content-Transfer-Encoding: 7bit From: Tobias Burnus To: gcc-cvs@gcc.gnu.org Subject: [gcc/devel/omp/gcc-12] 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/devel/omp/gcc-12 X-Git-Oldrev: 8a04333f11401685a8c06cf10afcde70c5e4ec6f X-Git-Newrev: c8d2d18fc5d03590f6d7e452084b16cce5e044e4 Message-Id: <20220912135030.44623385C317@sourceware.org> Date: Mon, 12 Sep 2022 13:50:30 +0000 (GMT) List-Id: https://gcc.gnu.org/g:c8d2d18fc5d03590f6d7e452084b16cce5e044e4 commit c8d2d18fc5d03590f6d7e452084b16cce5e044e4 Author: Tobias Burnus Date: Mon Sep 12 15:47:20 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. (cherry picked from commit 6b43f556f392a7165582aca36a19fe7389d995b2) Diff: --- gcc/ChangeLog.omp | 11 +++++++++++ gcc/config/gcn/mkoffload.cc | 2 +- gcc/config/nvptx/mkoffload.cc | 18 ++++++++++++++---- libgomp/ChangeLog.omp | 15 +++++++++++++++ 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 + 11 files changed, 65 insertions(+), 5 deletions(-) diff --git a/gcc/ChangeLog.omp b/gcc/ChangeLog.omp index a2d1222bccf..4f80bcbd356 100644 --- a/gcc/ChangeLog.omp +++ b/gcc/ChangeLog.omp @@ -1,3 +1,14 @@ +2022-09-12 Tobias Burnus + + Backport from mainline: + 2022-09-12 Tobias Burnus + + * 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. + 2022-09-12 Tobias Burnus Backport from mainline: diff --git a/gcc/config/gcn/mkoffload.cc b/gcc/config/gcn/mkoffload.cc index 2c7a1feafd2..14d7eedc790 100644 --- a/gcc/config/gcn/mkoffload.cc +++ b/gcc/config/gcn/mkoffload.cc @@ -844,7 +844,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/ChangeLog.omp b/libgomp/ChangeLog.omp index 44913b6e2ac..c9827ba53c2 100644 --- a/libgomp/ChangeLog.omp +++ b/libgomp/ChangeLog.omp @@ -1,3 +1,18 @@ +2022-09-12 Tobias Burnus + + Backport from mainline: + 2022-09-12 Tobias Burnus + + * 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. + 2022-09-12 Tobias Burnus Backport from mainline: diff --git a/libgomp/libgomp.texi b/libgomp/libgomp.texi index c342aa6f5cc..0069b34f445 100644 --- a/libgomp/libgomp.texi +++ b/libgomp/libgomp.texi @@ -4381,6 +4381,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: