[PATCH][OG12] amdgcn: Support AMD-specific 'isa' and 'arch' traits in OpenMP context selectors

[Paul-Antoine Arras <pa@codesourcery.com>]

[-- Attachment #1: Type: text/plain, Size: 489 bytes --]

Hi all,

This patch adds or fixes support for various AMD 'isa' and 'arch' trait 
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 
'metadirective' construct.

This patch is closely related to 
https://gcc.gnu.org/r13-4403-g1fd508744eccda but cannot be committed to 
mainline because metadirectives and dynamic context selectors have not 
landed there yet.

Can this be committed to OG12?

Thanks,

[-- Attachment #2: 0001-amdgcn-Support-AMD-specific-isa-and-arch-traits-in-O.patch --]
[-- Type: text/plain, Size: 3953 bytes --]

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