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

[Tom de Vries <tdevries@suse.de>]

[-- 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 >> $@