* [PATCH][OG12] amdgcn: Support AMD-specific 'isa' and 'arch' traits in OpenMP context selectors
@ 2022-11-30 15:32 Paul-Antoine Arras
2022-11-30 18:50 ` Kwok Cheung Yeung
0 siblings, 1 reply; 8+ messages in thread
From: Paul-Antoine Arras @ 2022-11-30 15:32 UTC (permalink / raw)
To: gcc-patches
[-- 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
^ permalink raw reply [flat|nested] 8+ messages in thread
* Re: [PATCH][OG12] amdgcn: Support AMD-specific 'isa' and 'arch' traits in OpenMP context selectors
2022-11-30 15:32 [PATCH][OG12] amdgcn: Support AMD-specific 'isa' and 'arch' traits in OpenMP context selectors Paul-Antoine Arras
@ 2022-11-30 18:50 ` Kwok Cheung Yeung
2022-12-01 11:10 ` Paul-Antoine Arras
0 siblings, 1 reply; 8+ messages in thread
From: Kwok Cheung Yeung @ 2022-11-30 18:50 UTC (permalink / raw)
To: pa; +Cc: gcc-patches
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
^ permalink raw reply [flat|nested] 8+ messages in thread
* Re: [PATCH][OG12] amdgcn: Support AMD-specific 'isa' and 'arch' traits in OpenMP context selectors
2022-11-30 18:50 ` Kwok Cheung Yeung
@ 2022-12-01 11:10 ` Paul-Antoine Arras
2022-12-01 12:45 ` Andrew Stubbs
0 siblings, 1 reply; 8+ messages in thread
From: Paul-Antoine Arras @ 2022-12-01 11:10 UTC (permalink / raw)
To: Kwok Cheung Yeung; +Cc: gcc-patches
[-- Attachment #1: Type: text/plain, Size: 1493 bytes --]
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,
--
PA
[-- Attachment #2: 0001-amdgcn-Support-AMD-specific-isa-and-arch-traits-in-O.patch --]
[-- Type: text/plain, Size: 7203 bytes --]
From f846292d2ce953a633fe400226277cf0cb0d6243 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.
---
gcc/config/gcn/gcn-opts.h | 6 +++
gcc/config/gcn/gcn.h | 37 ++++++++------
libgomp/config/gcn/selector.c | 24 ++++++++--
.../libgomp.c-c++-common/metadirective-6.c | 48 +++++++++++++++++++
4 files changed, 96 insertions(+), 19 deletions(-)
create mode 100644 libgomp/testsuite/libgomp.c-c++-common/metadirective-6.c
diff --git gcc/config/gcn/gcn-opts.h gcc/config/gcn/gcn-opts.h
index 07ddc79cda3..fb7e5d9a5e9 100644
--- gcc/config/gcn/gcn-opts.h
+++ gcc/config/gcn/gcn-opts.h
@@ -27,6 +27,12 @@ enum processor_type
PROCESSOR_GFX90a
};
+#define TARGET_FIJI (gcn_arch == PROCESSOR_FIJI)
+#define TARGET_VEGA10 (gcn_arch == PROCESSOR_VEGA10)
+#define TARGET_VEGA20 (gcn_arch == PROCESSOR_VEGA20)
+#define TARGET_GFX908 (gcn_arch == PROCESSOR_GFX908)
+#define TARGET_GFX90a (gcn_arch == PROCESSOR_GFX90a)
+
/* Set in gcn_option_override. */
extern enum gcn_isa {
ISA_UNKNOWN,
diff --git gcc/config/gcn/gcn.h gcc/config/gcn/gcn.h
index 38f7212db59..22a95ba6609 100644
--- gcc/config/gcn/gcn.h
+++ gcc/config/gcn/gcn.h
@@ -16,20 +16,29 @@
#include "config/gcn/gcn-opts.h"
-#define TARGET_CPU_CPP_BUILTINS() \
- do \
- { \
- builtin_define ("__AMDGCN__"); \
- if (TARGET_GCN3) \
- builtin_define ("__GCN3__"); \
- else if (TARGET_GCN5) \
- builtin_define ("__GCN5__"); \
- else if (TARGET_CDNA1) \
- builtin_define ("__CDNA1__"); \
- else if (TARGET_CDNA2) \
- builtin_define ("__CDNA2__"); \
- } \
- while(0)
+#define TARGET_CPU_CPP_BUILTINS() \
+ do \
+ { \
+ builtin_define ("__AMDGCN__"); \
+ if (TARGET_GCN3) \
+ builtin_define ("__GCN3__"); \
+ else if (TARGET_GCN5) \
+ builtin_define ("__GCN5__"); \
+ else if (TARGET_CDNA1) \
+ builtin_define ("__CDNA1__"); \
+ else if (TARGET_CDNA2) \
+ builtin_define ("__CDNA2__"); \
+ 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)
/* Support for a compile-time default architecture and tuning.
The rules are:
diff --git libgomp/config/gcn/selector.c libgomp/config/gcn/selector.c
index 60793fc05d3..b07bdb9b7c7 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 __FIJI__
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 __VEGA10__
+ if (strcmp (isa, "gfx900") == 0)
+ return true;
+#endif
+
+#ifdef __VEGA20__
+ 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
^ permalink raw reply [flat|nested] 8+ messages in thread
* Re: [PATCH][OG12] amdgcn: Support AMD-specific 'isa' and 'arch' traits in OpenMP context selectors
2022-12-01 11:10 ` Paul-Antoine Arras
@ 2022-12-01 12:45 ` Andrew Stubbs
2022-12-01 14:35 ` [PATCH] amdgcn: Add preprocessor builtins for every processor type Paul-Antoine Arras
2022-12-01 15:48 ` [PATCH][OG12] amdgcn: Support AMD-specific 'isa' and 'arch' traits in OpenMP context selectors Paul-Antoine Arras
0 siblings, 2 replies; 8+ messages in thread
From: Andrew Stubbs @ 2022-12-01 12:45 UTC (permalink / raw)
To: Paul-Antoine Arras, Kwok Cheung Yeung; +Cc: gcc-patches
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.
^ permalink raw reply [flat|nested] 8+ messages in thread
* [PATCH] amdgcn: Add preprocessor builtins for every processor type
2022-12-01 12:45 ` Andrew Stubbs
@ 2022-12-01 14:35 ` Paul-Antoine Arras
2022-12-01 14:42 ` Andrew Stubbs
2022-12-01 15:48 ` [PATCH][OG12] amdgcn: Support AMD-specific 'isa' and 'arch' traits in OpenMP context selectors Paul-Antoine Arras
1 sibling, 1 reply; 8+ messages in thread
From: Paul-Antoine Arras @ 2022-12-01 14:35 UTC (permalink / raw)
To: Andrew Stubbs, Kwok Cheung Yeung; +Cc: gcc-patches
[-- Attachment #1: Type: text/plain, Size: 2272 bytes --]
Hi Andrew, all,
On 01/12/2022 13:45, Andrew Stubbs wrote:
> 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).
>
> [...]
>
> 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.
I believe this patch addresses your comments regarding the GCN bits.
The new builtins are consistent with the LLVM naming convention (lower
case, canonical name). For gfx803, I also kept '__fiji__' to be
consistent with -march=fiji.
Is it OK for mainline?
Thanks,
--
PA
[-- Attachment #2: 0001-amdgcn-Add-preprocessor-builtins-for-every-processor.patch --]
[-- Type: text/plain, Size: 3995 bytes --]
From 238e8e131741fc962fe87482d1e9a6eb1252c75c Mon Sep 17 00:00:00 2001
From: Paul-Antoine Arras <pa@codesourcery.com>
Date: Thu, 1 Dec 2022 15:09:54 +0100
Subject: [PATCH] amdgcn: Add preprocessor builtins for every processor type
Provide a specific builtin for each possible value of '-march'.
gcc/ChangeLog:
* config/gcn/gcn-opts.h (TARGET_FIJI): -march=fiji.
(TARGET_VEGA10): -march=gfx900.
(TARGET_VEGA20): -march=gfx906.
(TARGET_GFX908): -march=gfx908.
(TARGET_GFX90a): -march=gfx90a.
* config/gcn/gcn.h (TARGET_CPU_CPP_BUILTINS): Define a builtin that uniquely maps to '-march'.
---
gcc/config/gcn/gcn-opts.h | 6 ++++++
gcc/config/gcn/gcn.h | 40 +++++++++++++++++++++++++--------------
2 files changed, 32 insertions(+), 14 deletions(-)
diff --git gcc/config/gcn/gcn-opts.h gcc/config/gcn/gcn-opts.h
index b62dfb45f59..b54eae79faf 100644
--- gcc/config/gcn/gcn-opts.h
+++ gcc/config/gcn/gcn-opts.h
@@ -27,6 +27,12 @@ enum processor_type
PROCESSOR_GFX90a
};
+#define TARGET_FIJI (gcn_arch == PROCESSOR_FIJI)
+#define TARGET_VEGA10 (gcn_arch == PROCESSOR_VEGA10)
+#define TARGET_VEGA20 (gcn_arch == PROCESSOR_VEGA20)
+#define TARGET_GFX908 (gcn_arch == PROCESSOR_GFX908)
+#define TARGET_GFX90a (gcn_arch == PROCESSOR_GFX90a)
+
/* Set in gcn_option_override. */
extern enum gcn_isa {
ISA_UNKNOWN,
diff --git gcc/config/gcn/gcn.h gcc/config/gcn/gcn.h
index 38f7212db59..1cc5981d904 100644
--- gcc/config/gcn/gcn.h
+++ gcc/config/gcn/gcn.h
@@ -16,20 +16,32 @@
#include "config/gcn/gcn-opts.h"
-#define TARGET_CPU_CPP_BUILTINS() \
- do \
- { \
- builtin_define ("__AMDGCN__"); \
- if (TARGET_GCN3) \
- builtin_define ("__GCN3__"); \
- else if (TARGET_GCN5) \
- builtin_define ("__GCN5__"); \
- else if (TARGET_CDNA1) \
- builtin_define ("__CDNA1__"); \
- else if (TARGET_CDNA2) \
- builtin_define ("__CDNA2__"); \
- } \
- while(0)
+#define TARGET_CPU_CPP_BUILTINS() \
+ do \
+ { \
+ builtin_define ("__AMDGCN__"); \
+ if (TARGET_GCN3) \
+ builtin_define ("__GCN3__"); \
+ else if (TARGET_GCN5) \
+ builtin_define ("__GCN5__"); \
+ else if (TARGET_CDNA1) \
+ builtin_define ("__CDNA1__"); \
+ else if (TARGET_CDNA2) \
+ builtin_define ("__CDNA2__"); \
+ if (TARGET_FIJI) \
+ { \
+ builtin_define ("__fiji__"); \
+ builtin_define ("__gfx803__"); \
+ } \
+ else if (TARGET_VEGA10) \
+ builtin_define ("__gfx900__"); \
+ else if (TARGET_VEGA20) \
+ builtin_define ("__gfx906__"); \
+ else if (TARGET_GFX908) \
+ builtin_define ("__gfx908__"); \
+ else if (TARGET_GFX90a) \
+ builtin_define ("__gfx90a__"); \
+ } while (0)
/* Support for a compile-time default architecture and tuning.
The rules are:
--
2.31.1
^ permalink raw reply [flat|nested] 8+ messages in thread
* Re: [PATCH] amdgcn: Add preprocessor builtins for every processor type
2022-12-01 14:35 ` [PATCH] amdgcn: Add preprocessor builtins for every processor type Paul-Antoine Arras
@ 2022-12-01 14:42 ` Andrew Stubbs
0 siblings, 0 replies; 8+ messages in thread
From: Andrew Stubbs @ 2022-12-01 14:42 UTC (permalink / raw)
To: Paul-Antoine Arras, Kwok Cheung Yeung; +Cc: gcc-patches
On 01/12/2022 14:35, Paul-Antoine Arras wrote:
> I believe this patch addresses your comments regarding the GCN bits.
>
> The new builtins are consistent with the LLVM naming convention (lower
> case, canonical name). For gfx803, I also kept '__fiji__' to be
> consistent with -march=fiji.
>
> Is it OK for mainline?
You need to wrap the long line in the changelog (I'm not sure it'll even
let you push it like that), but otherwise it looks fine.
OK.
Andrew
^ permalink raw reply [flat|nested] 8+ messages in thread
* Re: [PATCH][OG12] amdgcn: Support AMD-specific 'isa' and 'arch' traits in OpenMP context selectors
2022-12-01 12:45 ` Andrew Stubbs
2022-12-01 14:35 ` [PATCH] amdgcn: Add preprocessor builtins for every processor type Paul-Antoine Arras
@ 2022-12-01 15:48 ` Paul-Antoine Arras
2022-12-02 16:51 ` Kwok Cheung Yeung
1 sibling, 1 reply; 8+ messages in thread
From: Paul-Antoine Arras @ 2022-12-01 15:48 UTC (permalink / raw)
To: Andrew Stubbs, Kwok Cheung Yeung; +Cc: gcc-patches
[-- Attachment #1: Type: text/plain, Size: 420 bytes --]
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
[-- Attachment #2: 0001-amdgcn-Support-AMD-specific-isa-and-arch-traits-in-O.patch --]
[-- Type: text/plain, Size: 4686 bytes --]
From 494a815af459b13da6fe9bf5a84b94d4b1f94915 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 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 <pa@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 <tobias@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
^ permalink raw reply [flat|nested] 8+ messages in thread
* Re: [PATCH][OG12] amdgcn: Support AMD-specific 'isa' and 'arch' traits in OpenMP context selectors
2022-12-01 15:48 ` [PATCH][OG12] amdgcn: Support AMD-specific 'isa' and 'arch' traits in OpenMP context selectors Paul-Antoine Arras
@ 2022-12-02 16:51 ` Kwok Cheung Yeung
0 siblings, 0 replies; 8+ messages in thread
From: Kwok Cheung Yeung @ 2022-12-02 16:51 UTC (permalink / raw)
To: pa; +Cc: gcc-patches
> 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
^ permalink raw reply [flat|nested] 8+ messages in thread
end of thread, other threads:[~2022-12-02 16:51 UTC | newest]
Thread overview: 8+ messages (download: mbox.gz / follow: Atom feed)
-- links below jump to the message on this page --
2022-11-30 15:32 [PATCH][OG12] amdgcn: Support AMD-specific 'isa' and 'arch' traits in OpenMP context selectors Paul-Antoine Arras
2022-11-30 18:50 ` Kwok Cheung Yeung
2022-12-01 11:10 ` Paul-Antoine Arras
2022-12-01 12:45 ` Andrew Stubbs
2022-12-01 14:35 ` [PATCH] amdgcn: Add preprocessor builtins for every processor type Paul-Antoine Arras
2022-12-01 14:42 ` Andrew Stubbs
2022-12-01 15:48 ` [PATCH][OG12] amdgcn: Support AMD-specific 'isa' and 'arch' traits in OpenMP context selectors Paul-Antoine Arras
2022-12-02 16:51 ` Kwok Cheung Yeung
This is a public inbox, see mirroring instructions
for how to clone and mirror all data and code used for this inbox;
as well as URLs for read-only IMAP folder(s) and NNTP newsgroup(s).