public inbox for gcc-patches@gcc.gnu.org
 help / color / mirror / Atom feed
* [PATCH][OG12] amdgcn: Support AMD-specific 'isa' and 'arch' traits in OpenMP context selectors
@ 2022-11-30 15:32 Paul-Antoine Arras
  2022-11-30 18:50 ` Kwok Cheung Yeung
  0 siblings, 1 reply; 8+ messages in thread
From: Paul-Antoine Arras @ 2022-11-30 15:32 UTC (permalink / raw)
  To: gcc-patches

[-- Attachment #1: Type: text/plain, Size: 489 bytes --]

Hi all,

This patch adds or fixes support for various AMD 'isa' and 'arch' trait 
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 
'metadirective' construct.

This patch is closely related to 
https://gcc.gnu.org/r13-4403-g1fd508744eccda but cannot be committed to 
mainline because metadirectives and dynamic context selectors have not 
landed there yet.

Can this be committed to OG12?

Thanks,

[-- Attachment #2: 0001-amdgcn-Support-AMD-specific-isa-and-arch-traits-in-O.patch --]
[-- Type: text/plain, Size: 3953 bytes --]

From 88522107dd39ba3ff8465cf688fe4438fa3b77b4 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.
---
 libgomp/config/gcn/selector.c                 | 15 ++++--
 .../libgomp.c-c++-common/metadirective-6.c    | 48 +++++++++++++++++++
 2 files changed, 60 insertions(+), 3 deletions(-)
 create mode 100644 libgomp/testsuite/libgomp.c-c++-common/metadirective-6.c

diff --git libgomp/config/gcn/selector.c libgomp/config/gcn/selector.c
index 60793fc05d3..c948497c538 100644
--- 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;
 
   if (!isa)
@@ -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
 
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


^ permalink raw reply	[flat|nested] 8+ messages in thread
end of thread, other threads:[~2022-12-02 16:51 UTC | newest]

Thread overview: 8+ messages (download: mbox.gz / follow: Atom feed)
-- links below jump to the message on this page --
2022-11-30 15:32 [PATCH][OG12] amdgcn: Support AMD-specific 'isa' and 'arch' traits in OpenMP context selectors Paul-Antoine Arras
2022-11-30 18:50 ` Kwok Cheung Yeung
2022-12-01 11:10   ` Paul-Antoine Arras
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

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).