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 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 <tdevr...@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 <tob...@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 >> $@
