Re: [PATCH][libgomp, testsuite, nvptx] Add libgomp.c/declare-variant-3-sm*.c

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

[-- Attachment #1: Type: text/plain, Size: 1547 bytes --]

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

[-- Attachment #2: 0006-libgomp-testsuite-nvptx-Add-libgomp.c-declare-variant-3-sm-.c.patch --]
[-- Type: text/x-patch, Size: 5833 bytes --]

[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  <tdevries@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;
+}