public inbox for gcc-cvs@sourceware.org
help / color / mirror / Atom feed
From: Paul-Antoine Arras <parras@gcc.gnu.org>
To: gcc-cvs@gcc.gnu.org
Subject: [gcc/devel/omp/gcc-12] amdgcn: Support AMD-specific 'isa' traits in OpenMP context selectors
Date: Wed, 30 Nov 2022 11:38:06 +0000 (GMT) [thread overview]
Message-ID: <20221130113806.B858A3858D28@sourceware.org> (raw)
https://gcc.gnu.org/g:b60c566317ae14916249544604bd7d6e46837be5
commit b60c566317ae14916249544604bd7d6e46837be5
Author: Paul-Antoine Arras <pa@codesourcery.com>
Date: Tue Nov 29 16:22:07 2022 +0100
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.
(cherry picked from commit 1fd508744eccda9ad9c6d6fcce5b2ea9c568818d)
Diff:
---
gcc/ChangeLog.omp | 8 +++
gcc/config/gcn/gcn.cc | 2 +-
gcc/config/gcn/t-omp-device | 2 +-
libgomp/ChangeLog.omp | 13 +++++
.../testsuite/libgomp.c/declare-variant-4-fiji.c | 8 +++
.../testsuite/libgomp.c/declare-variant-4-gfx803.c | 7 +++
.../testsuite/libgomp.c/declare-variant-4-gfx900.c | 7 +++
.../testsuite/libgomp.c/declare-variant-4-gfx906.c | 7 +++
.../testsuite/libgomp.c/declare-variant-4-gfx908.c | 7 +++
.../testsuite/libgomp.c/declare-variant-4-gfx90a.c | 7 +++
libgomp/testsuite/libgomp.c/declare-variant-4.h | 63 ++++++++++++++++++++++
11 files changed, 129 insertions(+), 2 deletions(-)
diff --git a/gcc/ChangeLog.omp b/gcc/ChangeLog.omp
index 48443438e39..b6204f7c760 100644
--- a/gcc/ChangeLog.omp
+++ b/gcc/ChangeLog.omp
@@ -1,3 +1,11 @@
+2022-11-30 Paul-Antoine Arras <pa@codesourcery.com>
+
+ Backported from master:
+ 2022-11-30 Paul-Antoine Arras <pa@codesourcery.com>
+
+ * config/gcn/gcn.cc (gcn_omp_device_kind_arch_isa): Add gfx803.
+ * config/gcn/t-omp-device: Add gfx803.
+
2022-11-28 Tobias Burnus <tobias@codesourcery.com>
Backported from master:
diff --git a/gcc/config/gcn/gcn.cc b/gcc/config/gcn/gcn.cc
index 85f20a46d67..b9a18613e15 100644
--- a/gcc/config/gcn/gcn.cc
+++ b/gcc/config/gcn/gcn.cc
@@ -3016,7 +3016,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 a/gcc/config/gcn/t-omp-device b/gcc/config/gcn/t-omp-device
index 27d36db894b..538624f7ec7 100644
--- a/gcc/config/gcn/t-omp-device
+++ b/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 a/libgomp/ChangeLog.omp b/libgomp/ChangeLog.omp
index 499f4ecd6f0..4985ae73858 100644
--- a/libgomp/ChangeLog.omp
+++ b/libgomp/ChangeLog.omp
@@ -1,3 +1,16 @@
+2022-11-30 Paul-Antoine Arras <pa@codesourcery.com>
+
+ Backported from master:
+ 2022-11-30 Paul-Antoine Arras <pa@codesourcery.com>
+
+ * 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.
+
2022-11-28 Tobias Burnus <tobias@codesourcery.com>
Backported from master:
diff --git a/libgomp/testsuite/libgomp.c/declare-variant-4-fiji.c b/libgomp/testsuite/libgomp.c/declare-variant-4-fiji.c
new file mode 100644
index 00000000000..ae2af1cc00c
--- /dev/null
+++ b/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 a/libgomp/testsuite/libgomp.c/declare-variant-4-gfx803.c b/libgomp/testsuite/libgomp.c/declare-variant-4-gfx803.c
new file mode 100644
index 00000000000..e0437a04d65
--- /dev/null
+++ b/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 a/libgomp/testsuite/libgomp.c/declare-variant-4-gfx900.c b/libgomp/testsuite/libgomp.c/declare-variant-4-gfx900.c
new file mode 100644
index 00000000000..8de03725dec
--- /dev/null
+++ b/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 a/libgomp/testsuite/libgomp.c/declare-variant-4-gfx906.c b/libgomp/testsuite/libgomp.c/declare-variant-4-gfx906.c
new file mode 100644
index 00000000000..be6f193ed3a
--- /dev/null
+++ b/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 a/libgomp/testsuite/libgomp.c/declare-variant-4-gfx908.c b/libgomp/testsuite/libgomp.c/declare-variant-4-gfx908.c
new file mode 100644
index 00000000000..311fad9074d
--- /dev/null
+++ b/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 a/libgomp/testsuite/libgomp.c/declare-variant-4-gfx90a.c b/libgomp/testsuite/libgomp.c/declare-variant-4-gfx90a.c
new file mode 100644
index 00000000000..96cc14ca0a3
--- /dev/null
+++ b/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 a/libgomp/testsuite/libgomp.c/declare-variant-4.h b/libgomp/testsuite/libgomp.c/declare-variant-4.h
new file mode 100644
index 00000000000..2d7c1ef1a5a
--- /dev/null
+++ b/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;
+}
reply other threads:[~2022-11-30 11:38 UTC|newest]
Thread overview: [no followups] expand[flat|nested] mbox.gz Atom feed
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=20221130113806.B858A3858D28@sourceware.org \
--to=parras@gcc.gnu.org \
--cc=gcc-cvs@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).