From: Paul-Antoine Arras <pa@codesourcery.com>
To: Kwok Cheung Yeung <kcyeung77@gmail.com>
Cc: gcc-patches <gcc-patches@gcc.gnu.org>
Subject: Re: [PATCH][OG12] amdgcn: Support AMD-specific 'isa' and 'arch' traits in OpenMP context selectors
Date: Thu, 1 Dec 2022 12:10:37 +0100 [thread overview]
Message-ID: <82a884ed-ea1d-5116-fedf-42de6e22e730@codesourcery.com> (raw)
In-Reply-To: <202ee18b-90c8-830b-82a5-b705a9ec5c73@gmail.com>
[-- 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
next prev parent reply other threads:[~2022-12-01 11:10 UTC|newest]
Thread overview: 8+ messages / expand[flat|nested] mbox.gz Atom feed top
2022-11-30 15:32 Paul-Antoine Arras
2022-11-30 18:50 ` Kwok Cheung Yeung
2022-12-01 11:10 ` Paul-Antoine Arras [this message]
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
Reply instructions:
You may reply publicly to this message via plain-text email
using any one of the following methods:
* Save the following mbox file, import it into your mail client,
and reply-to-all from there: mbox
Avoid top-posting and favor interleaved quoting:
https://en.wikipedia.org/wiki/Posting_style#Interleaved_style
* Reply using the --to, --cc, and --in-reply-to
switches of git-send-email(1):
git send-email \
--in-reply-to=82a884ed-ea1d-5116-fedf-42de6e22e730@codesourcery.com \
--to=pa@codesourcery.com \
--cc=gcc-patches@gcc.gnu.org \
--cc=kcyeung77@gmail.com \
/path/to/YOUR_REPLY
https://kernel.org/pub/software/scm/git/docs/git-send-email.html
* If your mail client supports setting the In-Reply-To header
via mailto: links, try the mailto: link
Be sure your reply has a Subject: header at the top and a blank line
before the message body.
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).