Hi Kwok,

On 30/11/2022 19:50, Kwok Cheung Yeung wrote:
Hello PA,

--- 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;

The logic here looks wrong to me - surely it should return false if arch is not 'gcn' AND it is not 'amdgcn'?

Sure. Fixed in revised patch.

@@ -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

Okay for gfx908 and gfx90a, but is there any way of distinguishing between 'gfx900' and 'gfx906' ISAs? I don't think these are mutually compatible.


Since I did not find any existing builtin to check the exact ISA, I added all of them for consistency. Let me know if that looks good to you.

Thanks,
--
PA
From f846292d2ce953a633fe400226277cf0cb0d6243 Mon Sep 17 00:00:00 2001
From: Paul-Antoine Arras <p...@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.
---
 gcc/config/gcn/gcn-opts.h                     |  6 +++
 gcc/config/gcn/gcn.h                          | 37 ++++++++------
 libgomp/config/gcn/selector.c                 | 24 ++++++++--
 .../libgomp.c-c++-common/metadirective-6.c    | 48 +++++++++++++++++++
 4 files changed, 96 insertions(+), 19 deletions(-)
 create mode 100644 libgomp/testsuite/libgomp.c-c++-common/metadirective-6.c

diff --git gcc/config/gcn/gcn-opts.h gcc/config/gcn/gcn-opts.h
index 07ddc79cda3..fb7e5d9a5e9 100644
--- gcc/config/gcn/gcn-opts.h
+++ gcc/config/gcn/gcn-opts.h
@@ -27,6 +27,12 @@ enum processor_type
   PROCESSOR_GFX90a
 };
 
+#define TARGET_FIJI (gcn_arch == PROCESSOR_FIJI)
+#define TARGET_VEGA10 (gcn_arch == PROCESSOR_VEGA10)
+#define TARGET_VEGA20 (gcn_arch == PROCESSOR_VEGA20)
+#define TARGET_GFX908 (gcn_arch == PROCESSOR_GFX908)
+#define TARGET_GFX90a (gcn_arch == PROCESSOR_GFX90a)
+
 /* Set in gcn_option_override.  */
 extern enum gcn_isa {
   ISA_UNKNOWN,
diff --git gcc/config/gcn/gcn.h gcc/config/gcn/gcn.h
index 38f7212db59..22a95ba6609 100644
--- gcc/config/gcn/gcn.h
+++ gcc/config/gcn/gcn.h
@@ -16,20 +16,29 @@
 
 #include "config/gcn/gcn-opts.h"
 
-#define TARGET_CPU_CPP_BUILTINS()      \
-  do                                   \
-    {                                  \
-      builtin_define ("__AMDGCN__");   \
-      if (TARGET_GCN3)                 \
-       builtin_define ("__GCN3__");    \
-      else if (TARGET_GCN5)            \
-       builtin_define ("__GCN5__");    \
-      else if (TARGET_CDNA1)           \
-       builtin_define ("__CDNA1__");   \
-      else if (TARGET_CDNA2)           \
-       builtin_define ("__CDNA2__");   \
-    }                                  \
-  while(0)
+#define TARGET_CPU_CPP_BUILTINS()                                              
\
+  do                                                                           
\
+    {                                                                          
\
+      builtin_define ("__AMDGCN__");                                           
\
+      if (TARGET_GCN3)                                                         
\
+       builtin_define ("__GCN3__");                                           \
+      else if (TARGET_GCN5)                                                    
\
+       builtin_define ("__GCN5__");                                           \
+      else if (TARGET_CDNA1)                                                   
\
+       builtin_define ("__CDNA1__");                                          \
+      else if (TARGET_CDNA2)                                                   
\
+       builtin_define ("__CDNA2__");                                          \
+      if (TARGET_FIJI)                                                         
\
+       builtin_define ("__FIJI__");                                           \
+      else if (TARGET_VEGA10)                                                  
\
+       builtin_define ("__VEGA10__");                                         \
+      else if (TARGET_VEGA20)                                                  
\
+       builtin_define ("__VEGA20__");                                         \
+      else if (TARGET_GFX908)                                                  
\
+       builtin_define ("__GFX908__");                                         \
+      else if (TARGET_GFX90a)                                                  
\
+       builtin_define ("__GFX90a__");                                         \
+  } while (0)
 
 /* Support for a compile-time default architecture and tuning.
    The rules are:
diff --git libgomp/config/gcn/selector.c libgomp/config/gcn/selector.c
index 60793fc05d3..b07bdb9b7c7 100644
--- libgomp/config/gcn/selector.c
+++ libgomp/config/gcn/selector.c
@@ -36,20 +36,34 @@ 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)
     return true;
 
-#ifdef __GCN3__
+#ifdef __FIJI__
   if (strcmp (isa, "fiji") == 0 || strcmp (isa, "gfx803") == 0)
     return true;
 #endif
 
-#ifdef __GCN5__
-  if (strcmp (isa, "gfx900") == 0 || strcmp (isa, "gfx906") != 0
-      || strcmp (isa, "gfx908") == 0)
+#ifdef __VEGA10__
+  if (strcmp (isa, "gfx900") == 0)
+    return true;
+#endif
+
+#ifdef __VEGA20__
+  if (strcmp (isa, "gfx906") != 0)
+    return true;
+#endif
+
+#ifdef __GFX908__
+  if (strcmp (isa, "gfx908") == 0)
+    return true;
+#endif
+
+#ifdef __GFX90a__
+  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

Reply via email to