* [PATCH] amdgcn: Support AMD-specific 'isa' traits in OpenMP context selectors
@ 2022-11-29 15:56 Paul-Antoine Arras
2022-11-29 17:51 ` Andrew Stubbs
` (2 more replies)
0 siblings, 3 replies; 6+ messages in thread
From: Paul-Antoine Arras @ 2022-11-29 15:56 UTC (permalink / raw)
To: gcc-patches
[-- 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
^ permalink raw reply [flat|nested] 6+ messages in thread
* Re: [PATCH] amdgcn: Support AMD-specific 'isa' traits in OpenMP context selectors
2022-11-29 15:56 [PATCH] amdgcn: Support AMD-specific 'isa' traits in OpenMP context selectors Paul-Antoine Arras
@ 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
2023-11-30 15:15 ` Fix 'libgomp.c/declare-variant-4-*.c', add 'libgomp.c/declare-variant-4.c' (was: [PATCH] " Thomas Schwinge
2 siblings, 0 replies; 6+ messages in thread
From: Andrew Stubbs @ 2022-11-29 17:51 UTC (permalink / raw)
To: Paul-Antoine Arras, gcc-patches
On 29/11/2022 15:56, Paul-Antoine Arras wrote:
> 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?
OK
Andrew
^ permalink raw reply [flat|nested] 6+ messages in thread
* [Patch] libgomp.texi: List GCN's 'gfx803' under OpenMP Context Selectors (was: amdgcn: Support AMD-specific 'isa' traits in OpenMP context selectors)
2022-11-29 15:56 [PATCH] amdgcn: Support AMD-specific 'isa' traits in OpenMP context selectors Paul-Antoine Arras
2022-11-29 17:51 ` Andrew Stubbs
@ 2022-11-29 18:26 ` Tobias Burnus
2022-11-30 9:43 ` Andrew Stubbs
2023-11-30 15:15 ` Fix 'libgomp.c/declare-variant-4-*.c', add 'libgomp.c/declare-variant-4.c' (was: [PATCH] " Thomas Schwinge
2 siblings, 1 reply; 6+ messages in thread
From: Tobias Burnus @ 2022-11-29 18:26 UTC (permalink / raw)
To: Paul-Antoine Arras, gcc-patches, Andrew Stubbs
[-- Attachment #1: Type: text/plain, Size: 713 bytes --]
Hi PA, hi Andrew, hi Jakub, hi all,
On 29.11.22 16:56, Paul-Antoine Arras wrote:
> This patch adds support for 'gfx803' as an alias for 'fiji' in OpenMP
> context selectors, [...]
I think this should be documented somewhere. We have
https://gcc.gnu.org/onlinedocs/libgomp/OpenMP-Context-Selectors.html
For GCN and ISA, it refers to -march= and gfx803 is only a context
selector. Hence:
How about the attached patch?
Tobias
-----------------
Siemens Electronic Design Automation GmbH; Anschrift: Arnulfstraße 201, 80634 München; Gesellschaft mit beschränkter Haftung; Geschäftsführer: Thomas Heurung, Frank Thürauf; Sitz der Gesellschaft: München; Registergericht München, HRB 106955
[-- Attachment #2: gfx803.diff --]
[-- Type: text/x-patch, Size: 677 bytes --]
libgomp.texi: List GCN's 'gfx803' under OpenMP Context Selectors
libgomp/ChangeLog:
* libgomp.texi (OpenMP Context Selectors): Add 'gfx803' to gcn's isa.
diff --git a/libgomp/libgomp.texi b/libgomp/libgomp.texi
index 85cae742cd4..0066d41fdc5 100644
--- a/libgomp/libgomp.texi
+++ b/libgomp/libgomp.texi
@@ -4378,5 +4378,6 @@ offloading devices (it's not clear if they should be):
@item @code{amdgcn}, @code{gcn}
@tab @code{gpu}
- @tab See @code{-march=} in ``AMD GCN Options''
+ @tab See @code{-march=} in ``AMD GCN Options''@footnote{Additionally
+ supported is @code{gfx803} as an alias for @code{fiji}.}
@item @code{nvptx}
@tab @code{gpu}
^ permalink raw reply [flat|nested] 6+ messages in thread
* Re: [Patch] libgomp.texi: List GCN's 'gfx803' under OpenMP Context Selectors (was: amdgcn: Support AMD-specific 'isa' traits in OpenMP context selectors)
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
0 siblings, 1 reply; 6+ messages in thread
From: Andrew Stubbs @ 2022-11-30 9:43 UTC (permalink / raw)
To: Tobias Burnus, Paul-Antoine Arras, gcc-patches
On 29/11/2022 18:26, Tobias Burnus wrote:
> Hi PA, hi Andrew, hi Jakub, hi all,
>
> On 29.11.22 16:56, Paul-Antoine Arras wrote:
>> This patch adds support for 'gfx803' as an alias for 'fiji' in OpenMP
>> context selectors, [...]
>
> I think this should be documented somewhere. We have
> https://gcc.gnu.org/onlinedocs/libgomp/OpenMP-Context-Selectors.html
>
> For GCN and ISA, it refers to -march= and gfx803 is only a context
> selector. Hence:
>
> How about the attached patch?
The wording is a little odd.
How about "Additionally, gfx908 is supported as an alias for fiji"?
Andrew
^ permalink raw reply [flat|nested] 6+ messages in thread
* Re: [Patch] libgomp.texi: List GCN's 'gfx803' under OpenMP Context Selectors (was: amdgcn: Support AMD-specific 'isa' traits in OpenMP context selectors)
2022-11-30 9:43 ` Andrew Stubbs
@ 2022-11-30 10:27 ` Tobias Burnus
0 siblings, 0 replies; 6+ messages in thread
From: Tobias Burnus @ 2022-11-30 10:27 UTC (permalink / raw)
To: Andrew Stubbs, Paul-Antoine Arras, gcc-patches
On 30.11.22 10:43, Andrew Stubbs wrote:
> On 29/11/2022 18:26, Tobias Burnus wrote:
>> On 29.11.22 16:56, Paul-Antoine Arras wrote:
>>> This patch adds support for 'gfx803' as an alias for 'fiji' in OpenMP
>>> context selectors, [...]
PA committed that patch as
https://gcc.gnu.org/r13-4403-g1fd508744eccda9ad9c6d6fcce5b2ea9c568818d
(thanks!)
>> I think this should be documented somewhere. We have
>> https://gcc.gnu.org/onlinedocs/libgomp/OpenMP-Context-Selectors.html
> The wording is a little odd.
> How about "Additionally, gfx908 is supported as an alias for fiji"?
Committed with the suggested wording:
https://gcc.gnu.org/r13-4404-ge0b95c2e8b771b53876321a6a0a9497619af73cd
Thanks,
Tobias
PS: It does not help with finding a good wording if that's the last task
before calling it a day...
-----------------
Siemens Electronic Design Automation GmbH; Anschrift: Arnulfstraße 201, 80634 München; Gesellschaft mit beschränkter Haftung; Geschäftsführer: Thomas Heurung, Frank Thürauf; Sitz der Gesellschaft: München; Registergericht München, HRB 106955
^ permalink raw reply [flat|nested] 6+ messages in thread
* Fix 'libgomp.c/declare-variant-4-*.c', add 'libgomp.c/declare-variant-4.c' (was: [PATCH] amdgcn: Support AMD-specific 'isa' traits in OpenMP context selectors)
2022-11-29 15:56 [PATCH] amdgcn: Support AMD-specific 'isa' traits in OpenMP context selectors Paul-Antoine Arras
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
@ 2023-11-30 15:15 ` Thomas Schwinge
2 siblings, 0 replies; 6+ messages in thread
From: Thomas Schwinge @ 2023-11-30 15:15 UTC (permalink / raw)
To: Paul-Antoine Arras, gcc-patches; +Cc: Jakub Jelinek, Tobias Burnus
[-- Attachment #1: Type: text/plain, Size: 1340 bytes --]
Hi!
On 2022-11-29T16:56:21+0100, Paul-Antoine Arras <pa@codesourcery.com> wrote:
> [...] also adds test
> cases checking all supported AMD ISAs are properly recognised when used
> in a 'declare variant' construct.
\o/ Yay for test cases!
> --- /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" } } */
Etc. Unfortunately, I only ever see:
UNSUPPORTED: libgomp.c/declare-variant-4-fiji.c
UNSUPPORTED: libgomp.c/declare-variant-4-gfx803.c
UNSUPPORTED: libgomp.c/declare-variant-4-gfx900.c
UNSUPPORTED: libgomp.c/declare-variant-4-gfx906.c
UNSUPPORTED: libgomp.c/declare-variant-4-gfx908.c
UNSUPPORTED: libgomp.c/declare-variant-4-gfx90a.c
Pushed to master branch commit aae57a9e19ba5d2016fa8a88db6d7ff15a40ca35
"Fix 'libgomp.c/declare-variant-4-*.c', add 'libgomp.c/declare-variant-4.c'",
see attached.
Grüße
Thomas
-----------------
Siemens Electronic Design Automation GmbH; Anschrift: Arnulfstraße 201, 80634 München; Gesellschaft mit beschränkter Haftung; Geschäftsführer: Thomas Heurung, Frank Thürauf; Sitz der Gesellschaft: München; Registergericht München, HRB 106955
[-- Warning: decoded text below may be mangled, UTF-8 assumed --]
[-- Attachment #2: 0001-Fix-libgomp.c-declare-variant-4-.c-add-libgomp.c-dec.patch --]
[-- Type: text/x-diff, Size: 7591 bytes --]
From aae57a9e19ba5d2016fa8a88db6d7ff15a40ca35 Mon Sep 17 00:00:00 2001
From: Thomas Schwinge <thomas@codesourcery.com>
Date: Tue, 21 Nov 2023 16:36:48 +0100
Subject: [PATCH] Fix 'libgomp.c/declare-variant-4-*.c', add
'libgomp.c/declare-variant-4.c'
These test cases being 'dg-skip-if [...] { ! amdgcn-*-* }' meant they were
only ever considered for 'amdgcn-*-*' target -- that is, GCN *target*, not
GCN *offload target*, what is intended to be tested here, but for which they
thus always were UNSUPPORTED. Use the same style of 'dg-[...]' directives
as for the nvptx offloading test cases, 'libgomp.c/declare-variant-3*.c'.
Fix-up for commit 1fd508744eccda9ad9c6d6fcce5b2ea9c568818d
"amdgcn: Support AMD-specific 'isa' traits in OpenMP context selectors".
libgomp/
* testsuite/libgomp.c/declare-variant-4-fiji.c: Adjust.
* testsuite/libgomp.c/declare-variant-4-gfx803.c: Likewise.
* testsuite/libgomp.c/declare-variant-4-gfx900.c: Likewise.
* testsuite/libgomp.c/declare-variant-4-gfx906.c: Likewise.
* testsuite/libgomp.c/declare-variant-4-gfx908.c: Likewise.
* testsuite/libgomp.c/declare-variant-4-gfx90a.c: Likewise.
* testsuite/libgomp.c/declare-variant-4.h: Likewise.
* testsuite/libgomp.c/declare-variant-4.c: New.
---
libgomp/testsuite/libgomp.c/declare-variant-4-fiji.c | 5 +++--
libgomp/testsuite/libgomp.c/declare-variant-4-gfx803.c | 5 +++--
libgomp/testsuite/libgomp.c/declare-variant-4-gfx900.c | 5 +++--
libgomp/testsuite/libgomp.c/declare-variant-4-gfx906.c | 5 +++--
libgomp/testsuite/libgomp.c/declare-variant-4-gfx908.c | 5 +++--
libgomp/testsuite/libgomp.c/declare-variant-4-gfx90a.c | 5 +++--
libgomp/testsuite/libgomp.c/declare-variant-4.c | 8 ++++++++
libgomp/testsuite/libgomp.c/declare-variant-4.h | 5 +++++
8 files changed, 31 insertions(+), 12 deletions(-)
create mode 100644 libgomp/testsuite/libgomp.c/declare-variant-4.c
diff --git a/libgomp/testsuite/libgomp.c/declare-variant-4-fiji.c b/libgomp/testsuite/libgomp.c/declare-variant-4-fiji.c
index 8a4e0f4728c..a138fb092f8 100644
--- a/libgomp/testsuite/libgomp.c/declare-variant-4-fiji.c
+++ b/libgomp/testsuite/libgomp.c/declare-variant-4-fiji.c
@@ -1,5 +1,6 @@
-/* { dg-do run { target { offload_target_amdgcn } } } */
-/* { dg-skip-if "fiji/gfx803 only" { ! amdgcn-*-* } { "*" } { "-foffload=-march=fiji" } } */
+/* { dg-do link { target { offload_target_amdgcn } } } */
+/* { dg-additional-options -foffload=amdgcn-amdhsa } */
+/* { dg-additional-options -foffload=-march=fiji } */
/* { dg-additional-options "-foffload=-fdump-tree-optimized" } */
#define USE_FIJI_FOR_GFX803
diff --git a/libgomp/testsuite/libgomp.c/declare-variant-4-gfx803.c b/libgomp/testsuite/libgomp.c/declare-variant-4-gfx803.c
index 050d7c9dd79..03dffddac49 100644
--- a/libgomp/testsuite/libgomp.c/declare-variant-4-gfx803.c
+++ b/libgomp/testsuite/libgomp.c/declare-variant-4-gfx803.c
@@ -1,5 +1,6 @@
-/* { dg-do run { target { offload_target_amdgcn } } } */
-/* { dg-skip-if "fiji/gfx803 only" { ! amdgcn-*-* } { "*" } { "-foffload=-march=fiji" } } */
+/* { dg-do link { target { offload_target_amdgcn } } } */
+/* { dg-additional-options -foffload=amdgcn-amdhsa } */
+/* { dg-additional-options -foffload=-march=fiji } */
/* { dg-additional-options "-foffload=-fdump-tree-optimized" } */
#include "declare-variant-4.h"
diff --git a/libgomp/testsuite/libgomp.c/declare-variant-4-gfx900.c b/libgomp/testsuite/libgomp.c/declare-variant-4-gfx900.c
index 2eeb4a248c1..f3f5244c7bc 100644
--- a/libgomp/testsuite/libgomp.c/declare-variant-4-gfx900.c
+++ b/libgomp/testsuite/libgomp.c/declare-variant-4-gfx900.c
@@ -1,5 +1,6 @@
-/* { dg-do run { target { offload_target_amdgcn } } } */
-/* { dg-skip-if "gfx900 only" { ! amdgcn-*-* } { "*" } { "-foffload=-march=gfx900" } } */
+/* { dg-do link { target { offload_target_amdgcn } } } */
+/* { dg-additional-options -foffload=amdgcn-amdhsa } */
+/* { dg-additional-options -foffload=-march=gfx900 } */
/* { dg-additional-options "-foffload=-fdump-tree-optimized" } */
#include "declare-variant-4.h"
diff --git a/libgomp/testsuite/libgomp.c/declare-variant-4-gfx906.c b/libgomp/testsuite/libgomp.c/declare-variant-4-gfx906.c
index 73a670dcc2a..ac43388a59f 100644
--- a/libgomp/testsuite/libgomp.c/declare-variant-4-gfx906.c
+++ b/libgomp/testsuite/libgomp.c/declare-variant-4-gfx906.c
@@ -1,5 +1,6 @@
-/* { dg-do run { target { offload_target_amdgcn } } } */
-/* { dg-skip-if "gfx906 only" { ! amdgcn-*-* } { "*" } { "-foffload=-march=gfx906" } } */
+/* { dg-do link { target { offload_target_amdgcn } } } */
+/* { dg-additional-options -foffload=amdgcn-amdhsa } */
+/* { dg-additional-options -foffload=-march=gfx906 } */
/* { dg-additional-options "-foffload=-fdump-tree-optimized" } */
#include "declare-variant-4.h"
diff --git a/libgomp/testsuite/libgomp.c/declare-variant-4-gfx908.c b/libgomp/testsuite/libgomp.c/declare-variant-4-gfx908.c
index ead330f9f2c..f60741f202f 100644
--- a/libgomp/testsuite/libgomp.c/declare-variant-4-gfx908.c
+++ b/libgomp/testsuite/libgomp.c/declare-variant-4-gfx908.c
@@ -1,5 +1,6 @@
-/* { dg-do run { target { offload_target_amdgcn } } } */
-/* { dg-skip-if "gfx908 only" { ! amdgcn-*-* } { "*" } { "-foffload=-march=gfx908" } } */
+/* { dg-do link { target { offload_target_amdgcn } } } */
+/* { dg-additional-options -foffload=amdgcn-amdhsa } */
+/* { dg-additional-options -foffload=-march=gfx908 } */
/* { dg-additional-options "-foffload=-fdump-tree-optimized" } */
#include "declare-variant-4.h"
diff --git a/libgomp/testsuite/libgomp.c/declare-variant-4-gfx90a.c b/libgomp/testsuite/libgomp.c/declare-variant-4-gfx90a.c
index a9b2d62a49d..832d174e4a5 100644
--- a/libgomp/testsuite/libgomp.c/declare-variant-4-gfx90a.c
+++ b/libgomp/testsuite/libgomp.c/declare-variant-4-gfx90a.c
@@ -1,5 +1,6 @@
-/* { dg-do run { target { offload_target_amdgcn } } } */
-/* { dg-skip-if "gfx90a only" { ! amdgcn-*-* } { "*" } { "-foffload=-march=gfx90a" } } */
+/* { dg-do link { target { offload_target_amdgcn } } } */
+/* { dg-additional-options -foffload=amdgcn-amdhsa } */
+/* { dg-additional-options -foffload=-march=gfx90a } */
/* { dg-additional-options "-foffload=-fdump-tree-optimized" } */
#include "declare-variant-4.h"
diff --git a/libgomp/testsuite/libgomp.c/declare-variant-4.c b/libgomp/testsuite/libgomp.c/declare-variant-4.c
new file mode 100644
index 00000000000..3c8802a6081
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c/declare-variant-4.c
@@ -0,0 +1,8 @@
+/* { dg-additional-options -DOFFLOAD_DEVICE_GCN { target offload_device_gcn } } */
+/* { dg-additional-options {-fdump-tree-optimized -foffload-options=-fdump-tree-optimized} } */
+
+#include "declare-variant-4.h"
+
+/* { dg-final { scan-tree-dump "= f \\(\\);" "optimized" } }
+ { dg-final { only_for_offload_target amdgcn-amdhsa scan-offload-tree-dump "= gfx\[^ \]+ \\(\\);" "optimized" { target offload_target_amdgcn } } }
+ { dg-final { only_for_offload_target nvptx-none scan-offload-tree-dump "= f \\(\\);" "optimized" { target offload_target_nvptx } } } */
diff --git a/libgomp/testsuite/libgomp.c/declare-variant-4.h b/libgomp/testsuite/libgomp.c/declare-variant-4.h
index 47517b75ee7..a70352430c2 100644
--- a/libgomp/testsuite/libgomp.c/declare-variant-4.h
+++ b/libgomp/testsuite/libgomp.c/declare-variant-4.h
@@ -61,10 +61,15 @@ main (void)
#pragma omp target map(from : v)
v = f ();
+#ifdef OFFLOAD_DEVICE_GCN
if (v == 0)
__builtin_abort ();
__builtin_printf ("AMDGCN accelerator: gfx%x\n", v);
+#else
+ if (v != 0)
+ __builtin_abort ();
+#endif
return 0;
}
--
2.34.1
^ permalink raw reply [flat|nested] 6+ messages in thread
end of thread, other threads:[~2023-11-30 15:15 UTC | newest]
Thread overview: 6+ messages (download: mbox.gz / follow: Atom feed)
-- links below jump to the message on this page --
2022-11-29 15:56 [PATCH] amdgcn: Support AMD-specific 'isa' traits in OpenMP context selectors Paul-Antoine Arras
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
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).