From f846292d2ce953a633fe400226277cf0cb0d6243 Mon Sep 17 00:00:00 2001 From: Paul-Antoine Arras 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