[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.

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" } } */

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" } } */

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" } } */

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" } } */

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" } } */

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" } } */

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