From: Thomas Schwinge <thomas@codesourcery.com>
To: Paul-Antoine Arras <pa@codesourcery.com>, <gcc-patches@gcc.gnu.org>
Cc: Jakub Jelinek <jakub@redhat.com>,
Tobias Burnus <tobias@codesourcery.com>
Subject: 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)
Date: Thu, 30 Nov 2023 16:15:04 +0100 [thread overview]
Message-ID: <8734wn2ryf.fsf@dem-tschwing-1.ger.mentorg.com> (raw)
In-Reply-To: <dd63de7d-171c-bc9b-a3c5-5a3254c1c8a2@codesourcery.com>
[-- 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
prev parent reply other threads:[~2023-11-30 15:15 UTC|newest]
Thread overview: 6+ messages / expand[flat|nested] mbox.gz Atom feed top
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 ` Thomas Schwinge [this message]
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=8734wn2ryf.fsf@dem-tschwing-1.ger.mentorg.com \
--to=thomas@codesourcery.com \
--cc=gcc-patches@gcc.gnu.org \
--cc=jakub@redhat.com \
--cc=pa@codesourcery.com \
--cc=tobias@codesourcery.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).