[OG12] amdgcn: Support AMD-specific 'isa' and 'arch' traits in OpenMP context selectors
Commit Message
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 <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
Comments
Hello PA,
> --- 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;
The logic here looks wrong to me - surely it should return false if arch
is not 'gcn' AND it is not 'amdgcn'?
> @@ -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
Okay for gfx908 and gfx90a, but is there any way of distinguishing
between 'gfx900' and 'gfx906' ISAs? I don't think these are mutually
compatible.
Thanks
Kwok
Hi Kwok,
On 30/11/2022 19:50, Kwok Cheung Yeung wrote:
> Hello PA,
>
>> --- 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;
>
> The logic here looks wrong to me - surely it should return false if arch
> is not 'gcn' AND it is not 'amdgcn'?
Sure. Fixed in revised patch.
>> @@ -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
>
> Okay for gfx908 and gfx90a, but is there any way of distinguishing
> between 'gfx900' and 'gfx906' ISAs? I don't think these are mutually
> compatible.
>
Since I did not find any existing builtin to check the exact ISA, I
added all of them for consistency. Let me know if that looks good to you.
Thanks,
On 01/12/2022 11:10, Paul-Antoine Arras wrote:
> + if (TARGET_FIJI) \
> + builtin_define ("__FIJI__"); \
> + else if (TARGET_VEGA10) \
> + builtin_define ("__VEGA10__"); \
> + else if (TARGET_VEGA20) \
> + builtin_define ("__VEGA20__"); \
> + else if (TARGET_GFX908) \
> + builtin_define ("__GFX908__"); \
> + else if (TARGET_GFX90a) \
> + builtin_define ("__GFX90a__"); \
> + } while (0)
>
I don't think it makes sense to say __VEGA10__ when the user asked for
-march=gfx900.
This whole naming thing is a bit of a mess already, so I think we'd do
better to either keep the same names throughout or match what LLVM does
(since it got to these first).
Please use "__gfx900__" etc. (lower case).
I'm half tempted to do a global search and replace on the internal
names, but since they're not externally visible that would probably just
be making merge conflicts for the sake of it.
Thanks
Andrew
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.
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,
> 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?
Looks good to me, thanks!
Kwok
@@ -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
new file mode 100644
@@ -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" } } } } */