[Patch] nvptx: Add -mptx=6.0 + -misa=sm_70

[Patch] nvptx: Add -mptx=6.0 + -misa=sm_70

[Tobias Burnus]

[-- Attachment #1: Type: text/plain, Size: 4161 bytes --]

This patch exposes two -m* option values which are already
internally available. I think it makes sense to expose them
explicitly to the user (see below), but there are also arguments
against. Thoughts?


PTX version (-mptx=)
[patch adds -mptx=6.0 as option]

* Currently supported internally are 3.1 (CUDA 5.0, used by GCC <= 11),
   6.0 (CUDA 9.0, current GCC 12 default), 6.3 (CUDA 10.0), 7.0 (CUDA 11.0)
* -mptx= supports 3.1, 6.3, 7.0 – but not the internal default 6.0

First, I think all versions make sense:
* 3.1 is the previous default and permits running with older CUDA (if need)
* 6.0 is for CUDA 9 - and if we want to support it, it has to stay.
   6.0 is the default since commit
   https://gcc.gnu.org/pipermail/gcc-patches/2022-February/590007.html
* 6.3 is CUDA 10.0. In that PTX version, a lot of nice features
   were added like .alias
* 7.0 is CUDA 11.0. This adds support for sm_80 (honored in code gen).

PTX >= 6.0 makes sense as it permits newer sm_* (in particular: sm_53 and sm_70)
and
+  /* Pick at least 6.0, to enable using bar.warp.sync to have a way to force
+     warp convergence.  */
On the other hand, for older systems, CUDA 10.0 might be too new and we still
want to support CUDA 9. (At least that's how I understood one of nvpx gcc
emails, which I cannot find at the moment.)

Assuming we don't want to change the default minimal version from PTX 6.0
back to 6.3, it looks as both should stay.
Downside: we probably need one lib{c,gomp,gfortran,...} per PTX version,
i.e. 4 versions (3.1, 6.0, 6.3, 7.0).

I think it makes sense to expose the 6.0 value to the user and not
only use it internally behind the scenes. As it is already used internally,
the change is tiny but user visible. Thus, it has to stay when we will
bump the default in later GCC versions; on the other hand, if we bump
the default, it might be also a good reason to have it to permit the
user to have a backward compatible PTX output for linking libraries.

  * * *

SM version (-misa=)
[Patch adds -misa=sm_70]

* The compiler supports internally: SM_30, SM_35, SM_53, SM_70, SM_75, SM_80.
* GCC <= 11 only had sm_30 and sm_35 (supported since PTX 3.1/CUDA 5.0)
* GCC 12 exposes
   - sm_30, sm_35,
   - sm_53 (PTX 4.2, CUDA 7.0),
   - sm_75 (PTX 6.3, CUDA 10.0)
   - sm_80 (PTX 7.0, CUDA 11.0)
   but it does not permit using -misa=sm_70 (PTX 6.0, CUDA 9.0).
* Note: sm_75 + sm_80 imply a newer PTX version, which
   the compiler defaults to (if no -mptx= has been specified).

I think it makes sense to have sm_70 in addition:
* sm_70 enables several new features (see PTX documentation)
* sm_70 is the highest supported for CUDA 9 (default PTX version);
   as sm_75 will require CUDA 10, currently only sm_53 can be used with CUDA 9.
* The current code actually does generate different code for >= sm_70
   already.

  * * *

This patch updates -misa= and -mptx= documentation to match what actually has
been implemented. I think that makes sense as:
* The currently documented default for -mptx= is no longer true.
* The available values are already exposed via the diagnostic
* The multilib issue already occurs when the user explicitly specifies -mptx=6.3
   (or -mptx=3.1).
* If needed, we could note that certain PTX or ISA values are experimental.

I think besides > sm_35 being experimental, there is no reason that higher sm_*
should not be used. Except for the pre-existing multilib issue and for the ICE
when bootstrapping with sm_53 (instead of sm_35) as default ISA version.
But that's solved by Roger's patch (pending ME (and then BE) review),
https://gcc.gnu.org/pipermail/gcc-patches/2022-February/590545.html

* * *

Comments to any of those three patches (-mptx=6.0, -misa=sm_70, documentation)?
(Lightly tested on x86-64 with nvptx offloading.)
OK? (All, some?)

Tobias
-----------------
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

[-- Attachment #2: sm70-ptx-6.diff --]
[-- Type: text/x-patch, Size: 4015 bytes --]

nvptx: Add -mptx=6.0 + -misa=sm_70

gcc/ChangeLog:

	* config/nvptx/nvptx-c.cc (nvptx_cpu_cpp_builtins): Handle SM70.
	* gcc/config/nvptx/nvptx.cc (first_ptx_version_supporting_sm):
	Likewise.
	* config/nvptx/nvptx.opt (misa): Add sm_70 alias PTX_ISA_SM70.
	(mptx): Add 6.0 alias PTX_VERSION_6_0.
	* config/nvptx/t-omp-device: Add sm_53, sm_70, sm_75, sm_80.
	* doc/invoke.texi (-misa, -mptx): Update for new values and
	defaults.

 gcc/config/nvptx/nvptx-c.cc   |  2 ++
 gcc/config/nvptx/nvptx.cc     |  2 ++
 gcc/config/nvptx/nvptx.opt    |  6 ++++++
 gcc/config/nvptx/t-omp-device |  2 +-
 gcc/doc/invoke.texi           | 17 +++++++++++------
 5 files changed, 22 insertions(+), 7 deletions(-)

diff --git a/gcc/config/nvptx/nvptx-c.cc b/gcc/config/nvptx/nvptx-c.cc
index d68b9910d7e..b2375fb5b16 100644
--- a/gcc/config/nvptx/nvptx-c.cc
+++ b/gcc/config/nvptx/nvptx-c.cc
@@ -43,6 +43,8 @@ nvptx_cpu_cpp_builtins (void)
     cpp_define (parse_in, "__PTX_SM__=800");
   else if (TARGET_SM75)
     cpp_define (parse_in, "__PTX_SM__=750");
+  else if (TARGET_SM70)
+    cpp_define (parse_in, "__PTX_SM__=700");
   else if (TARGET_SM53)
     cpp_define (parse_in, "__PTX_SM__=530");
   else if (TARGET_SM35)
diff --git a/gcc/config/nvptx/nvptx.cc b/gcc/config/nvptx/nvptx.cc
index afbad5bdde6..f3f3201a7ba 100644
--- a/gcc/config/nvptx/nvptx.cc
+++ b/gcc/config/nvptx/nvptx.cc
@@ -217,6 +217,8 @@ first_ptx_version_supporting_sm (enum ptx_isa sm)
       return PTX_VERSION_3_1;
     case PTX_ISA_SM53:
       return PTX_VERSION_4_2;
+    case PTX_ISA_SM70:
+      return PTX_VERSION_6_0;
     case PTX_ISA_SM75:
       return PTX_VERSION_6_3;
     case PTX_ISA_SM80:
diff --git a/gcc/config/nvptx/nvptx.opt b/gcc/config/nvptx/nvptx.opt
index e3f65b2d0b1..c00af111829 100644
--- a/gcc/config/nvptx/nvptx.opt
+++ b/gcc/config/nvptx/nvptx.opt
@@ -64,6 +64,9 @@ Enum(ptx_isa) String(sm_35) Value(PTX_ISA_SM35)
 EnumValue
 Enum(ptx_isa) String(sm_53) Value(PTX_ISA_SM53)
 
+EnumValue
+Enum(ptx_isa) String(sm_70) Value(PTX_ISA_SM70)
+
 EnumValue
 Enum(ptx_isa) String(sm_75) Value(PTX_ISA_SM75)
 
@@ -82,6 +85,9 @@ Known PTX versions (for use with the -mptx= option):
 EnumValue
 Enum(ptx_version) String(3.1) Value(PTX_VERSION_3_1)
 
+EnumValue
+Enum(ptx_version) String(6.0) Value(PTX_VERSION_6_0)
+
 EnumValue
 Enum(ptx_version) String(6.3) Value(PTX_VERSION_6_3)
 
diff --git a/gcc/config/nvptx/t-omp-device b/gcc/config/nvptx/t-omp-device
index 8765d9f1881..4228218a424 100644
--- a/gcc/config/nvptx/t-omp-device
+++ b/gcc/config/nvptx/t-omp-device
@@ -1,4 +1,4 @@
 omp-device-properties-nvptx: $(srcdir)/config/nvptx/nvptx.cc
 	echo kind: gpu > $@
 	echo arch: nvptx >> $@
-	echo isa: sm_30 sm_35 >> $@
+	echo isa: sm_30 sm_35 sm_53 sm_70 sm_75 sm_80 >> $@
diff --git a/gcc/doc/invoke.texi b/gcc/doc/invoke.texi
index e1a00c80307..54ca0070356 100644
--- a/gcc/doc/invoke.texi
+++ b/gcc/doc/invoke.texi
@@ -27250,15 +27250,20 @@ supported.
 
 @item -misa=@var{ISA-string}
 @opindex march
-Generate code for given the specified PTX ISA (e.g.@: @samp{sm_35}).  ISA
-strings must be lower-case.  Valid ISA strings include @samp{sm_30} and
-@samp{sm_35}.  The default ISA is sm_35.
+Generate code for given the specified PTX ISA (e.g.@: @samp{sm_70}).  ISA
+strings must be lower-case.  Valid ISA strings include @samp{sm_30},
+@samp{sm_35}, @samp{sm_53}, @samp{sm_70}, @samp{sm_75}, and @samp{sm_80}.
+The default ISA is sm_35.
 
 @item -mptx=@var{version-string}
 @opindex mptx
-Generate code for given the specified PTX version (e.g.@: @samp{6.3}).
-Valid version strings include @samp{3.1} and @samp{6.3}.  The default PTX
-version is 3.1.
+Generate code for given the specified PTX version (e.g.@: @samp{7.0}).
+Valid version strings include @samp{3.1}, @samp{6.0}, @samp{6.3}, and
+@samp{7.0}.  The default PTX version is 6.0, unless a higher minimal
+version is required for specified PTX ISA via option @option{-misa=}.
+
+the lowest supported
+PTX version supporting
 
 @item -mmainkernel
 @opindex mmainkernel

Re: [Patch] nvptx: Add -mptx=6.0 + -misa=sm_70

[Tom de Vries]

[-- Attachment #1: Type: text/plain, Size: 1583 bytes --]

On 2/17/22 18:24, Tobias Burnus wrote:
> PTX version (-mptx=)
> [patch adds -mptx=6.0 as option]
> 
> * Currently supported internally are 3.1 (CUDA 5.0, used by GCC <= 11),
>    6.0 (CUDA 9.0, current GCC 12 default), 6.3 (CUDA 10.0), 7.0 (CUDA 11.0)
> * -mptx= supports 3.1, 6.3, 7.0 – but not the internal default 6.0
> 

I tend not to think in terms of CUDA versions, but supported driver 
versions.

In the end, drivers are used to translate ptx to SASS for execution, 
CUDA is just used for build time verification (or not, if it's not in 
the path).

And a driver may or may not be supported.  F.i. 390.x still may receive 
updates from nvidia, but there are JIT bugs that we've reported that 
they've decided not to fix, so from that point of view 390.x is unsupported.

> I think it makes sense to expose the 6.0 value to the user and not
> only use it internally behind the scenes. As it is already used internally,
> the change is tiny but user visible. 

Sure, I've committed this (with a somewhat shorter commit log).

>Thus, it has to stay when we will
> bump the default in later GCC versions; on the other hand, if we bump
> the default, it might be also a good reason to have it to permit the
> user to have a backward compatible PTX output for linking libraries.
> 

FWIW, I think that it's possible to link different versions of ptx isa 
together (though perhaps there are specific scenarios where that's not 
possible, I'm not sure).  But mixing versions restricts the range of 
drivers you can use, so it may make sense to just use one version.

Thanks,
- Tom

[-- Attachment #2: 0001-nvptx-Add-mptx-6.0.patch --]
[-- Type: text/x-patch, Size: 1765 bytes --]

nvptx: Add -mptx=6.0

Currently supported internally are 3.1, 6.0, 6.3 and 7.0.

However, -mptx= supports 3.1, 6.3, 7.0 – but not the internal default 6.0.

Add -mptx=6.0 for consistency.

Tested on nvptx.

gcc/ChangeLog:

	* config/nvptx/nvptx.opt (mptx): Add 6.0 alias PTX_VERSION_6_0.
	* doc/invoke.texi (-mptx): Update for new values and defaults.

Co-Authored-By: Tom de Vries <tdevries@suse.de>

---
 gcc/config/nvptx/nvptx.opt | 3 +++
 gcc/doc/invoke.texi        | 7 ++++---
 2 files changed, 7 insertions(+), 3 deletions(-)

diff --git a/gcc/config/nvptx/nvptx.opt b/gcc/config/nvptx/nvptx.opt
index e56ec9288da..97e127cc4fb 100644
--- a/gcc/config/nvptx/nvptx.opt
+++ b/gcc/config/nvptx/nvptx.opt
@@ -82,6 +82,9 @@ Known PTX versions (for use with the -mptx= option):
 EnumValue
 Enum(ptx_version) String(3.1) Value(PTX_VERSION_3_1)
 
+EnumValue
+Enum(ptx_version) String(6.0) Value(PTX_VERSION_6_0)
+
 EnumValue
 Enum(ptx_version) String(6.3) Value(PTX_VERSION_6_3)
 
diff --git a/gcc/doc/invoke.texi b/gcc/doc/invoke.texi
index 635c5f79278..56f3a01de44 100644
--- a/gcc/doc/invoke.texi
+++ b/gcc/doc/invoke.texi
@@ -27286,9 +27286,10 @@ strings must be lower-case.  Valid ISA strings include @samp{sm_30} and
 
 @item -mptx=@var{version-string}
 @opindex mptx
-Generate code for given the specified PTX version (e.g.@: @samp{6.3}).
-Valid version strings include @samp{3.1} and @samp{6.3}.  The default PTX
-version is 3.1.
+Generate code for given the specified PTX version (e.g.@: @samp{7.0}).
+Valid version strings include @samp{3.1}, @samp{6.0}, @samp{6.3}, and
+@samp{7.0}.  The default PTX version is 6.0, unless a higher minimal
+version is required for specified PTX ISA via option @option{-misa=}.
 
 @item -mmainkernel
 @opindex mmainkernel

Re: [Patch] nvptx: Add -mptx=6.0 + -misa=sm_70

[Tom de Vries]

[-- Attachment #1: Type: text/plain, Size: 527 bytes --]

On 2/17/22 18:24, Tobias Burnus wrote:
> SM version (-misa=)
> [Patch adds -misa=sm_70]
> 
> * The compiler supports internally: SM_30, SM_35, SM_53, SM_70, SM_75, 
> SM_80.

I'd formulate it like: it uses SM_70 internally to accurately formulate 
when certain insns can be used.

> I think it makes sense to have sm_70 in addition:
> * The current code actually does generate different code for >= sm_70
>    already.

Agreed.

I've committed this (with a somewhat shorter commit log), and a 
test-case update.

Thanks,
- Tom

[-- Attachment #2: 0001-nvptx-Add-misa-sm_70.patch --]
[-- Type: text/x-patch, Size: 3172 bytes --]

nvptx: Add -misa=sm_70

Add -misa=sm_70, and use it to specify the misa value in test-case
gcc.target/nvptx/atomic-store-2.c.

Tested on nvptx.

gcc/ChangeLog:

	* config/nvptx/nvptx-c.cc (nvptx_cpu_cpp_builtins): Handle SM70.
	* config/nvptx/nvptx.cc (first_ptx_version_supporting_sm):
	Likewise.
	* config/nvptx/nvptx.opt (misa): Add sm_70 alias PTX_ISA_SM70.

gcc/testsuite/ChangeLog:

2022-02-22  Tom de Vries  <tdevries@suse.de>

	* gcc.target/nvptx/atomic-store-2.c: Use -misa=sm_70.
	* gcc.target/nvptx/uniform-simt-3.c: Same.

Co-Authored-By: Tom de Vries <tdevries@suse.de>

---
 gcc/config/nvptx/nvptx-c.cc                     | 2 ++
 gcc/config/nvptx/nvptx.cc                       | 2 ++
 gcc/config/nvptx/nvptx.opt                      | 3 +++
 gcc/testsuite/gcc.target/nvptx/atomic-store-2.c | 2 +-
 gcc/testsuite/gcc.target/nvptx/uniform-simt-3.c | 2 +-
 5 files changed, 9 insertions(+), 2 deletions(-)

diff --git a/gcc/config/nvptx/nvptx-c.cc b/gcc/config/nvptx/nvptx-c.cc
index d68b9910d7e..b2375fb5b16 100644
--- a/gcc/config/nvptx/nvptx-c.cc
+++ b/gcc/config/nvptx/nvptx-c.cc
@@ -43,6 +43,8 @@ nvptx_cpu_cpp_builtins (void)
     cpp_define (parse_in, "__PTX_SM__=800");
   else if (TARGET_SM75)
     cpp_define (parse_in, "__PTX_SM__=750");
+  else if (TARGET_SM70)
+    cpp_define (parse_in, "__PTX_SM__=700");
   else if (TARGET_SM53)
     cpp_define (parse_in, "__PTX_SM__=530");
   else if (TARGET_SM35)
diff --git a/gcc/config/nvptx/nvptx.cc b/gcc/config/nvptx/nvptx.cc
index 981b91f7095..858789e6df7 100644
--- a/gcc/config/nvptx/nvptx.cc
+++ b/gcc/config/nvptx/nvptx.cc
@@ -217,6 +217,8 @@ first_ptx_version_supporting_sm (enum ptx_isa sm)
       return PTX_VERSION_3_1;
     case PTX_ISA_SM53:
       return PTX_VERSION_4_2;
+    case PTX_ISA_SM70:
+      return PTX_VERSION_6_0;
     case PTX_ISA_SM75:
       return PTX_VERSION_6_3;
     case PTX_ISA_SM80:
diff --git a/gcc/config/nvptx/nvptx.opt b/gcc/config/nvptx/nvptx.opt
index 97e127cc4fb..9776c3b9a1f 100644
--- a/gcc/config/nvptx/nvptx.opt
+++ b/gcc/config/nvptx/nvptx.opt
@@ -64,6 +64,9 @@ Enum(ptx_isa) String(sm_35) Value(PTX_ISA_SM35)
 EnumValue
 Enum(ptx_isa) String(sm_53) Value(PTX_ISA_SM53)
 
+EnumValue
+Enum(ptx_isa) String(sm_70) Value(PTX_ISA_SM70)
+
 EnumValue
 Enum(ptx_isa) String(sm_75) Value(PTX_ISA_SM75)
 
diff --git a/gcc/testsuite/gcc.target/nvptx/atomic-store-2.c b/gcc/testsuite/gcc.target/nvptx/atomic-store-2.c
index cd5e4c38267..b58f33f2abd 100644
--- a/gcc/testsuite/gcc.target/nvptx/atomic-store-2.c
+++ b/gcc/testsuite/gcc.target/nvptx/atomic-store-2.c
@@ -2,7 +2,7 @@
    shared state space.  */
 
 /* { dg-do compile } */
-/* { dg-options "-misa=sm_75" } */
+/* { dg-options "-misa=sm_70" } */
 
 enum memmodel
 {
diff --git a/gcc/testsuite/gcc.target/nvptx/uniform-simt-3.c b/gcc/testsuite/gcc.target/nvptx/uniform-simt-3.c
index 532fa825161..b61b8ba9d5b 100644
--- a/gcc/testsuite/gcc.target/nvptx/uniform-simt-3.c
+++ b/gcc/testsuite/gcc.target/nvptx/uniform-simt-3.c
@@ -1,4 +1,4 @@
 /* { dg-do compile } */
-/* { dg-options "-O2 -muniform-simt -misa=sm_75" } */
+/* { dg-options "-O2 -muniform-simt -misa=sm_70" } */
 
 #include "atomic-store-2.c"

Re: [Patch] nvptx: Add -mptx=6.0 + -misa=sm_70

[Tom de Vries]

On 2/17/22 18:24, Tobias Burnus wrote:
> diff --git a/gcc/config/nvptx/t-omp-device b/gcc/config/nvptx/t-omp-device
> index 8765d9f1881..4228218a424 100644
> --- a/gcc/config/nvptx/t-omp-device
> +++ b/gcc/config/nvptx/t-omp-device
> @@ -1,4 +1,4 @@
>  omp-device-properties-nvptx: $(srcdir)/config/nvptx/nvptx.cc
>  	echo kind: gpu > $@
>  	echo arch: nvptx >> $@
> -	echo isa: sm_30 sm_35 >> $@
> +	echo isa: sm_30 sm_35 sm_53 sm_70 sm_75 sm_80 >> $@

I'm not sure I understand how this is used.  Is this user-visible?  Is 
there a libgomp test-case where we can observe a difference?

Thanks,
- Tom

Re: [Patch] nvptx: Add -mptx=6.0 + -misa=sm_70

[Tobias Burnus]

Hi Tom,

On 22.02.22 15:43, Tom de Vries wrote:
> On 2/17/22 18:24, Tobias Burnus wrote:
>> --- a/gcc/config/nvptx/t-omp-device
>> +++ b/gcc/config/nvptx/t-omp-device
>> @@ -1,4 +1,4 @@
>>      echo kind: gpu > $@
>>      echo arch: nvptx >> $@
>> -    echo isa: sm_30 sm_35 >> $@
>> +    echo isa: sm_30 sm_35 sm_53 sm_70 sm_75 sm_80 >> $@
>
> I'm not sure I understand how this is used.  Is this user-visible?  Is
> there a libgomp test-case where we can observe a difference?

That's used for OpenMP context selectors like; that way, one can generate,
e.g. one code used with nvptx and one with gcn as with:

#pragma omp declare variant (on_nvptx) match(construct={target},device={arch(nvptx)})
#pragma omp declare variant (on_gcn) match(construct={target},device={arch(gcn)})
...
   #pragma omp target map(from:v)
   v = on ();
which then either calls 'on' or 'on_nvptx' or 'on_gcn'
(from libgomp/testsuite/libgomp.c/target-42.c)


The following testcases use 'arch(nvptx)':

libgomp/testsuite/libgomp.c-c++-common/on_device_arch.h
libgomp/testsuite/libgomp.c/target-42.c
libgomp/testsuite/libgomp.c/usleep.h
libgomp/testsuite/libgomp.fortran/declare-variant-1.f90

For ISA, there is only one run-time test:

libgomp/testsuite/libgomp.c/declare-variant-1.c

but only for x86-64: match (device={isa("avx512f")})

The sm_35 also appears, but only in the compile-time tests:
gcc/testsuite/{c-c++-common,gfortran.dg}/gomp/declare-variant-{9,10}.*

Tobias

-----------------
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

Re: [Patch] nvptx: Add -mptx=6.0 + -misa=sm_70

[Tom de Vries]

[-- Attachment #1: Type: text/plain, Size: 1782 bytes --]

On 2/22/22 17:03, Tobias Burnus wrote:
> Hi Tom,
> 
> On 22.02.22 15:43, Tom de Vries wrote:
>> On 2/17/22 18:24, Tobias Burnus wrote:
>>> --- a/gcc/config/nvptx/t-omp-device
>>> +++ b/gcc/config/nvptx/t-omp-device
>>> @@ -1,4 +1,4 @@
>>>      echo kind: gpu > $@
>>>      echo arch: nvptx >> $@
>>> -    echo isa: sm_30 sm_35 >> $@
>>> +    echo isa: sm_30 sm_35 sm_53 sm_70 sm_75 sm_80 >> $@
>>
>> I'm not sure I understand how this is used.  Is this user-visible?  Is
>> there a libgomp test-case where we can observe a difference?
> 
> That's used for OpenMP context selectors like; that way, one can generate,
> e.g. one code used with nvptx and one with gcn as with:
> 
> #pragma omp declare variant (on_nvptx) 
> match(construct={target},device={arch(nvptx)})
> #pragma omp declare variant (on_gcn) 
> match(construct={target},device={arch(gcn)})
> ...
>    #pragma omp target map(from:v)
>    v = on ();
> which then either calls 'on' or 'on_nvptx' or 'on_gcn'
> (from libgomp/testsuite/libgomp.c/target-42.c)
> 
> 
> The following testcases use 'arch(nvptx)':
> 
> libgomp/testsuite/libgomp.c-c++-common/on_device_arch.h
> libgomp/testsuite/libgomp.c/target-42.c
> libgomp/testsuite/libgomp.c/usleep.h
> libgomp/testsuite/libgomp.fortran/declare-variant-1.f90
> 
> For ISA, there is only one run-time test:
> 
> libgomp/testsuite/libgomp.c/declare-variant-1.c
> 
> but only for x86-64: match (device={isa("avx512f")})
> 
> The sm_35 also appears, but only in the compile-time tests:
> gcc/testsuite/{c-c++-common,gfortran.dg}/gomp/declare-variant-{9,10}.*
> 

Thanks for the explanation.

I've updated the patch to include changes to 
nvptx_omp_device_kind_arch_isa, and committed.

I'll try to submit a patch with one or more test-cases.

Thanks,
- Tom


[-- Attachment #2: 0003-nvptx-Add-missing-t-omp-device-isas.patch --]
[-- Type: text/x-patch, Size: 1866 bytes --]

[nvptx] Add missing t-omp-device isas

In t-omp-device we list isas that can be used in omp declare variant like so:
...
  #pragma omp declare variant (f30) match (device={isa("sm_30")})
...
and in nvptx_omp_device_kind_arch_isa we handle them.

Update both to reflect the current list of isas.

Tested on x86_64-linux with nvptx accelerator.

gcc/ChangeLog:

2022-02-23  Tom de Vries  <tdevries@suse.de>

	* config/nvptx/nvptx.cc (nvptx_omp_device_kind_arch_isa): Handle
	sm_70, sm_75 and sm_80.
	* config/nvptx/t-omp-device: Add sm_53, sm_70, sm_75 and sm_80.

Co-Authored-By: Tobias Burnus <tobias@codesourcery.com>

---
 gcc/config/nvptx/nvptx.cc     | 8 +++++++-
 gcc/config/nvptx/t-omp-device | 2 +-
 2 files changed, 8 insertions(+), 2 deletions(-)

diff --git a/gcc/config/nvptx/nvptx.cc b/gcc/config/nvptx/nvptx.cc
index 6f6d592e462..b9451c2ed09 100644
--- a/gcc/config/nvptx/nvptx.cc
+++ b/gcc/config/nvptx/nvptx.cc
@@ -6181,7 +6181,13 @@ nvptx_omp_device_kind_arch_isa (enum omp_device_kind_arch_isa trait,
       if (strcmp (name, "sm_35") == 0)
 	return TARGET_SM35 && !TARGET_SM53;
       if (strcmp (name, "sm_53") == 0)
-	return TARGET_SM53;
+	return TARGET_SM53 && !TARGET_SM70;
+      if (strcmp (name, "sm_70") == 0)
+	return TARGET_SM70 && !TARGET_SM75;
+      if (strcmp (name, "sm_75") == 0)
+	return TARGET_SM75 && !TARGET_SM80;
+      if (strcmp (name, "sm_80") == 0)
+	return TARGET_SM80;
       return 0;
     default:
       gcc_unreachable ();
diff --git a/gcc/config/nvptx/t-omp-device b/gcc/config/nvptx/t-omp-device
index 8765d9f1881..4228218a424 100644
--- a/gcc/config/nvptx/t-omp-device
+++ b/gcc/config/nvptx/t-omp-device
@@ -1,4 +1,4 @@
 omp-device-properties-nvptx: $(srcdir)/config/nvptx/nvptx.cc
 	echo kind: gpu > $@
 	echo arch: nvptx >> $@
-	echo isa: sm_30 sm_35 >> $@
+	echo isa: sm_30 sm_35 sm_53 sm_70 sm_75 sm_80 >> $@

[PATCH][libgomp, testsuite, nvptx] Add libgomp.c/declare-variant-3-sm*.c

[Tom de Vries]

[-- Attachment #1: Type: text/plain, Size: 281 bytes --]

[ was: Re: [Patch] nvptx: Add -mptx=6.0 + -misa=sm_70 ]

On 2/24/22 09:29, Tom de Vries wrote:
> I'll try to submit a patch with one or more test-cases.

Hi,

These test-cases exercise the omp declare variant construct using the 
available nvptx isas.

OK for trunk?

Thanks,
- Tom

[-- Attachment #2: 0006-libgomp-testsuite-nvptx-Add-libgomp.c-declare-variant-3-sm-.c.patch --]
[-- Type: text/x-patch, Size: 8085 bytes --]

[libgomp, testsuite, nvptx] Add libgomp.c/declare-variant-3-sm*.c

Add openmp test-cases that test the omp declare variant construct:
...
  #pragma omp declare variant (f30) match (device={isa("sm_30")})
...
using the available nvptx isas.

On a Pascal board GT 1030 with sm_61, we have these unsupported:
...
UNSUPPORTED: libgomp.c/declare-variant-3-sm70.c
UNSUPPORTED: libgomp.c/declare-variant-3-sm75.c
UNSUPPORTED: libgomp.c/declare-variant-3-sm80.c
...
and on a Turing board T400 with sm_75, we have this only this one:
...
UNSUPPORTED: libgomp.c/declare-variant-3-sm80.c
...

Tested on x86_64 with nvptx accelerator.

libgomp/ChangeLog:

2022-02-24  Tom de Vries  <tdevries@suse.de>

	* testsuite/lib/libgomp.exp
	(check_effective_target_offload_device_nvptx_sm_xx)
	(check_effective_target_offload_device_nvptx_sm_30)
	(check_effective_target_offload_device_nvptx_sm_35)
	(check_effective_target_offload_device_nvptx_sm_53)
	(check_effective_target_offload_device_nvptx_sm_70)
	(check_effective_target_offload_device_nvptx_sm_75)
	(check_effective_target_offload_device_nvptx_sm_80): New proc.
	* testsuite/libgomp.c/declare-variant-3-sm30.c: New test.
	* testsuite/libgomp.c/declare-variant-3-sm35.c: New test.
	* testsuite/libgomp.c/declare-variant-3-sm53.c: New test.
	* testsuite/libgomp.c/declare-variant-3-sm70.c: New test.
	* testsuite/libgomp.c/declare-variant-3-sm75.c: New test.
	* testsuite/libgomp.c/declare-variant-3-sm80.c: New test.
	* testsuite/libgomp.c/declare-variant-3.h: New header file.

---
 libgomp/testsuite/lib/libgomp.exp                  | 46 +++++++++++++++
 .../testsuite/libgomp.c/declare-variant-3-sm30.c   |  5 ++
 .../testsuite/libgomp.c/declare-variant-3-sm35.c   |  5 ++
 .../testsuite/libgomp.c/declare-variant-3-sm53.c   |  5 ++
 .../testsuite/libgomp.c/declare-variant-3-sm70.c   |  5 ++
 .../testsuite/libgomp.c/declare-variant-3-sm75.c   |  5 ++
 .../testsuite/libgomp.c/declare-variant-3-sm80.c   |  5 ++
 libgomp/testsuite/libgomp.c/declare-variant-3.h    | 66 ++++++++++++++++++++++
 8 files changed, 142 insertions(+)

diff --git a/libgomp/testsuite/lib/libgomp.exp b/libgomp/testsuite/lib/libgomp.exp
index 8c5ecfff0ac..d664863b15c 100644
--- a/libgomp/testsuite/lib/libgomp.exp
+++ b/libgomp/testsuite/lib/libgomp.exp
@@ -426,6 +426,52 @@ proc check_effective_target_offload_device_nvptx { } {
     } ]
 }
 
+# Return 1 if using nvptx offload device which supports -misa=sm_$SM.
+proc check_effective_target_offload_device_nvptx_sm_xx { sm } {
+    if { ![check_effective_target_offload_device_nvptx] } {
+	return 0
+    }
+    return [check_runtime_nocache offload_device_nvptx_sm_$sm {
+      int main ()
+	{
+	  int x = 1;
+	  #pragma omp target map(tofrom: x)
+	    x--;
+	  return x;
+	}
+    } "-foffload=-misa=sm_$sm" ]
+}
+
+# See check_effective_target_offload_device_nvptx_sm_xx.
+proc check_effective_target_offload_device_nvptx_sm_30 { } {
+    return [check_effective_target_offload_device_nvptx_sm_xx 30]
+}
+
+# See check_effective_target_offload_device_nvptx_sm_xx.
+proc check_effective_target_offload_device_nvptx_sm_35 { } {
+    return [check_effective_target_offload_device_nvptx_sm_xx 35]
+}
+
+# See check_effective_target_offload_device_nvptx_sm_xx.
+proc check_effective_target_offload_device_nvptx_sm_53 { } {
+    return [check_effective_target_offload_device_nvptx_sm_xx 53]
+}
+
+# See check_effective_target_offload_device_nvptx_sm_xx.
+proc check_effective_target_offload_device_nvptx_sm_70 { } {
+    return [check_effective_target_offload_device_nvptx_sm_xx 70]
+}
+
+# See check_effective_target_offload_device_nvptx_sm_xx.
+proc check_effective_target_offload_device_nvptx_sm_75 { } {
+    return [check_effective_target_offload_device_nvptx_sm_xx 75]
+}
+
+# See check_effective_target_offload_device_nvptx_sm_xx.
+proc check_effective_target_offload_device_nvptx_sm_80 { } {
+    return [check_effective_target_offload_device_nvptx_sm_xx 80]
+}
+
 # Return 1 if at least one Nvidia GPU is accessible.
 
 proc check_effective_target_openacc_nvidia_accel_present { } {
diff --git a/libgomp/testsuite/libgomp.c/declare-variant-3-sm30.c b/libgomp/testsuite/libgomp.c/declare-variant-3-sm30.c
new file mode 100644
index 00000000000..7c680b07a94
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c/declare-variant-3-sm30.c
@@ -0,0 +1,5 @@
+/* { dg-do run { target { offload_target_nvptx } } } */
+/* { dg-require-effective-target offload_device_nvptx_sm_30 } */
+/* { dg-additional-options "-foffload=-misa=sm_30" } */
+
+#include "declare-variant-3.h"
diff --git a/libgomp/testsuite/libgomp.c/declare-variant-3-sm35.c b/libgomp/testsuite/libgomp.c/declare-variant-3-sm35.c
new file mode 100644
index 00000000000..b8b2a714248
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c/declare-variant-3-sm35.c
@@ -0,0 +1,5 @@
+/* { dg-do run { target { offload_target_nvptx } } } */
+/* { dg-require-effective-target offload_device_nvptx_sm_35 } */
+/* { dg-additional-options "-foffload=-misa=sm_35" } */
+
+#include "declare-variant-3.h"
diff --git a/libgomp/testsuite/libgomp.c/declare-variant-3-sm53.c b/libgomp/testsuite/libgomp.c/declare-variant-3-sm53.c
new file mode 100644
index 00000000000..cfc7ee6f137
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c/declare-variant-3-sm53.c
@@ -0,0 +1,5 @@
+/* { dg-do run { target { offload_target_nvptx } } } */
+/* { dg-require-effective-target offload_device_nvptx_sm_53 } */
+/* { dg-additional-options "-foffload=-misa=sm_53" } */
+
+#include "declare-variant-3.h"
diff --git a/libgomp/testsuite/libgomp.c/declare-variant-3-sm70.c b/libgomp/testsuite/libgomp.c/declare-variant-3-sm70.c
new file mode 100644
index 00000000000..4527cc1376c
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c/declare-variant-3-sm70.c
@@ -0,0 +1,5 @@
+/* { dg-do run { target { offload_target_nvptx } } } */
+/* { dg-require-effective-target offload_device_nvptx_sm_70 } */
+/* { dg-additional-options "-foffload=-misa=sm_70" } */
+
+#include "declare-variant-3.h"
diff --git a/libgomp/testsuite/libgomp.c/declare-variant-3-sm75.c b/libgomp/testsuite/libgomp.c/declare-variant-3-sm75.c
new file mode 100644
index 00000000000..8e7da369fc5
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c/declare-variant-3-sm75.c
@@ -0,0 +1,5 @@
+/* { dg-do run { target { offload_target_nvptx } } } */
+/* { dg-require-effective-target offload_device_nvptx_sm_75 } */
+/* { dg-additional-options "-foffload=-misa=sm_75" } */
+
+#include "declare-variant-3.h"
diff --git a/libgomp/testsuite/libgomp.c/declare-variant-3-sm80.c b/libgomp/testsuite/libgomp.c/declare-variant-3-sm80.c
new file mode 100644
index 00000000000..5cf5b2867a2
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c/declare-variant-3-sm80.c
@@ -0,0 +1,5 @@
+/* { dg-do run { target { offload_target_nvptx } } } */
+/* { dg-require-effective-target offload_device_nvptx_sm_80 } */
+/* { dg-additional-options "-foffload=-misa=sm_80" } */
+
+#include "declare-variant-3.h"
diff --git a/libgomp/testsuite/libgomp.c/declare-variant-3.h b/libgomp/testsuite/libgomp.c/declare-variant-3.h
new file mode 100644
index 00000000000..772fc20a519
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c/declare-variant-3.h
@@ -0,0 +1,66 @@
+#pragma omp declare target
+int
+f30 (void)
+{
+  return 30;
+}
+
+int
+f35 (void)
+{
+  return 35;
+}
+
+int
+f53 (void)
+{
+  return 53;
+}
+
+int
+f70 (void)
+{
+  return 70;
+}
+
+int
+f75 (void)
+{
+  return 75;
+}
+
+int
+f80 (void)
+{
+  return 80;
+}
+
+#pragma omp declare variant (f30) match (device={isa("sm_30")})
+#pragma omp declare variant (f35) match (device={isa("sm_35")})
+#pragma omp declare variant (f53) match (device={isa("sm_53")})
+#pragma omp declare variant (f70) match (device={isa("sm_70")})
+#pragma omp declare variant (f75) match (device={isa("sm_75")})
+#pragma omp declare variant (f80) match (device={isa("sm_80")})
+int
+f (void)
+{
+  return 0;
+}
+
+#pragma omp end declare target
+
+int
+main (void)
+{
+  int v = 0;
+
+  #pragma omp target map(from:v)
+  v = f ();
+
+  if (v == 0)
+    __builtin_abort ();
+
+  __builtin_printf ("Nvptx accelerator: sm_%d\n", v);
+
+  return 0;
+}

Re: [PATCH][libgomp, testsuite, nvptx] Add libgomp.c/declare-variant-3-sm*.c

[Jakub Jelinek]

On Thu, Feb 24, 2022 at 11:01:22AM +0100, Tom de Vries wrote:
> [ was: Re: [Patch] nvptx: Add -mptx=6.0 + -misa=sm_70 ]
> 
> On 2/24/22 09:29, Tom de Vries wrote:
> > I'll try to submit a patch with one or more test-cases.
> 
> Hi,
> 
> These test-cases exercise the omp declare variant construct using the
> available nvptx isas.
> 
> OK for trunk?
> 
> Thanks,
> - Tom

> [libgomp, testsuite, nvptx] Add libgomp.c/declare-variant-3-sm*.c
> 
> Add openmp test-cases that test the omp declare variant construct:
> ...
>   #pragma omp declare variant (f30) match (device={isa("sm_30")})
> ...
> using the available nvptx isas.
> 
> On a Pascal board GT 1030 with sm_61, we have these unsupported:
> ...
> UNSUPPORTED: libgomp.c/declare-variant-3-sm70.c
> UNSUPPORTED: libgomp.c/declare-variant-3-sm75.c
> UNSUPPORTED: libgomp.c/declare-variant-3-sm80.c
> ...
> and on a Turing board T400 with sm_75, we have this only this one:
> ...
> UNSUPPORTED: libgomp.c/declare-variant-3-sm80.c
> ...
> 
> Tested on x86_64 with nvptx accelerator.

I think testing it through dg-do link tests with -fdump-tree-optimized
or so would be better, you wouldn't need access to actual hardware level
and checking in the dump what function is actually called for each case is
easy.

	Jakub

Re: [PATCH][libgomp, testsuite, nvptx] Add libgomp.c/declare-variant-3-sm*.c

[Tom de Vries]

[-- Attachment #1: Type: text/plain, Size: 1547 bytes --]

On 2/24/22 11:09, Jakub Jelinek wrote:
> On Thu, Feb 24, 2022 at 11:01:22AM +0100, Tom de Vries wrote:
>> [ was: Re: [Patch] nvptx: Add -mptx=6.0 + -misa=sm_70 ]
>>
>> On 2/24/22 09:29, Tom de Vries wrote:
>>> I'll try to submit a patch with one or more test-cases.
>>
>> Hi,
>>
>> These test-cases exercise the omp declare variant construct using the
>> available nvptx isas.
>>
>> OK for trunk?
>>
>> Thanks,
>> - Tom
> 
>> [libgomp, testsuite, nvptx] Add libgomp.c/declare-variant-3-sm*.c
>>
>> Add openmp test-cases that test the omp declare variant construct:
>> ...
>>    #pragma omp declare variant (f30) match (device={isa("sm_30")})
>> ...
>> using the available nvptx isas.
>>
>> On a Pascal board GT 1030 with sm_61, we have these unsupported:
>> ...
>> UNSUPPORTED: libgomp.c/declare-variant-3-sm70.c
>> UNSUPPORTED: libgomp.c/declare-variant-3-sm75.c
>> UNSUPPORTED: libgomp.c/declare-variant-3-sm80.c
>> ...
>> and on a Turing board T400 with sm_75, we have this only this one:
>> ...
>> UNSUPPORTED: libgomp.c/declare-variant-3-sm80.c
>> ...
>>
>> Tested on x86_64 with nvptx accelerator.
> 
> I think testing it through dg-do link tests with -fdump-tree-optimized
> or so would be better, you wouldn't need access to actual hardware level
> and checking in the dump what function is actually called for each case is
> easy.
> 

Done, expect for the sm_30 test which is still dg-do run (although I've 
added the compile time test) which should pass on all boards (since we 
don't support below sm_30).

OK for trunk?

Thanks,
- Tom

[-- Attachment #2: 0006-libgomp-testsuite-nvptx-Add-libgomp.c-declare-variant-3-sm-.c.patch --]
[-- Type: text/x-patch, Size: 5833 bytes --]

[libgomp, testsuite, nvptx] Add libgomp.c/declare-variant-3-sm*.c

Add openmp test-cases that test the omp declare variant construct:
...
  #pragma omp declare variant (f30) match (device={isa("sm_30")})
...
using the available nvptx isas.

Only the one for sm_30 is a dg-do run test-case, the other ones are dg-do
link.

Tested on x86_64 with nvptx accelerator.

libgomp/ChangeLog:

2022-02-24  Tom de Vries  <tdevries@suse.de>

	* testsuite/libgomp.c/declare-variant-3-sm30.c: New test.
	* testsuite/libgomp.c/declare-variant-3-sm35.c: New test.
	* testsuite/libgomp.c/declare-variant-3-sm53.c: New test.
	* testsuite/libgomp.c/declare-variant-3-sm70.c: New test.
	* testsuite/libgomp.c/declare-variant-3-sm75.c: New test.
	* testsuite/libgomp.c/declare-variant-3-sm80.c: New test.
	* testsuite/libgomp.c/declare-variant-3.h: New header file.

---
 .../testsuite/libgomp.c/declare-variant-3-sm30.c   |  7 +++
 .../testsuite/libgomp.c/declare-variant-3-sm35.c   |  7 +++
 .../testsuite/libgomp.c/declare-variant-3-sm53.c   |  7 +++
 .../testsuite/libgomp.c/declare-variant-3-sm70.c   |  7 +++
 .../testsuite/libgomp.c/declare-variant-3-sm75.c   |  7 +++
 .../testsuite/libgomp.c/declare-variant-3-sm80.c   |  7 +++
 libgomp/testsuite/libgomp.c/declare-variant-3.h    | 66 ++++++++++++++++++++++
 7 files changed, 108 insertions(+)

diff --git a/libgomp/testsuite/libgomp.c/declare-variant-3-sm30.c b/libgomp/testsuite/libgomp.c/declare-variant-3-sm30.c
new file mode 100644
index 00000000000..ad1602c13cd
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c/declare-variant-3-sm30.c
@@ -0,0 +1,7 @@
+/* { dg-do run { target { offload_target_nvptx } } } */
+/* { dg-additional-options "-foffload=-misa=sm_30" } */
+/* { dg-additional-options "-foffload=-fdump-tree-optimized" } */
+
+#include "declare-variant-3.h"
+
+/* { dg-final { scan-offload-tree-dump "= f30 \\(\\);" "optimized" } } */
diff --git a/libgomp/testsuite/libgomp.c/declare-variant-3-sm35.c b/libgomp/testsuite/libgomp.c/declare-variant-3-sm35.c
new file mode 100644
index 00000000000..1a7cda2456b
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c/declare-variant-3-sm35.c
@@ -0,0 +1,7 @@
+/* { dg-do link { target { offload_target_nvptx } } } */
+/* { dg-additional-options "-foffload=-misa=sm_35" } */
+/* { dg-additional-options "-foffload=-fdump-tree-optimized" } */
+
+#include "declare-variant-3.h"
+
+/* { dg-final { scan-offload-tree-dump "= f35 \\(\\);" "optimized" } } */
diff --git a/libgomp/testsuite/libgomp.c/declare-variant-3-sm53.c b/libgomp/testsuite/libgomp.c/declare-variant-3-sm53.c
new file mode 100644
index 00000000000..a37b5fdaa28
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c/declare-variant-3-sm53.c
@@ -0,0 +1,7 @@
+/* { dg-do link { target { offload_target_nvptx } } } */
+/* { dg-additional-options "-foffload=-misa=sm_53" } */
+/* { dg-additional-options "-foffload=-fdump-tree-optimized" } */
+
+#include "declare-variant-3.h"
+
+/* { dg-final { scan-offload-tree-dump "= f53 \\(\\);" "optimized" } } */
diff --git a/libgomp/testsuite/libgomp.c/declare-variant-3-sm70.c b/libgomp/testsuite/libgomp.c/declare-variant-3-sm70.c
new file mode 100644
index 00000000000..ab022cd79f9
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c/declare-variant-3-sm70.c
@@ -0,0 +1,7 @@
+/* { dg-do link { target { offload_target_nvptx } } } */
+/* { dg-additional-options "-foffload=-misa=sm_70" } */
+/* { dg-additional-options "-foffload=-fdump-tree-optimized" } */
+
+#include "declare-variant-3.h"
+
+/* { dg-final { scan-offload-tree-dump "= f70 \\(\\);" "optimized" } } */
diff --git a/libgomp/testsuite/libgomp.c/declare-variant-3-sm75.c b/libgomp/testsuite/libgomp.c/declare-variant-3-sm75.c
new file mode 100644
index 00000000000..7d09195d9c4
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c/declare-variant-3-sm75.c
@@ -0,0 +1,7 @@
+/* { dg-do link { target { offload_target_nvptx } } } */
+/* { dg-additional-options "-foffload=-misa=sm_75" } */
+/* { dg-additional-options "-foffload=-fdump-tree-optimized" } */
+
+#include "declare-variant-3.h"
+
+/* { dg-final { scan-offload-tree-dump "= f75 \\(\\);" "optimized" } } */
diff --git a/libgomp/testsuite/libgomp.c/declare-variant-3-sm80.c b/libgomp/testsuite/libgomp.c/declare-variant-3-sm80.c
new file mode 100644
index 00000000000..898ae6e4da8
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c/declare-variant-3-sm80.c
@@ -0,0 +1,7 @@
+/* { dg-do link { target { offload_target_nvptx } } } */
+/* { dg-additional-options "-foffload=-misa=sm_80" } */
+/* { dg-additional-options "-foffload=-fdump-tree-optimized" } */
+
+#include "declare-variant-3.h"
+
+/* { dg-final { scan-offload-tree-dump "= f80 \\(\\);" "optimized" } } */
diff --git a/libgomp/testsuite/libgomp.c/declare-variant-3.h b/libgomp/testsuite/libgomp.c/declare-variant-3.h
new file mode 100644
index 00000000000..772fc20a519
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c/declare-variant-3.h
@@ -0,0 +1,66 @@
+#pragma omp declare target
+int
+f30 (void)
+{
+  return 30;
+}
+
+int
+f35 (void)
+{
+  return 35;
+}
+
+int
+f53 (void)
+{
+  return 53;
+}
+
+int
+f70 (void)
+{
+  return 70;
+}
+
+int
+f75 (void)
+{
+  return 75;
+}
+
+int
+f80 (void)
+{
+  return 80;
+}
+
+#pragma omp declare variant (f30) match (device={isa("sm_30")})
+#pragma omp declare variant (f35) match (device={isa("sm_35")})
+#pragma omp declare variant (f53) match (device={isa("sm_53")})
+#pragma omp declare variant (f70) match (device={isa("sm_70")})
+#pragma omp declare variant (f75) match (device={isa("sm_75")})
+#pragma omp declare variant (f80) match (device={isa("sm_80")})
+int
+f (void)
+{
+  return 0;
+}
+
+#pragma omp end declare target
+
+int
+main (void)
+{
+  int v = 0;
+
+  #pragma omp target map(from:v)
+  v = f ();
+
+  if (v == 0)
+    __builtin_abort ();
+
+  __builtin_printf ("Nvptx accelerator: sm_%d\n", v);
+
+  return 0;
+}

Re: [PATCH][libgomp, testsuite, nvptx] Add libgomp.c/declare-variant-3-sm*.c

[Jakub Jelinek]

On Thu, Feb 24, 2022 at 11:32:53AM +0100, Tom de Vries wrote:
> libgomp/ChangeLog:
> 
> 2022-02-24  Tom de Vries  <tdevries@suse.de>
> 
> 	* testsuite/libgomp.c/declare-variant-3-sm30.c: New test.
> 	* testsuite/libgomp.c/declare-variant-3-sm35.c: New test.
> 	* testsuite/libgomp.c/declare-variant-3-sm53.c: New test.
> 	* testsuite/libgomp.c/declare-variant-3-sm70.c: New test.
> 	* testsuite/libgomp.c/declare-variant-3-sm75.c: New test.
> 	* testsuite/libgomp.c/declare-variant-3-sm80.c: New test.
> 	* testsuite/libgomp.c/declare-variant-3.h: New header file.

LGTM, thanks.

	Jakub

Fix 'libgomp.c/declare-variant-3-*.c' compilation for configurations where GCN offloading is enabled in addition to nvptx (was: [PATCH][libgomp, testsuite, nvptx] Add libgomp.c/declare-variant-3-sm*.c)

[Thomas Schwinge]

Hi!

On 2022-02-24T11:32:53+0100, Tom de Vries via Gcc-patches <gcc-patches@gcc.gnu.org> wrote:
> [libgomp, testsuite, nvptx] Add libgomp.c/declare-variant-3-sm*.c

\o/ Yay for test cases!

> --- /dev/null
> +++ b/libgomp/testsuite/libgomp.c/declare-variant-3-sm30.c
> @@ -0,0 +1,7 @@
> +/* { dg-do run { target { offload_target_nvptx } } } */
> +/* { dg-additional-options "-foffload=-misa=sm_30" } */
> +/* { dg-additional-options "-foffload=-fdump-tree-optimized" } */

Etc.

Pushed to master branch commit 3f5a3b7539e066b539e81b901687facdea4e1bac
"Fix 'libgomp.c/declare-variant-3-*.c' compilation for configurations where GCN offloading is enabled in addition to nvptx",
see attached.


Grüße
 Thomas
-----------------
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

Re: Fix 'libgomp.c/declare-variant-3-*.c' compilation for configurations where GCN offloading is enabled in addition to nvptx (was: [PATCH][libgomp, testsuite, nvptx] Add libgomp.c/declare-variant-3-sm*.c)

[Thomas Schwinge]

[-- Attachment #1: Type: text/plain, Size: 1080 bytes --]

Hi!

On 2023-11-30T15:48:25+0100, I wrote:
> On 2022-02-24T11:32:53+0100, Tom de Vries via Gcc-patches <gcc-patches@gcc.gnu.org> wrote:
>> [libgomp, testsuite, nvptx] Add libgomp.c/declare-variant-3-sm*.c
>
> \o/ Yay for test cases!
>
>> --- /dev/null
>> +++ b/libgomp/testsuite/libgomp.c/declare-variant-3-sm30.c
>> @@ -0,0 +1,7 @@
>> +/* { dg-do run { target { offload_target_nvptx } } } */
>> +/* { dg-additional-options "-foffload=-misa=sm_30" } */
>> +/* { dg-additional-options "-foffload=-fdump-tree-optimized" } */
>
> Etc.
>
> Pushed to master branch commit 3f5a3b7539e066b539e81b901687facdea4e1bac
> "Fix 'libgomp.c/declare-variant-3-*.c' compilation for configurations where GCN offloading is enabled in addition to nvptx",
> see attached.

..., now.


Grüße
 Thomas


-----------------
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

[-- Warning: decoded text below may be mangled, UTF-8 assumed --]
[-- Attachment #2: 0001-Fix-libgomp.c-declare-variant-3-.c-compilation-for-c.patch --]
[-- Type: text/x-diff, Size: 4647 bytes --]

From 3f5a3b7539e066b539e81b901687facdea4e1bac Mon Sep 17 00:00:00 2001
From: Thomas Schwinge <thomas@codesourcery.com>
Date: Fri, 25 Feb 2022 12:55:05 +0100
Subject: [PATCH] Fix 'libgomp.c/declare-variant-3-*.c' compilation for
 configurations where GCN offloading is enabled in addition to nvptx

The GCN offloading compiler doesn't like '-misa=sm_30' etc.; restrict to
'-foffload=nvptx-none' compilation only.

Fix-up for commit 59b8ade88774b4dcf1691a8f650cdbb86cc30862
"[libgomp, testsuite, nvptx] Add libgomp.c/declare-variant-3-sm*.c".

	libgomp/
	* testsuite/libgomp.c/declare-variant-3-sm30.c:
	'dg-additional-options -foffload=nvptx-none'.
	* testsuite/libgomp.c/declare-variant-3-sm35.c: Likewise.
	* testsuite/libgomp.c/declare-variant-3-sm53.c: Likewise.
	* testsuite/libgomp.c/declare-variant-3-sm70.c: Likewise.
	* testsuite/libgomp.c/declare-variant-3-sm75.c: Likewise.
	* testsuite/libgomp.c/declare-variant-3-sm80.c: Likewise.
---
 libgomp/testsuite/libgomp.c/declare-variant-3-sm30.c | 1 +
 libgomp/testsuite/libgomp.c/declare-variant-3-sm35.c | 1 +
 libgomp/testsuite/libgomp.c/declare-variant-3-sm53.c | 1 +
 libgomp/testsuite/libgomp.c/declare-variant-3-sm70.c | 1 +
 libgomp/testsuite/libgomp.c/declare-variant-3-sm75.c | 1 +
 libgomp/testsuite/libgomp.c/declare-variant-3-sm80.c | 1 +
 6 files changed, 6 insertions(+)

diff --git a/libgomp/testsuite/libgomp.c/declare-variant-3-sm30.c b/libgomp/testsuite/libgomp.c/declare-variant-3-sm30.c
index a49bc12064a..be5edbaf0fd 100644
--- a/libgomp/testsuite/libgomp.c/declare-variant-3-sm30.c
+++ b/libgomp/testsuite/libgomp.c/declare-variant-3-sm30.c
@@ -1,4 +1,5 @@
 /* { dg-do run { target { offload_target_nvptx } } } */
+/* { dg-additional-options -foffload=nvptx-none } */
 /* { dg-additional-options "-foffload=-misa=sm_30 -foffload=-mptx=_" } */
 /* { dg-additional-options "-foffload=-fdump-tree-optimized" } */
 
diff --git a/libgomp/testsuite/libgomp.c/declare-variant-3-sm35.c b/libgomp/testsuite/libgomp.c/declare-variant-3-sm35.c
index 9f71acb8738..79cd8fc940d 100644
--- a/libgomp/testsuite/libgomp.c/declare-variant-3-sm35.c
+++ b/libgomp/testsuite/libgomp.c/declare-variant-3-sm35.c
@@ -1,4 +1,5 @@
 /* { dg-do link { target { offload_target_nvptx } } } */
+/* { dg-additional-options -foffload=nvptx-none } */
 /* { dg-additional-options "-foffload=-misa=sm_35 -foffload=-mptx=_" } */
 /* { dg-additional-options "-foffload=-fdump-tree-optimized" } */
 
diff --git a/libgomp/testsuite/libgomp.c/declare-variant-3-sm53.c b/libgomp/testsuite/libgomp.c/declare-variant-3-sm53.c
index fa713920ce0..8d05e9efc62 100644
--- a/libgomp/testsuite/libgomp.c/declare-variant-3-sm53.c
+++ b/libgomp/testsuite/libgomp.c/declare-variant-3-sm53.c
@@ -1,4 +1,5 @@
 /* { dg-do link { target { offload_target_nvptx } } } */
+/* { dg-additional-options -foffload=nvptx-none } */
 /* { dg-additional-options "-foffload=-misa=sm_53 -foffload=-mptx=_" } */
 /* { dg-additional-options "-foffload=-fdump-tree-optimized" } */
 
diff --git a/libgomp/testsuite/libgomp.c/declare-variant-3-sm70.c b/libgomp/testsuite/libgomp.c/declare-variant-3-sm70.c
index 90f0116c582..e102c11ab44 100644
--- a/libgomp/testsuite/libgomp.c/declare-variant-3-sm70.c
+++ b/libgomp/testsuite/libgomp.c/declare-variant-3-sm70.c
@@ -1,4 +1,5 @@
 /* { dg-do link { target { offload_target_nvptx } } } */
+/* { dg-additional-options -foffload=nvptx-none } */
 /* { dg-additional-options "-foffload=-misa=sm_70 -foffload=-mptx=_" } */
 /* { dg-additional-options "-foffload=-fdump-tree-optimized" } */
 
diff --git a/libgomp/testsuite/libgomp.c/declare-variant-3-sm75.c b/libgomp/testsuite/libgomp.c/declare-variant-3-sm75.c
index 86f2e72866a..7122ccc22d1 100644
--- a/libgomp/testsuite/libgomp.c/declare-variant-3-sm75.c
+++ b/libgomp/testsuite/libgomp.c/declare-variant-3-sm75.c
@@ -1,4 +1,5 @@
 /* { dg-do link { target { offload_target_nvptx } } } */
+/* { dg-additional-options -foffload=nvptx-none } */
 /* { dg-additional-options "-foffload=-misa=sm_75 -foffload=-mptx=_" } */
 /* { dg-additional-options "-foffload=-fdump-tree-optimized" } */
 
diff --git a/libgomp/testsuite/libgomp.c/declare-variant-3-sm80.c b/libgomp/testsuite/libgomp.c/declare-variant-3-sm80.c
index de208d9bdd1..5e13d876500 100644
--- a/libgomp/testsuite/libgomp.c/declare-variant-3-sm80.c
+++ b/libgomp/testsuite/libgomp.c/declare-variant-3-sm80.c
@@ -1,4 +1,5 @@
 /* { dg-do link { target { offload_target_nvptx } } } */
+/* { dg-additional-options -foffload=nvptx-none } */
 /* { dg-additional-options "-foffload=-misa=sm_80 -foffload=-mptx=_" } */
 /* { dg-additional-options "-foffload=-fdump-tree-optimized" } */
 
-- 
2.34.1

In 'libgomp.c/declare-variant-{3,4}-*.c', restrict 'scan-offload-tree-dump's to 'only_for_offload_target [...]' (was: [PATCH][libgomp, testsuite, nvptx] Add libgomp.c/declare-variant-3-sm*.c)

[Thomas Schwinge]

[-- Attachment #1: Type: text/plain, Size: 1518 bytes --]

Hi!

On 2022-02-24T11:32:53+0100, Tom de Vries via Gcc-patches <gcc-patches@gcc.gnu.org> wrote:
> --- /dev/null
> +++ b/libgomp/testsuite/libgomp.c/declare-variant-3-sm30.c
> @@ -0,0 +1,7 @@
> +/* { dg-do run { target { offload_target_nvptx } } } */

This means, the test case is active if nvptx offloading compilation is
enabled.  But, consider the case that GCN offloading compilation also
is enabled:

| +/* { dg-additional-options -foffload=nvptx-none } */
> +/* { dg-additional-options "-foffload=-misa=sm_30" } */

This means, the test cases is restricted to nvptx offloading compilation,
but not GCN.

> +/* { dg-additional-options "-foffload=-fdump-tree-optimized" } */

This will thus produce a dump file for nvptx, only.  No dump file
produced for GCN.

> +[...]
> +/* { dg-final { scan-offload-tree-dump "= f30 \\(\\);" "optimized" } } */

..., and this will try to scan dump files for both GCN and nvptx.  The
former doesn't exist, resulting in UNRESOLVEDs for GCN.  Etc.  I've
pushed to master branch commit 186e22c5de8eb49603bb1e74ac5d0eba6cc40480
"In 'libgomp.c/declare-variant-{3,4}-*.c', restrict 'scan-offload-tree-dump's to 'only_for_offload_target [...]'",
see attached.


Grüße
 Thomas


-----------------
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

[-- Warning: decoded text below may be mangled, UTF-8 assumed --]
[-- Attachment #2: 0001-In-libgomp.c-declare-variant-3-4-.c-restrict-scan-of.patch --]
[-- Type: text/x-diff, Size: 8897 bytes --]

From 186e22c5de8eb49603bb1e74ac5d0eba6cc40480 Mon Sep 17 00:00:00 2001
From: Thomas Schwinge <thomas@codesourcery.com>
Date: Tue, 21 Nov 2023 17:31:37 +0100
Subject: [PATCH] In 'libgomp.c/declare-variant-{3,4}-*.c', restrict
 'scan-offload-tree-dump's to 'only_for_offload_target [...]'

... to care for the case where not just one but both of GCN and nvptx
offloading are enabled.  In that case, we currently get:

    UNRESOLVED: libgomp.c/declare-variant-3-sm30.c scan-amdgcn-amdhsa-offload-tree-dump optimized "= f30 \\(\\);"

... in addition to:

    PASS: libgomp.c/declare-variant-3-sm30.c scan-nvptx-none-offload-tree-dump optimized "= f30 \\(\\);"

Etc.

Fix-up for commit 59b8ade88774b4dcf1691a8f650cdbb86cc30862
"[libgomp, testsuite, nvptx] Add libgomp.c/declare-variant-3-sm*.c",
and commit 1fd508744eccda9ad9c6d6fcce5b2ea9c568818d
"amdgcn: Support AMD-specific 'isa' traits in OpenMP context selectors".

	libgomp/
	* testsuite/libgomp.c/declare-variant-3-sm30.c: Restrict
	'scan-offload-tree-dump' to 'only_for_offload_target nvptx-none'.
	* testsuite/libgomp.c/declare-variant-3-sm35.c: Likewise.
	* testsuite/libgomp.c/declare-variant-3-sm53.c: Likewise.
	* testsuite/libgomp.c/declare-variant-3-sm70.c: Likewise.
	* testsuite/libgomp.c/declare-variant-3-sm75.c: Likewise.
	* testsuite/libgomp.c/declare-variant-3-sm80.c: Likewise.
	* testsuite/libgomp.c/declare-variant-4-fiji.c: Restrict
	'scan-offload-tree-dump' to
	'only_for_offload_target amdgcn-amdhsa'.
	* testsuite/libgomp.c/declare-variant-4-gfx803.c: Likewise.
	* testsuite/libgomp.c/declare-variant-4-gfx900.c: Likewise.
	* testsuite/libgomp.c/declare-variant-4-gfx906.c: Likewise.
	* testsuite/libgomp.c/declare-variant-4-gfx908.c: Likewise.
	* testsuite/libgomp.c/declare-variant-4-gfx90a.c: Likewise.
---
 libgomp/testsuite/libgomp.c/declare-variant-3-sm30.c   | 2 +-
 libgomp/testsuite/libgomp.c/declare-variant-3-sm35.c   | 2 +-
 libgomp/testsuite/libgomp.c/declare-variant-3-sm53.c   | 2 +-
 libgomp/testsuite/libgomp.c/declare-variant-3-sm70.c   | 2 +-
 libgomp/testsuite/libgomp.c/declare-variant-3-sm75.c   | 2 +-
 libgomp/testsuite/libgomp.c/declare-variant-3-sm80.c   | 2 +-
 libgomp/testsuite/libgomp.c/declare-variant-4-fiji.c   | 2 +-
 libgomp/testsuite/libgomp.c/declare-variant-4-gfx803.c | 2 +-
 libgomp/testsuite/libgomp.c/declare-variant-4-gfx900.c | 2 +-
 libgomp/testsuite/libgomp.c/declare-variant-4-gfx906.c | 2 +-
 libgomp/testsuite/libgomp.c/declare-variant-4-gfx908.c | 2 +-
 libgomp/testsuite/libgomp.c/declare-variant-4-gfx90a.c | 2 +-
 12 files changed, 12 insertions(+), 12 deletions(-)

diff --git a/libgomp/testsuite/libgomp.c/declare-variant-3-sm30.c b/libgomp/testsuite/libgomp.c/declare-variant-3-sm30.c
index be5edbaf0fd..a373647bb33 100644
--- a/libgomp/testsuite/libgomp.c/declare-variant-3-sm30.c
+++ b/libgomp/testsuite/libgomp.c/declare-variant-3-sm30.c
@@ -5,4 +5,4 @@
 
 #include "declare-variant-3.h"
 
-/* { dg-final { scan-offload-tree-dump "= f30 \\(\\);" "optimized" } } */
+/* { dg-final { only_for_offload_target nvptx-none scan-offload-tree-dump "= f30 \\(\\);" "optimized" } } */
diff --git a/libgomp/testsuite/libgomp.c/declare-variant-3-sm35.c b/libgomp/testsuite/libgomp.c/declare-variant-3-sm35.c
index 79cd8fc940d..bf8dc3e2441 100644
--- a/libgomp/testsuite/libgomp.c/declare-variant-3-sm35.c
+++ b/libgomp/testsuite/libgomp.c/declare-variant-3-sm35.c
@@ -5,4 +5,4 @@
 
 #include "declare-variant-3.h"
 
-/* { dg-final { scan-offload-tree-dump "= f35 \\(\\);" "optimized" } } */
+/* { dg-final { only_for_offload_target nvptx-none scan-offload-tree-dump "= f35 \\(\\);" "optimized" } } */
diff --git a/libgomp/testsuite/libgomp.c/declare-variant-3-sm53.c b/libgomp/testsuite/libgomp.c/declare-variant-3-sm53.c
index 8d05e9efc62..1864a7a9277 100644
--- a/libgomp/testsuite/libgomp.c/declare-variant-3-sm53.c
+++ b/libgomp/testsuite/libgomp.c/declare-variant-3-sm53.c
@@ -5,4 +5,4 @@
 
 #include "declare-variant-3.h"
 
-/* { dg-final { scan-offload-tree-dump "= f53 \\(\\);" "optimized" } } */
+/* { dg-final { only_for_offload_target nvptx-none scan-offload-tree-dump "= f53 \\(\\);" "optimized" } } */
diff --git a/libgomp/testsuite/libgomp.c/declare-variant-3-sm70.c b/libgomp/testsuite/libgomp.c/declare-variant-3-sm70.c
index e102c11ab44..2249cd4c24d 100644
--- a/libgomp/testsuite/libgomp.c/declare-variant-3-sm70.c
+++ b/libgomp/testsuite/libgomp.c/declare-variant-3-sm70.c
@@ -5,4 +5,4 @@
 
 #include "declare-variant-3.h"
 
-/* { dg-final { scan-offload-tree-dump "= f70 \\(\\);" "optimized" } } */
+/* { dg-final { only_for_offload_target nvptx-none scan-offload-tree-dump "= f70 \\(\\);" "optimized" } } */
diff --git a/libgomp/testsuite/libgomp.c/declare-variant-3-sm75.c b/libgomp/testsuite/libgomp.c/declare-variant-3-sm75.c
index 7122ccc22d1..18ede59c541 100644
--- a/libgomp/testsuite/libgomp.c/declare-variant-3-sm75.c
+++ b/libgomp/testsuite/libgomp.c/declare-variant-3-sm75.c
@@ -5,4 +5,4 @@
 
 #include "declare-variant-3.h"
 
-/* { dg-final { scan-offload-tree-dump "= f75 \\(\\);" "optimized" } } */
+/* { dg-final { only_for_offload_target nvptx-none scan-offload-tree-dump "= f75 \\(\\);" "optimized" } } */
diff --git a/libgomp/testsuite/libgomp.c/declare-variant-3-sm80.c b/libgomp/testsuite/libgomp.c/declare-variant-3-sm80.c
index 5e13d876500..20048f1d702 100644
--- a/libgomp/testsuite/libgomp.c/declare-variant-3-sm80.c
+++ b/libgomp/testsuite/libgomp.c/declare-variant-3-sm80.c
@@ -5,4 +5,4 @@
 
 #include "declare-variant-3.h"
 
-/* { dg-final { scan-offload-tree-dump "= f80 \\(\\);" "optimized" } } */
+/* { dg-final { only_for_offload_target nvptx-none scan-offload-tree-dump "= f80 \\(\\);" "optimized" } } */
diff --git a/libgomp/testsuite/libgomp.c/declare-variant-4-fiji.c b/libgomp/testsuite/libgomp.c/declare-variant-4-fiji.c
index ae2af1cc00c..8a4e0f4728c 100644
--- a/libgomp/testsuite/libgomp.c/declare-variant-4-fiji.c
+++ b/libgomp/testsuite/libgomp.c/declare-variant-4-fiji.c
@@ -5,4 +5,4 @@
 #define USE_FIJI_FOR_GFX803
 #include "declare-variant-4.h"
 
-/* { dg-final { scan-offload-tree-dump "= gfx803 \\(\\);" "optimized" } } */
+/* { dg-final { only_for_offload_target amdgcn-amdhsa scan-offload-tree-dump "= gfx803 \\(\\);" "optimized" } } */
diff --git a/libgomp/testsuite/libgomp.c/declare-variant-4-gfx803.c b/libgomp/testsuite/libgomp.c/declare-variant-4-gfx803.c
index e0437a04d65..050d7c9dd79 100644
--- a/libgomp/testsuite/libgomp.c/declare-variant-4-gfx803.c
+++ b/libgomp/testsuite/libgomp.c/declare-variant-4-gfx803.c
@@ -4,4 +4,4 @@
 
 #include "declare-variant-4.h"
 
-/* { dg-final { scan-offload-tree-dump "= gfx803 \\(\\);" "optimized" } } */
+/* { dg-final { only_for_offload_target amdgcn-amdhsa scan-offload-tree-dump "= gfx803 \\(\\);" "optimized" } } */
diff --git a/libgomp/testsuite/libgomp.c/declare-variant-4-gfx900.c b/libgomp/testsuite/libgomp.c/declare-variant-4-gfx900.c
index 8de03725dec..2eeb4a248c1 100644
--- a/libgomp/testsuite/libgomp.c/declare-variant-4-gfx900.c
+++ b/libgomp/testsuite/libgomp.c/declare-variant-4-gfx900.c
@@ -4,4 +4,4 @@
 
 #include "declare-variant-4.h"
 
-/* { dg-final { scan-offload-tree-dump "= gfx900 \\(\\);" "optimized" } } */
+/* { dg-final { only_for_offload_target amdgcn-amdhsa scan-offload-tree-dump "= gfx900 \\(\\);" "optimized" } } */
diff --git a/libgomp/testsuite/libgomp.c/declare-variant-4-gfx906.c b/libgomp/testsuite/libgomp.c/declare-variant-4-gfx906.c
index be6f193ed3a..73a670dcc2a 100644
--- a/libgomp/testsuite/libgomp.c/declare-variant-4-gfx906.c
+++ b/libgomp/testsuite/libgomp.c/declare-variant-4-gfx906.c
@@ -4,4 +4,4 @@
 
 #include "declare-variant-4.h"
 
-/* { dg-final { scan-offload-tree-dump "= gfx906 \\(\\);" "optimized" } } */
+/* { dg-final { only_for_offload_target amdgcn-amdhsa scan-offload-tree-dump "= gfx906 \\(\\);" "optimized" } } */
diff --git a/libgomp/testsuite/libgomp.c/declare-variant-4-gfx908.c b/libgomp/testsuite/libgomp.c/declare-variant-4-gfx908.c
index 311fad9074d..ead330f9f2c 100644
--- a/libgomp/testsuite/libgomp.c/declare-variant-4-gfx908.c
+++ b/libgomp/testsuite/libgomp.c/declare-variant-4-gfx908.c
@@ -4,4 +4,4 @@
 
 #include "declare-variant-4.h"
 
-/* { dg-final { scan-offload-tree-dump "= gfx908 \\(\\);" "optimized" } } */
+/* { dg-final { only_for_offload_target amdgcn-amdhsa scan-offload-tree-dump "= gfx908 \\(\\);" "optimized" } } */
diff --git a/libgomp/testsuite/libgomp.c/declare-variant-4-gfx90a.c b/libgomp/testsuite/libgomp.c/declare-variant-4-gfx90a.c
index 96cc14ca0a3..a9b2d62a49d 100644
--- a/libgomp/testsuite/libgomp.c/declare-variant-4-gfx90a.c
+++ b/libgomp/testsuite/libgomp.c/declare-variant-4-gfx90a.c
@@ -4,4 +4,4 @@
 
 #include "declare-variant-4.h"
 
-/* { dg-final { scan-offload-tree-dump "= gfx90a \\(\\);" "optimized" } } */
+/* { dg-final { only_for_offload_target amdgcn-amdhsa scan-offload-tree-dump "= gfx90a \\(\\);" "optimized" } } */
-- 
2.34.1

Spin 'dg-do run' part of 'libgomp.c/declare-variant-3-sm30.c' off into new 'libgomp.c/declare-variant-3.c' (was: [PATCH][libgomp, testsuite, nvptx] Add libgomp.c/declare-variant-3-sm*.c)

[Thomas Schwinge]

[-- Attachment #1: Type: text/plain, Size: 759 bytes --]

Hi!

On 2022-02-24T11:32:53+0100, Tom de Vries via Gcc-patches <gcc-patches@gcc.gnu.org> wrote:
> the sm_30 test which is still dg-do run (although I've
> added the compile time test) which should pass on all boards (since we
> don't support below sm_30).

Pushed to master branch commit 95e6e32a85566a5510d21938d439e90e504e0ddc
"Spin 'dg-do run' part of 'libgomp.c/declare-variant-3-sm30.c' off into new 'libgomp.c/declare-variant-3.c'",
see attached.


Grüße
 Thomas


-----------------
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

[-- Warning: decoded text below may be mangled, UTF-8 assumed --]
[-- Attachment #2: 0001-Spin-dg-do-run-part-of-libgomp.c-declare-variant-3-s.patch --]
[-- Type: text/x-diff, Size: 3104 bytes --]

From 95e6e32a85566a5510d21938d439e90e504e0ddc Mon Sep 17 00:00:00 2001
From: Thomas Schwinge <thomas@codesourcery.com>
Date: Tue, 21 Nov 2023 19:03:47 +0100
Subject: [PATCH] Spin 'dg-do run' part of 'libgomp.c/declare-variant-3-sm30.c'
 off into new 'libgomp.c/declare-variant-3.c'

Having nvptx offloading configured doesn't imply being able to run nvptx
offloading test cases on the test host.

Also, make 'libgomp.c/declare-variant-3.c' work for all non-offloading and
offloading cases.

Fix-up for commit 59b8ade88774b4dcf1691a8f650cdbb86cc30862
"[libgomp, testsuite, nvptx] Add libgomp.c/declare-variant-3-sm*.c".

	libgomp/
	* testsuite/libgomp.c/declare-variant-3-sm30.c: Turn 'dg-do run'
	into 'dg-do link'.
	* testsuite/libgomp.c/declare-variant-3.c: New.
	* testsuite/libgomp.c/declare-variant-3.h: Extend.
---
 libgomp/testsuite/libgomp.c/declare-variant-3-sm30.c | 2 +-
 libgomp/testsuite/libgomp.c/declare-variant-3.c      | 8 ++++++++
 libgomp/testsuite/libgomp.c/declare-variant-3.h      | 5 +++++
 3 files changed, 14 insertions(+), 1 deletion(-)
 create mode 100644 libgomp/testsuite/libgomp.c/declare-variant-3.c

diff --git a/libgomp/testsuite/libgomp.c/declare-variant-3-sm30.c b/libgomp/testsuite/libgomp.c/declare-variant-3-sm30.c
index a373647bb33..d2ffa5637c5 100644
--- a/libgomp/testsuite/libgomp.c/declare-variant-3-sm30.c
+++ b/libgomp/testsuite/libgomp.c/declare-variant-3-sm30.c
@@ -1,4 +1,4 @@
-/* { dg-do run { target { offload_target_nvptx } } } */
+/* { dg-do link { target { offload_target_nvptx } } } */
 /* { dg-additional-options -foffload=nvptx-none } */
 /* { dg-additional-options "-foffload=-misa=sm_30 -foffload=-mptx=_" } */
 /* { dg-additional-options "-foffload=-fdump-tree-optimized" } */
diff --git a/libgomp/testsuite/libgomp.c/declare-variant-3.c b/libgomp/testsuite/libgomp.c/declare-variant-3.c
new file mode 100644
index 00000000000..62c1fa766ba
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c/declare-variant-3.c
@@ -0,0 +1,8 @@
+/* { dg-additional-options -DOFFLOAD_DEVICE_NVPTX { target offload_device_nvptx } } */
+/* { dg-additional-options {-fdump-tree-optimized -foffload-options=-fdump-tree-optimized} } */
+
+#include "declare-variant-3.h"
+
+/* { dg-final { scan-tree-dump "= f \\(\\);" "optimized" } }
+   { dg-final { only_for_offload_target amdgcn-amdhsa scan-offload-tree-dump "= f \\(\\);" "optimized" { target offload_target_amdgcn } } }
+   { dg-final { only_for_offload_target nvptx-none scan-offload-tree-dump "= f\[0-9\]+ \\(\\);" "optimized" { target offload_target_nvptx } } } */
diff --git a/libgomp/testsuite/libgomp.c/declare-variant-3.h b/libgomp/testsuite/libgomp.c/declare-variant-3.h
index 646e15e5311..38ee257e42d 100644
--- a/libgomp/testsuite/libgomp.c/declare-variant-3.h
+++ b/libgomp/testsuite/libgomp.c/declare-variant-3.h
@@ -65,10 +65,15 @@ main (void)
   #pragma omp target map(from:v)
   v = f ();
 
+#ifdef OFFLOAD_DEVICE_NVPTX
   if (v == 0)
     __builtin_abort ();
 
   __builtin_printf ("Nvptx accelerator: sm_%d\n", v);
+#else
+  if (v != 0)
+    __builtin_abort ();
+#endif
 
   return 0;
 }
-- 
2.34.1