Fix 'libgomp.c/declare-variant-4-*.c', add 'libgomp.c/declare-variant-4.c'

These test cases being 'dg-skip-if [...] { ! amdgcn-*-* }' meant they were
only ever considered for 'amdgcn-*-*' target -- that is, GCN *target*, not
GCN *offload target*, what is intended to be tested here, but for which they
thus always were UNSUPPORTED.  Use the same style of 'dg-[...]' directives
as for the nvptx offloading test cases, 'libgomp.c/declare-variant-3*.c'.

Fix-up for commit 1fd508744eccda9ad9c6d6fcce5b2ea9c568818d
"amdgcn: Support AMD-specific 'isa' traits in OpenMP context selectors".

	libgomp/
	* testsuite/libgomp.c/declare-variant-4-fiji.c: Adjust.
	* testsuite/libgomp.c/declare-variant-4-gfx803.c: Likewise.
	* testsuite/libgomp.c/declare-variant-4-gfx900.c: Likewise.
	* testsuite/libgomp.c/declare-variant-4-gfx906.c: Likewise.
	* testsuite/libgomp.c/declare-variant-4-gfx908.c: Likewise.
	* testsuite/libgomp.c/declare-variant-4-gfx90a.c: Likewise.
	* testsuite/libgomp.c/declare-variant-4.h: Likewise.
	* testsuite/libgomp.c/declare-variant-4.c: New.

diff --git a/libgomp/testsuite/libgomp.c/declare-variant-4-fiji.c b/libgomp/testsuite/libgomp.c/declare-variant-4-fiji.c
index 8a4e0f4728c0958ec7861a1da51e8b0b651301bc..a138fb092f8276606b83d98ed572168c9fc2f92e 100644 (file)
--- a/libgomp/testsuite/libgomp.c/declare-variant-4-fiji.c
+++ b/libgomp/testsuite/libgomp.c/declare-variant-4-fiji.c
@@ -1,5 +1,6 @@
-/* { dg-do run { target { offload_target_amdgcn } } } */
-/* { dg-skip-if "fiji/gfx803 only" { ! amdgcn-*-* } { "*" } { "-foffload=-march=fiji" } } */
+/* { dg-do link { target { offload_target_amdgcn } } } */
+/* { dg-additional-options -foffload=amdgcn-amdhsa } */
+/* { dg-additional-options -foffload=-march=fiji } */
 /* { dg-additional-options "-foffload=-fdump-tree-optimized" } */
 
 #define USE_FIJI_FOR_GFX803

diff --git a/libgomp/testsuite/libgomp.c/declare-variant-4-gfx803.c b/libgomp/testsuite/libgomp.c/declare-variant-4-gfx803.c
index 050d7c9dd79578c2bea88f6c1b80830da7676ff8..03dffddac49588d750d289c67c14d27cc07a6c67 100644 (file)
--- a/libgomp/testsuite/libgomp.c/declare-variant-4-gfx803.c
+++ b/libgomp/testsuite/libgomp.c/declare-variant-4-gfx803.c
@@ -1,5 +1,6 @@
-/* { dg-do run { target { offload_target_amdgcn } } } */
-/* { dg-skip-if "fiji/gfx803 only" { ! amdgcn-*-* } { "*" } { "-foffload=-march=fiji" } } */
+/* { dg-do link { target { offload_target_amdgcn } } } */
+/* { dg-additional-options -foffload=amdgcn-amdhsa } */
+/* { dg-additional-options -foffload=-march=fiji } */
 /* { dg-additional-options "-foffload=-fdump-tree-optimized" } */
 
 #include "declare-variant-4.h"

diff --git a/libgomp/testsuite/libgomp.c/declare-variant-4-gfx900.c b/libgomp/testsuite/libgomp.c/declare-variant-4-gfx900.c
index 2eeb4a248c184fb820fac356c8b3caf5e9eba66b..f3f5244c7bc779109dcb7b4e395622c00c0b5346 100644 (file)
--- a/libgomp/testsuite/libgomp.c/declare-variant-4-gfx900.c
+++ b/libgomp/testsuite/libgomp.c/declare-variant-4-gfx900.c
@@ -1,5 +1,6 @@
-/* { dg-do run { target { offload_target_amdgcn } } } */
-/* { dg-skip-if "gfx900 only" { ! amdgcn-*-* } { "*" } { "-foffload=-march=gfx900" } } */
+/* { dg-do link { target { offload_target_amdgcn } } } */
+/* { dg-additional-options -foffload=amdgcn-amdhsa } */
+/* { dg-additional-options -foffload=-march=gfx900 } */
 /* { dg-additional-options "-foffload=-fdump-tree-optimized" } */
 
 #include "declare-variant-4.h"

