From 88522107dd39ba3ff8465cf688fe4438fa3b77b4 Mon Sep 17 00:00:00 2001
From: Paul-Antoine Arras <pa@codesourcery.com>
Date: Wed, 30 Nov 2022 14:52:55 +0100
Subject: [PATCH] amdgcn: Support AMD-specific 'isa' and 'arch' traits in
 OpenMP context selectors

Add or fix libgomp support for 'amdgcn' as arch, and 'gfx908' and 'gfx90a' as isa traits.
Add test case for all supported 'isa' values used as context selectors in a metadirective construct..

libgomp/ChangeLog:

	* config/gcn/selector.c (GOMP_evaluate_current_device): Recognise 'amdgcn' as arch, and 'gfx908' and
	'gfx90a' as isa traits.
	* testsuite/libgomp.c-c++-common/metadirective-6.c: New test.
---
 libgomp/config/gcn/selector.c                 | 15 ++++--
 .../libgomp.c-c++-common/metadirective-6.c    | 48 +++++++++++++++++++
 2 files changed, 60 insertions(+), 3 deletions(-)
 create mode 100644 libgomp/testsuite/libgomp.c-c++-common/metadirective-6.c

diff --git libgomp/config/gcn/selector.c libgomp/config/gcn/selector.c
index 60793fc05d3..c948497c538 100644
--- libgomp/config/gcn/selector.c
+++ libgomp/config/gcn/selector.c
@@ -36,7 +36,7 @@ GOMP_evaluate_current_device (const char *kind, const char *arch,
   if (kind && strcmp (kind, "gpu") != 0)
     return false;
 
-  if (arch && strcmp (arch, "gcn") != 0)
+  if (arch && (strcmp (arch, "gcn") != 0 || strcmp (arch, "amdgcn") != 0))
     return false;
 
   if (!isa)
@@ -48,8 +48,17 @@ GOMP_evaluate_current_device (const char *kind, const char *arch,
 #endif
 
 #ifdef __GCN5__
-  if (strcmp (isa, "gfx900") == 0 || strcmp (isa, "gfx906") != 0
-      || strcmp (isa, "gfx908") == 0)
+  if (strcmp (isa, "gfx900") == 0 || strcmp (isa, "gfx906") != 0)
+    return true;
+#endif
+
+#ifdef __CDNA1__
+  if (strcmp (isa, "gfx908") == 0)
+    return true;
+#endif
+
+#ifdef __CDNA2__
+  if (strcmp (isa, "gfx90a") == 0)
     return true;
 #endif
 
diff --git libgomp/testsuite/libgomp.c-c++-common/metadirective-6.c libgomp/testsuite/libgomp.c-c++-common/metadirective-6.c
new file mode 100644
index 00000000000..6d169001db1
--- /dev/null
+++ libgomp/testsuite/libgomp.c-c++-common/metadirective-6.c
@@ -0,0 +1,48 @@
+/* { dg-do link { target { offload_target_amdgcn } } } */
+/* { dg-additional-options "-foffload=-fdump-tree-omp_expand_metadirective" } */
+
+#define N 100
+
+void f (int x[], int y[], int z[])
+{
+  int i;
+
+  #pragma omp target map(to: x, y) map(from: z)
+    #pragma omp metadirective \
+      when (device={isa("gfx803")}: teams num_teams(512)) \
+      when (device={isa("gfx900")}: teams num_teams(256)) \
+      when (device={isa("gfx906")}: teams num_teams(128)) \
+      when (device={isa("gfx908")}: teams num_teams(64)) \
+      when (device={isa("gfx90a")}: teams num_teams(32)) \
+      default (teams num_teams(4))
+	for (i = 0; i < N; i++)
+	  z[i] = x[i] * y[i];
+}
+
+int main (void)
+{
+  int x[N], y[N], z[N];
+  int i;
+
+  for (i = 0; i < N; i++)
+    {
+      x[i] = i;
+      y[i] = -i;
+    }
+
+  f (x, y, z);
+
+  for (i = 0; i < N; i++)
+    if (z[i] != x[i] * y[i])
+      return 1;
+
+  return 0;
+}
+
+/* The metadirective should be resolved after Gimplification.  */
+
+/* { dg-final { scan-offload-tree-dump "__builtin_GOMP_teams4 \\(512, 512" "omp_expand_metadirective" { target { any-opts "-foffload=-march=fiji" } } } } */
+/* { dg-final { scan-offload-tree-dump "__builtin_GOMP_teams4 \\(256, 256" "omp_expand_metadirective" { target { any-opts "-foffload=-march=gfx900" } } } } */
+/* { dg-final { scan-offload-tree-dump "__builtin_GOMP_teams4 \\(128, 128" "omp_expand_metadirective" { target { any-opts "-foffload=-march=gfx906" } } } } */
+/* { dg-final { scan-offload-tree-dump "__builtin_GOMP_teams4 \\(64, 64" "omp_expand_metadirective" { target { any-opts "-foffload=-march=gfx908" } } } } */
+/* { dg-final { scan-offload-tree-dump "__builtin_GOMP_teams4 \\(32, 32" "omp_expand_metadirective" { target { any-opts "-foffload=-march=gfx90a" } } } } */
-- 
2.31.1