diff mbox series

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

Message ID dd63de7d-171c-bc9b-a3c5-5a3254c1c8a2@codesourcery.com
State Accepted
Headers
Series amdgcn: Support AMD-specific 'isa' traits in OpenMP context selectors |

Checks

Context Check Description
snail/gcc-patch-check success Github commit url

Commit Message

Paul-Antoine Arras Nov. 29, 2022, 3:56 p.m. UTC 
  Hi all,

This patch adds support for 'gfx803' as an alias for 'fiji' in OpenMP 
context selectors, so as to be consistent with LLVM. It also adds test 
cases checking all supported AMD ISAs are properly recognised when used 
in a 'declare variant' construct.

Is it OK for mainline?

Thanks,

Comments

Andrew Stubbs Nov. 29, 2022, 5:51 p.m. UTC | #1 
On 29/11/2022 15:56, Paul-Antoine Arras wrote:
> Hi all,
> 
> This patch adds support for 'gfx803' as an alias for 'fiji' in OpenMP 
> context selectors, so as to be consistent with LLVM. It also adds test 
> cases checking all supported AMD ISAs are properly recognised when used 
> in a 'declare variant' construct.
> 
> Is it OK for mainline?

OK

Andrew
diff mbox series

Patch

diff --git gcc/config/gcn/gcn.cc gcc/config/gcn/gcn.cc
index c74fa007a21..39e93aeaeef 100644
--- gcc/config/gcn/gcn.cc
+++ 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 gcc/config/gcn/t-omp-device gcc/config/gcn/t-omp-device
index 27d36db894b..538624f7ec7 100644
--- gcc/config/gcn/t-omp-device
+++ 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 libgomp/testsuite/libgomp.c/declare-variant-4-fiji.c libgomp/testsuite/libgomp.c/declare-variant-4-fiji.c
new file mode 100644
index 00000000000..ae2af1cc00c
--- /dev/null
+++ 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 libgomp/testsuite/libgomp.c/declare-variant-4-gfx803.c libgomp/testsuite/libgomp.c/declare-variant-4-gfx803.c
new file mode 100644
index 00000000000..e0437a04d65
--- /dev/null
+++ 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 libgomp/testsuite/libgomp.c/declare-variant-4-gfx900.c libgomp/testsuite/libgomp.c/declare-variant-4-gfx900.c
new file mode 100644
index 00000000000..8de03725dec
--- /dev/null
+++ 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 libgomp/testsuite/libgomp.c/declare-variant-4-gfx906.c libgomp/testsuite/libgomp.c/declare-variant-4-gfx906.c
new file mode 100644
index 00000000000..be6f193ed3a
--- /dev/null
+++ 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 libgomp/testsuite/libgomp.c/declare-variant-4-gfx908.c libgomp/testsuite/libgomp.c/declare-variant-4-gfx908.c
new file mode 100644
index 00000000000..311fad9074d
--- /dev/null
+++ 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 libgomp/testsuite/libgomp.c/declare-variant-4-gfx90a.c libgomp/testsuite/libgomp.c/declare-variant-4-gfx90a.c
new file mode 100644
index 00000000000..96cc14ca0a3
--- /dev/null
+++ 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 libgomp/testsuite/libgomp.c/declare-variant-4.h libgomp/testsuite/libgomp.c/declare-variant-4.h
new file mode 100644
index 00000000000..2d7c1ef1a5a
--- /dev/null
+++ 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;
+}