diff --git a/libgomp/testsuite/libgomp.c/declare-variant-4-gfx906.c b/libgomp/testsuite/libgomp.c/declare-variant-4-gfx906.c
index 73a670dcc2a874e5b373d349688ad28657b634e5..ac43388a59f0ebee858a4deee0b2fb373a842c99 100644 (file)
--- a/libgomp/testsuite/libgomp.c/declare-variant-4-gfx906.c
+++ b/libgomp/testsuite/libgomp.c/declare-variant-4-gfx906.c
@@ -1,5 +1,6 @@
-/* { dg-do run { target { offload_target_amdgcn } } } */
-/* { dg-skip-if "gfx906 only" { ! amdgcn-*-* } { "*" } { "-foffload=-march=gfx906" } } */
+/* { dg-do link { target { offload_target_amdgcn } } } */
+/* { dg-additional-options -foffload=amdgcn-amdhsa } */
+/* { dg-additional-options -foffload=-march=gfx906 } */
 /* { dg-additional-options "-foffload=-fdump-tree-optimized" } */
 
 #include "declare-variant-4.h"

diff --git a/libgomp/testsuite/libgomp.c/declare-variant-4-gfx908.c b/libgomp/testsuite/libgomp.c/declare-variant-4-gfx908.c
index ead330f9f2cd8db7375ff8e3c29ff860ba998be1..f60741f202fac42e16dde0fba65485fa8b1b6171 100644 (file)
--- a/libgomp/testsuite/libgomp.c/declare-variant-4-gfx908.c
+++ b/libgomp/testsuite/libgomp.c/declare-variant-4-gfx908.c
@@ -1,5 +1,6 @@
-/* { dg-do run { target { offload_target_amdgcn } } } */
-/* { dg-skip-if "gfx908 only" { ! amdgcn-*-* } { "*" } { "-foffload=-march=gfx908" } } */
+/* { dg-do link { target { offload_target_amdgcn } } } */
+/* { dg-additional-options -foffload=amdgcn-amdhsa } */
+/* { dg-additional-options -foffload=-march=gfx908 } */
 /* { dg-additional-options "-foffload=-fdump-tree-optimized" } */
 
 #include "declare-variant-4.h"

diff --git a/libgomp/testsuite/libgomp.c/declare-variant-4-gfx90a.c b/libgomp/testsuite/libgomp.c/declare-variant-4-gfx90a.c
index a9b2d62a49da06a0cd6058a89d2ed17d55e99e24..832d174e4a5b2b7ee73ee8d606ced73f8bdf760c 100644 (file)
--- a/libgomp/testsuite/libgomp.c/declare-variant-4-gfx90a.c
+++ b/libgomp/testsuite/libgomp.c/declare-variant-4-gfx90a.c
@@ -1,5 +1,6 @@
-/* { dg-do run { target { offload_target_amdgcn } } } */
-/* { dg-skip-if "gfx90a only" { ! amdgcn-*-* } { "*" } { "-foffload=-march=gfx90a" } } */
+/* { dg-do link { target { offload_target_amdgcn } } } */
+/* { dg-additional-options -foffload=amdgcn-amdhsa } */
+/* { dg-additional-options -foffload=-march=gfx90a } */
 /* { dg-additional-options "-foffload=-fdump-tree-optimized" } */
 
 #include "declare-variant-4.h"

diff --git a/libgomp/testsuite/libgomp.c/declare-variant-4.c b/libgomp/testsuite/libgomp.c/declare-variant-4.c
new file mode 100644 (file)
index 0000000..3c8802a
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c/declare-variant-4.c
@@ -0,0 +1,8 @@
+/* { dg-additional-options -DOFFLOAD_DEVICE_GCN { target offload_device_gcn } } */
+/* { dg-additional-options {-fdump-tree-optimized -foffload-options=-fdump-tree-optimized} } */
+
+#include "declare-variant-4.h"
+
+/* { dg-final { scan-tree-dump "= f \\(\\);" "optimized" } }
+   { dg-final { only_for_offload_target amdgcn-amdhsa scan-offload-tree-dump "= gfx\[^ \]+ \\(\\);" "optimized" { target offload_target_amdgcn } } }
+   { dg-final { only_for_offload_target nvptx-none scan-offload-tree-dump "= f \\(\\);" "optimized" { target offload_target_nvptx } } } */

diff --git a/libgomp/testsuite/libgomp.c/declare-variant-4.h b/libgomp/testsuite/libgomp.c/declare-variant-4.h
index 47517b75ee7f8c2c71672489423157bb8f8146eb..a70352430c2727b51a003628fd58f7cfee4b8000 100644 (file)
--- a/libgomp/testsuite/libgomp.c/declare-variant-4.h
+++ b/libgomp/testsuite/libgomp.c/declare-variant-4.h
@@ -61,10 +61,15 @@ main (void)
 #pragma omp target map(from : v)
   v = f ();
 
+#ifdef OFFLOAD_DEVICE_GCN
   if (v == 0)
     __builtin_abort ();
 
   __builtin_printf ("AMDGCN accelerator: gfx%x\n", v);
+#else
+  if (v != 0)
+    __builtin_abort ();
+#endif
 
   return 0;
 }