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

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

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

[ 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

[-- Attachment #2: 0006-libgomp-testsuite-nvptx-Add-libgomp.c-declare-variant-3-sm-.c.patch --]
[-- Type: text/x-patch, Size: 8085 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.

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