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

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

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

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

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

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

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