On 2/24/22 11:09, Jakub Jelinek wrote:
On Thu, Feb 24, 2022 at 11:01:22AM +0100, Tom de Vries wrote:
[ 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.
I think testing it through dg-do link tests with -fdump-tree-optimized
or so would be better, you wouldn't need access to actual hardware level
and checking in the dump what function is actually called for each case is
easy.
Done, expect for the sm_30 test which is still dg-do run (although I've
added the compile time test) which should pass on all boards (since we
don't support below sm_30).
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.
Only the one for sm_30 is a dg-do run test-case, the other ones are dg-do
link.
Tested on x86_64 with nvptx accelerator.
libgomp/ChangeLog:
2022-02-24 Tom de Vries <tdevr...@suse.de>
* 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.
---
.../testsuite/libgomp.c/declare-variant-3-sm30.c | 7 +++
.../testsuite/libgomp.c/declare-variant-3-sm35.c | 7 +++
.../testsuite/libgomp.c/declare-variant-3-sm53.c | 7 +++
.../testsuite/libgomp.c/declare-variant-3-sm70.c | 7 +++
.../testsuite/libgomp.c/declare-variant-3-sm75.c | 7 +++
.../testsuite/libgomp.c/declare-variant-3-sm80.c | 7 +++
libgomp/testsuite/libgomp.c/declare-variant-3.h | 66 ++++++++++++++++++++++
7 files changed, 108 insertions(+)
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..ad1602c13cd
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c/declare-variant-3-sm30.c
@@ -0,0 +1,7 @@
+/* { dg-do run { target { offload_target_nvptx } } } */
+/* { dg-additional-options "-foffload=-misa=sm_30" } */
+/* { dg-additional-options "-foffload=-fdump-tree-optimized" } */
+
+#include "declare-variant-3.h"
+
+/* { dg-final { scan-offload-tree-dump "= f30 \\(\\);" "optimized" } } */
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..1a7cda2456b
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c/declare-variant-3-sm35.c
@@ -0,0 +1,7 @@
+/* { dg-do link { target { offload_target_nvptx } } } */
+/* { dg-additional-options "-foffload=-misa=sm_35" } */
+/* { dg-additional-options "-foffload=-fdump-tree-optimized" } */
+
+#include "declare-variant-3.h"
+
+/* { dg-final { scan-offload-tree-dump "= f35 \\(\\);" "optimized" } } */
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..a37b5fdaa28
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c/declare-variant-3-sm53.c
@@ -0,0 +1,7 @@
+/* { dg-do link { target { offload_target_nvptx } } } */
+/* { dg-additional-options "-foffload=-misa=sm_53" } */
+/* { dg-additional-options "-foffload=-fdump-tree-optimized" } */
+
+#include "declare-variant-3.h"
+
+/* { dg-final { scan-offload-tree-dump "= f53 \\(\\);" "optimized" } } */
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..ab022cd79f9
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c/declare-variant-3-sm70.c
@@ -0,0 +1,7 @@
+/* { dg-do link { target { offload_target_nvptx } } } */
+/* { dg-additional-options "-foffload=-misa=sm_70" } */
+/* { dg-additional-options "-foffload=-fdump-tree-optimized" } */
+
+#include "declare-variant-3.h"
+
+/* { dg-final { scan-offload-tree-dump "= f70 \\(\\);" "optimized" } } */
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..7d09195d9c4
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c/declare-variant-3-sm75.c
@@ -0,0 +1,7 @@
+/* { dg-do link { target { offload_target_nvptx } } } */
+/* { dg-additional-options "-foffload=-misa=sm_75" } */
+/* { dg-additional-options "-foffload=-fdump-tree-optimized" } */
+
+#include "declare-variant-3.h"
+
+/* { dg-final { scan-offload-tree-dump "= f75 \\(\\);" "optimized" } } */
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..898ae6e4da8
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c/declare-variant-3-sm80.c
@@ -0,0 +1,7 @@
+/* { dg-do link { target { offload_target_nvptx } } } */
+/* { dg-additional-options "-foffload=-misa=sm_80" } */
+/* { dg-additional-options "-foffload=-fdump-tree-optimized" } */
+
+#include "declare-variant-3.h"
+
+/* { dg-final { scan-offload-tree-dump "= f80 \\(\\);" "optimized" } } */
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;
+}
