On 01/12/2022 13:45, Andrew Stubbs wrote:
P.S. If you want to split the patch into the GCN bits and the bits that
depend on metadirectives then we can apply the first part to mainline
right away.
So this is the OG12-specific part (including metadirective and dynamic
context selectors) of the previous patch.
Once https://gcc.gnu.org/r13-4446-ge41b243302e996 is backported, is it
OK for OG12?
Thanks,
--
PA
From 494a815af459b13da6fe9bf5a84b94d4b1f94915 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 libgomp support for 'amdgcn' as arch, and for each processor type (as passed
to '-march') 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 '-march' values (as well as 'gfx803') as isa traits.
* testsuite/libgomp.c-c++-common/metadirective-6.c: New test.
---
libgomp/ChangeLog.omp | 6 +++
libgomp/config/gcn/selector.c | 24 ++++++++--
.../libgomp.c-c++-common/metadirective-6.c | 48 +++++++++++++++++++
3 files changed, 73 insertions(+), 5 deletions(-)
create mode 100644 libgomp/testsuite/libgomp.c-c++-common/metadirective-6.c
diff --git libgomp/ChangeLog.omp libgomp/ChangeLog.omp
index 74053a6eea0..a2f03914725 100644
--- libgomp/ChangeLog.omp
+++ libgomp/ChangeLog.omp
@@ -1,3 +1,9 @@
+2022-12-01 Paul-Antoine Arras <p...@codesourcery.com>
+
+ * config/gcn/selector.c (GOMP_evaluate_current_device): Recognise
'amdgcn'
+ as arch, and '-march' values (as well as 'gfx803') as isa traits.
+ * testsuite/libgomp.c-c++-common/metadirective-6.c: New test.
+
2022-11-30 Tobias Burnus <tob...@codesourcery.com>
Backported from master:
diff --git libgomp/config/gcn/selector.c libgomp/config/gcn/selector.c
index 60793fc05d3..570bc1e8ae6 100644
--- libgomp/config/gcn/selector.c
+++ libgomp/config/gcn/selector.c
@@ -36,20 +36,34 @@ 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)
return true;
-#ifdef __GCN3__
+#ifdef __gfx803__
if (strcmp (isa, "fiji") == 0 || strcmp (isa, "gfx803") == 0)
return true;
#endif
-#ifdef __GCN5__
- if (strcmp (isa, "gfx900") == 0 || strcmp (isa, "gfx906") != 0
- || strcmp (isa, "gfx908") == 0)
+#ifdef __gfx900__
+ if (strcmp (isa, "gfx900") == 0)
+ return true;
+#endif
+
+#ifdef __gfx906__
+ if (strcmp (isa, "gfx906") == 0)
+ return true;
+#endif
+
+#ifdef __gfx908__
+ if (strcmp (isa, "gfx908") == 0)
+ return true;
+#endif
+
+#ifdef __gfx90a__
+ 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
