amdgcn: Support AMD-specific 'isa' traits in OpenMP context selectors

Add support for gfx803 as an alias for fiji.
Add test cases for all supported 'isa' values.

gcc/ChangeLog:

	* config/gcn/gcn.cc (gcn_omp_device_kind_arch_isa): Add gfx803.
	* config/gcn/t-omp-device: Add gfx803.

libgomp/ChangeLog:

	* testsuite/libgomp.c/declare-variant-4-fiji.c: New test.
	* testsuite/libgomp.c/declare-variant-4-gfx803.c: New test.
	* testsuite/libgomp.c/declare-variant-4-gfx900.c: New test.
	* testsuite/libgomp.c/declare-variant-4-gfx906.c: New test.
	* testsuite/libgomp.c/declare-variant-4-gfx908.c: New test.
	* testsuite/libgomp.c/declare-variant-4-gfx90a.c: New test.
	* testsuite/libgomp.c/declare-variant-4.h: New header file.

diff --git a/gcc/config/gcn/gcn.cc b/gcc/config/gcn/gcn.cc
index c74fa007a211271c3031639f38f80536c10d166f..39e93aeaeef14e24831a464bae548b0848ce2b0f 100644 (file)
--- a/gcc/config/gcn/gcn.cc
+++ b/gcc/config/gcn/gcn.cc
@@ -2985,7 +2985,7 @@ gcn_omp_device_kind_arch_isa (enum omp_device_kind_arch_isa trait,
     case omp_device_arch:
       return strcmp (name, "amdgcn") == 0 || strcmp (name, "gcn") == 0;
     case omp_device_isa:
-      if (strcmp (name, "fiji") == 0)
+      if (strcmp (name, "fiji") == 0 || strcmp (name, "gfx803") == 0)
        return gcn_arch == PROCESSOR_FIJI;
       if (strcmp (name, "gfx900") == 0)
        return gcn_arch == PROCESSOR_VEGA10;

diff --git a/gcc/config/gcn/t-omp-device b/gcc/config/gcn/t-omp-device
index 27d36db894bed8845b240831867688a9aff44da2..538624f7ec7d77a1c94e399f85287f282740a905 100644 (file)
--- a/gcc/config/gcn/t-omp-device
+++ b/gcc/config/gcn/t-omp-device
@@ -1,4 +1,4 @@
 omp-device-properties-gcn: $(srcdir)/config/gcn/gcn.cc
        echo kind: gpu > $@
        echo arch: amdgcn gcn >> $@
-       echo isa: fiji gfx900 gfx906 gfx908 gfx90a >> $@
+       echo isa: fiji gfx803 gfx900 gfx906 gfx908 gfx90a >> $@

diff --git a/libgomp/testsuite/libgomp.c/declare-variant-4-fiji.c b/libgomp/testsuite/libgomp.c/declare-variant-4-fiji.c
new file mode 100644 (file)
index 0000000..ae2af1c
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c/declare-variant-4-fiji.c
@@ -0,0 +1,8 @@
+/* { dg-do run { target { offload_target_amdgcn } } } */
+/* { dg-skip-if "fiji/gfx803 only" { ! amdgcn-*-* } { "*" } { "-foffload=-march=fiji" } } */
+/* { dg-additional-options "-foffload=-fdump-tree-optimized" } */
+
+#define USE_FIJI_FOR_GFX803
+#include "declare-variant-4.h"
+
+/* { dg-final { scan-offload-tree-dump "= gfx803 \\(\\);" "optimized" } } */

diff --git a/libgomp/testsuite/libgomp.c/declare-variant-4-gfx803.c b/libgomp/testsuite/libgomp.c/declare-variant-4-gfx803.c
new file mode 100644 (file)
index 0000000..e0437a0
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c/declare-variant-4-gfx803.c
@@ -0,0 +1,7 @@
+/* { dg-do run { target { offload_target_amdgcn } } } */
+/* { dg-skip-if "fiji/gfx803 only" { ! amdgcn-*-* } { "*" } { "-foffload=-march=fiji" } } */
+/* { dg-additional-options "-foffload=-fdump-tree-optimized" } */
+
+#include "declare-variant-4.h"
+
+/* { dg-final { scan-offload-tree-dump "= gfx803 \\(\\);" "optimized" } } */

diff --git a/libgomp/testsuite/libgomp.c/declare-variant-4-gfx900.c b/libgomp/testsuite/libgomp.c/declare-variant-4-gfx900.c
new file mode 100644 (file)
index 0000000..8de0372
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c/declare-variant-4-gfx900.c
@@ -0,0 +1,7 @@
+/* { dg-do run { target { offload_target_amdgcn } } } */
+/* { dg-skip-if "gfx900 only" { ! amdgcn-*-* } { "*" } { "-foffload=-march=gfx900" } } */
+/* { dg-additional-options "-foffload=-fdump-tree-optimized" } */
+
+#include "declare-variant-4.h"
+
+/* { dg-final { scan-offload-tree-dump "= gfx900 \\(\\);" "optimized" } } */

diff --git a/libgomp/testsuite/libgomp.c/declare-variant-4-gfx906.c b/libgomp/testsuite/libgomp.c/declare-variant-4-gfx906.c
new file mode 100644 (file)
index 0000000..be6f193
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c/declare-variant-4-gfx906.c
@@ -0,0 +1,7 @@
+/* { dg-do run { target { offload_target_amdgcn } } } */
+/* { dg-skip-if "gfx906 only" { ! amdgcn-*-* } { "*" } { "-foffload=-march=gfx906" } } */
+/* { dg-additional-options "-foffload=-fdump-tree-optimized" } */
+
+#include "declare-variant-4.h"
+
+/* { dg-final { scan-offload-tree-dump "= gfx906 \\(\\);" "optimized" } } */

diff --git a/libgomp/testsuite/libgomp.c/declare-variant-4-gfx908.c b/libgomp/testsuite/libgomp.c/declare-variant-4-gfx908.c
new file mode 100644 (file)
index 0000000..311fad9
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c/declare-variant-4-gfx908.c
@@ -0,0 +1,7 @@
+/* { dg-do run { target { offload_target_amdgcn } } } */
+/* { dg-skip-if "gfx908 only" { ! amdgcn-*-* } { "*" } { "-foffload=-march=gfx908" } } */
+/* { dg-additional-options "-foffload=-fdump-tree-optimized" } */
+
+#include "declare-variant-4.h"
+
+/* { dg-final { scan-offload-tree-dump "= gfx908 \\(\\);" "optimized" } } */

diff --git a/libgomp/testsuite/libgomp.c/declare-variant-4-gfx90a.c b/libgomp/testsuite/libgomp.c/declare-variant-4-gfx90a.c
new file mode 100644 (file)
index 0000000..96cc14c
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c/declare-variant-4-gfx90a.c
@@ -0,0 +1,7 @@
+/* { dg-do run { target { offload_target_amdgcn } } } */
+/* { dg-skip-if "gfx90a only" { ! amdgcn-*-* } { "*" } { "-foffload=-march=gfx90a" } } */
+/* { dg-additional-options "-foffload=-fdump-tree-optimized" } */
+
+#include "declare-variant-4.h"
+
+/* { dg-final { scan-offload-tree-dump "= gfx90a \\(\\);" "optimized" } } */

diff --git a/libgomp/testsuite/libgomp.c/declare-variant-4.h b/libgomp/testsuite/libgomp.c/declare-variant-4.h
new file mode 100644 (file)
index 0000000..2d7c1ef
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c/declare-variant-4.h
@@ -0,0 +1,63 @@
+#pragma omp declare target
+int
+gfx803 (void)
+{
+  return 0x803;
+}
+
+int
+gfx900 (void)
+{
+  return 0x900;
+}
+
+int
+gfx906 (void)
+{
+  return 0x906;
+}
+
+int
+gfx908 (void)
+{
+  return 0x908;
+}
+
+int
+gfx90a (void)
+{
+  return 0x90a;
+}
+
+#ifdef USE_FIJI_FOR_GFX803
+#pragma omp declare variant(gfx803) match(device = {isa("fiji")})
+#else
+#pragma omp declare variant(gfx803) match(device = {isa("gfx803")})
+#endif
+#pragma omp declare variant(gfx900) match(device = {isa("gfx900")})
+#pragma omp declare variant(gfx906) match(device = {isa("gfx906")})
+#pragma omp declare variant(gfx908) match(device = {isa("gfx908")})
+#pragma omp declare variant(gfx90a) match(device = {isa("gfx90a")})
+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 ("AMDGCN accelerator: gfx%x\n", v);
+
+  return 0;
+}