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,
From 88522107dd39ba3ff8465cf688fe4438fa3b77b4 Mon Sep 17 00:00:00 2001
From: Paul-Antoine Arras <p...@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
