[ 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.
libgomp/ChangeLog:
2022-02-24 Tom de Vries <tdevr...@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;
+}
