From: Paul-Antoine Arras <pa@codesourcery.com>
To: <gcc-patches@gcc.gnu.org>
Subject: [PATCH] amdgcn: Support AMD-specific 'isa' traits in OpenMP context selectors
Date: Tue, 29 Nov 2022 16:56:21 +0100 [thread overview]
Message-ID: <dd63de7d-171c-bc9b-a3c5-5a3254c1c8a2@codesourcery.com> (raw)
[-- Attachment #1: Type: text/plain, Size: 298 bytes --]
Hi all,
This patch adds support for 'gfx803' as an alias for 'fiji' in OpenMP
context 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 'declare variant' construct.
Is it OK for mainline?
Thanks,
--
PA
[-- Attachment #2: 0001-amdgcn-Support-AMD-specific-isa-traits-in-OpenMP-con.patch --]
[-- Type: text/plain, Size: 7863 bytes --]
From 2523122f7fff806aca7f7f03109668064969aa2d Mon Sep 17 00:00:00 2001
From: Paul-Antoine Arras <pa@codesourcery.com>
Date: Tue, 29 Nov 2022 16:22:07 +0100
Subject: [PATCH] amdgcn: Support AMD-specific 'isa' traits in OpenMP context
selectors
Add support for gfx803 as an alias for fiji.
Add test cases for all supported 'isa' values.
gcc/ChangeLog:
* config/gcn/gcn.cc (gcn_omp_device_kind_arch_isa): Add gfx803.
* config/gcn/t-omp-device: Add gfx803.
libgomp/ChangeLog:
* testsuite/libgomp.c/declare-variant-4-fiji.c: New test.
* testsuite/libgomp.c/declare-variant-4-gfx803.c: New test.
* testsuite/libgomp.c/declare-variant-4-gfx900.c: New test.
* testsuite/libgomp.c/declare-variant-4-gfx906.c: New test.
* testsuite/libgomp.c/declare-variant-4-gfx908.c: New test.
* testsuite/libgomp.c/declare-variant-4-gfx90a.c: New test.
* testsuite/libgomp.c/declare-variant-4.h: New header file.
---
gcc/config/gcn/gcn.cc | 2 +-
gcc/config/gcn/t-omp-device | 2 +-
.../libgomp.c/declare-variant-4-fiji.c | 8 +++
.../libgomp.c/declare-variant-4-gfx803.c | 7 +++
.../libgomp.c/declare-variant-4-gfx900.c | 7 +++
.../libgomp.c/declare-variant-4-gfx906.c | 7 +++
.../libgomp.c/declare-variant-4-gfx908.c | 7 +++
.../libgomp.c/declare-variant-4-gfx90a.c | 7 +++
.../testsuite/libgomp.c/declare-variant-4.h | 63 +++++++++++++++++++
9 files changed, 108 insertions(+), 2 deletions(-)
create mode 100644 libgomp/testsuite/libgomp.c/declare-variant-4-fiji.c
create mode 100644 libgomp/testsuite/libgomp.c/declare-variant-4-gfx803.c
create mode 100644 libgomp/testsuite/libgomp.c/declare-variant-4-gfx900.c
create mode 100644 libgomp/testsuite/libgomp.c/declare-variant-4-gfx906.c
create mode 100644 libgomp/testsuite/libgomp.c/declare-variant-4-gfx908.c
create mode 100644 libgomp/testsuite/libgomp.c/declare-variant-4-gfx90a.c
create mode 100644 libgomp/testsuite/libgomp.c/declare-variant-4.h
diff --git gcc/config/gcn/gcn.cc gcc/config/gcn/gcn.cc
index c74fa007a21..39e93aeaeef 100644
--- gcc/config/gcn/gcn.cc
+++ gcc/config/gcn/gcn.cc
@@ -2985,7 +2985,7 @@ gcn_omp_device_kind_arch_isa (enum omp_device_kind_arch_isa trait,
case omp_device_arch:
return strcmp (name, "amdgcn") == 0 || strcmp (name, "gcn") == 0;
case omp_device_isa:
- if (strcmp (name, "fiji") == 0)
+ if (strcmp (name, "fiji") == 0 || strcmp (name, "gfx803") == 0)
return gcn_arch == PROCESSOR_FIJI;
if (strcmp (name, "gfx900") == 0)
return gcn_arch == PROCESSOR_VEGA10;
diff --git gcc/config/gcn/t-omp-device gcc/config/gcn/t-omp-device
index 27d36db894b..538624f7ec7 100644
--- gcc/config/gcn/t-omp-device
+++ gcc/config/gcn/t-omp-device
@@ -1,4 +1,4 @@
omp-device-properties-gcn: $(srcdir)/config/gcn/gcn.cc
echo kind: gpu > $@
echo arch: amdgcn gcn >> $@
- echo isa: fiji gfx900 gfx906 gfx908 gfx90a >> $@
+ echo isa: fiji gfx803 gfx900 gfx906 gfx908 gfx90a >> $@
diff --git libgomp/testsuite/libgomp.c/declare-variant-4-fiji.c libgomp/testsuite/libgomp.c/declare-variant-4-fiji.c
new file mode 100644
index 00000000000..ae2af1cc00c
--- /dev/null
+++ libgomp/testsuite/libgomp.c/declare-variant-4-fiji.c
@@ -0,0 +1,8 @@
+/* { dg-do run { target { offload_target_amdgcn } } } */
+/* { dg-skip-if "fiji/gfx803 only" { ! amdgcn-*-* } { "*" } { "-foffload=-march=fiji" } } */
+/* { dg-additional-options "-foffload=-fdump-tree-optimized" } */
+
+#define USE_FIJI_FOR_GFX803
+#include "declare-variant-4.h"
+
+/* { dg-final { scan-offload-tree-dump "= gfx803 \\(\\);" "optimized" } } */
diff --git libgomp/testsuite/libgomp.c/declare-variant-4-gfx803.c libgomp/testsuite/libgomp.c/declare-variant-4-gfx803.c
new file mode 100644
index 00000000000..e0437a04d65
--- /dev/null
+++ libgomp/testsuite/libgomp.c/declare-variant-4-gfx803.c
@@ -0,0 +1,7 @@
+/* { dg-do run { target { offload_target_amdgcn } } } */
+/* { dg-skip-if "fiji/gfx803 only" { ! amdgcn-*-* } { "*" } { "-foffload=-march=fiji" } } */
+/* { dg-additional-options "-foffload=-fdump-tree-optimized" } */
+
+#include "declare-variant-4.h"
+
+/* { dg-final { scan-offload-tree-dump "= gfx803 \\(\\);" "optimized" } } */
diff --git libgomp/testsuite/libgomp.c/declare-variant-4-gfx900.c libgomp/testsuite/libgomp.c/declare-variant-4-gfx900.c
new file mode 100644
index 00000000000..8de03725dec
--- /dev/null
+++ libgomp/testsuite/libgomp.c/declare-variant-4-gfx900.c
@@ -0,0 +1,7 @@
+/* { dg-do run { target { offload_target_amdgcn } } } */
+/* { dg-skip-if "gfx900 only" { ! amdgcn-*-* } { "*" } { "-foffload=-march=gfx900" } } */
+/* { dg-additional-options "-foffload=-fdump-tree-optimized" } */
+
+#include "declare-variant-4.h"
+
+/* { dg-final { scan-offload-tree-dump "= gfx900 \\(\\);" "optimized" } } */
diff --git libgomp/testsuite/libgomp.c/declare-variant-4-gfx906.c libgomp/testsuite/libgomp.c/declare-variant-4-gfx906.c
new file mode 100644
index 00000000000..be6f193ed3a
--- /dev/null
+++ libgomp/testsuite/libgomp.c/declare-variant-4-gfx906.c
@@ -0,0 +1,7 @@
+/* { dg-do run { target { offload_target_amdgcn } } } */
+/* { dg-skip-if "gfx906 only" { ! amdgcn-*-* } { "*" } { "-foffload=-march=gfx906" } } */
+/* { dg-additional-options "-foffload=-fdump-tree-optimized" } */
+
+#include "declare-variant-4.h"
+
+/* { dg-final { scan-offload-tree-dump "= gfx906 \\(\\);" "optimized" } } */
diff --git libgomp/testsuite/libgomp.c/declare-variant-4-gfx908.c libgomp/testsuite/libgomp.c/declare-variant-4-gfx908.c
new file mode 100644
index 00000000000..311fad9074d
--- /dev/null
+++ libgomp/testsuite/libgomp.c/declare-variant-4-gfx908.c
@@ -0,0 +1,7 @@
+/* { dg-do run { target { offload_target_amdgcn } } } */
+/* { dg-skip-if "gfx908 only" { ! amdgcn-*-* } { "*" } { "-foffload=-march=gfx908" } } */
+/* { dg-additional-options "-foffload=-fdump-tree-optimized" } */
+
+#include "declare-variant-4.h"
+
+/* { dg-final { scan-offload-tree-dump "= gfx908 \\(\\);" "optimized" } } */
diff --git libgomp/testsuite/libgomp.c/declare-variant-4-gfx90a.c libgomp/testsuite/libgomp.c/declare-variant-4-gfx90a.c
new file mode 100644
index 00000000000..96cc14ca0a3
--- /dev/null
+++ libgomp/testsuite/libgomp.c/declare-variant-4-gfx90a.c
@@ -0,0 +1,7 @@
+/* { dg-do run { target { offload_target_amdgcn } } } */
+/* { dg-skip-if "gfx90a only" { ! amdgcn-*-* } { "*" } { "-foffload=-march=gfx90a" } } */
+/* { dg-additional-options "-foffload=-fdump-tree-optimized" } */
+
+#include "declare-variant-4.h"
+
+/* { dg-final { scan-offload-tree-dump "= gfx90a \\(\\);" "optimized" } } */
diff --git libgomp/testsuite/libgomp.c/declare-variant-4.h libgomp/testsuite/libgomp.c/declare-variant-4.h
new file mode 100644
index 00000000000..2d7c1ef1a5a
--- /dev/null
+++ libgomp/testsuite/libgomp.c/declare-variant-4.h
@@ -0,0 +1,63 @@
+#pragma omp declare target
+int
+gfx803 (void)
+{
+ return 0x803;
+}
+
+int
+gfx900 (void)
+{
+ return 0x900;
+}
+
+int
+gfx906 (void)
+{
+ return 0x906;
+}
+
+int
+gfx908 (void)
+{
+ return 0x908;
+}
+
+int
+gfx90a (void)
+{
+ return 0x90a;
+}
+
+#ifdef USE_FIJI_FOR_GFX803
+#pragma omp declare variant(gfx803) match(device = {isa("fiji")})
+#else
+#pragma omp declare variant(gfx803) match(device = {isa("gfx803")})
+#endif
+#pragma omp declare variant(gfx900) match(device = {isa("gfx900")})
+#pragma omp declare variant(gfx906) match(device = {isa("gfx906")})
+#pragma omp declare variant(gfx908) match(device = {isa("gfx908")})
+#pragma omp declare variant(gfx90a) match(device = {isa("gfx90a")})
+int
+f (void)
+{
+ return 0;
+}
+
+#pragma omp end declare target
+
+int
+main (void)
+{
+ int v = 0;
+
+#pragma omp target map(from : v)
+ v = f ();
+
+ if (v == 0)
+ __builtin_abort ();
+
+ __builtin_printf ("AMDGCN accelerator: gfx%x\n", v);
+
+ return 0;
+}
--
2.31.1
next reply other threads:[~2022-11-29 15:56 UTC|newest]
Thread overview: 6+ messages / expand[flat|nested] mbox.gz Atom feed top
2022-11-29 15:56 Paul-Antoine Arras [this message]
2022-11-29 17:51 ` Andrew Stubbs
2022-11-29 18:26 ` [Patch] libgomp.texi: List GCN's 'gfx803' under OpenMP Context Selectors (was: amdgcn: Support AMD-specific 'isa' traits in OpenMP context selectors) Tobias Burnus
2022-11-30 9:43 ` Andrew Stubbs
2022-11-30 10:27 ` Tobias Burnus
2023-11-30 15:15 ` Fix 'libgomp.c/declare-variant-4-*.c', add 'libgomp.c/declare-variant-4.c' (was: [PATCH] " Thomas Schwinge
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=dd63de7d-171c-bc9b-a3c5-5a3254c1c8a2@codesourcery.com \
--to=pa@codesourcery.com \
--cc=gcc-patches@gcc.gnu.org \
/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).