Message ID | a17f985a-a003-8c5f-f406-121f54282aef@codesourcery.com |
---|---|
State | New |
Headers | show |
Series | nvptx: Add -mptx=6.0 + -misa=sm_70 | expand |
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
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