From mboxrd@z Thu Jan 1 00:00:00 1970 Return-Path: Received: by sourceware.org (Postfix, from userid 7899) id 6B9193864A3A; Tue, 6 Dec 2022 13:57:48 +0000 (GMT) DKIM-Filter: OpenDKIM Filter v2.11.0 sourceware.org 6B9193864A3A DKIM-Signature: v=1; a=rsa-sha256; c=relaxed/relaxed; d=gcc.gnu.org; s=default; t=1670335068; bh=7rtWKkHUWvzVwEH3z+n1TZOvCMlEggj4K9fs+GnoYOA=; h=From:To:Subject:Date:From; b=sSQRo3yaaRZesbN1WPBqey4uHI56iBOW8vZ/JDQ7C9NRmXGFr4z1FSqz8v4V62EOv Uwg8Z5KWIjIdrWo2tu5OZAgXZQS4yWdeAk2LLBwqL2LklwPgTQhWfnlJK+jCsM7Nk7 +ia9Uo9jM+mwyIz4FA3csvK6xyXf2bxEHMDr1S/Q= Content-Type: text/plain; charset="us-ascii" MIME-Version: 1.0 Content-Transfer-Encoding: 7bit From: Paul-Antoine Arras To: gcc-cvs@gcc.gnu.org Subject: [gcc/devel/omp/gcc-12] amdgcn: Support AMD-specific 'isa' and 'arch' traits in OpenMP context selectors X-Act-Checkin: gcc X-Git-Author: Paul-Antoine Arras X-Git-Refname: refs/heads/devel/omp/gcc-12 X-Git-Oldrev: a8db05dcf5819c7ddd5c71da40bb08a533551ae9 X-Git-Newrev: de21480cbed844883a02f063759710594c87589e Message-Id: <20221206135748.6B9193864A3A@sourceware.org> Date: Tue, 6 Dec 2022 13:57:48 +0000 (GMT) List-Id: https://gcc.gnu.org/g:de21480cbed844883a02f063759710594c87589e commit de21480cbed844883a02f063759710594c87589e Author: Paul-Antoine Arras Date: Wed Nov 30 14:52:55 2022 +0100 amdgcn: Support AMD-specific 'isa' and 'arch' traits in OpenMP context selectors Add libgomp support for 'amdgcn' as arch, and for each processor type (as passed to '-march') 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 '-march' values (as well as 'gfx803') as isa traits. * testsuite/libgomp.c-c++-common/metadirective-6.c: New test. Diff: --- libgomp/ChangeLog.omp | 6 +++ libgomp/config/gcn/selector.c | 24 ++++++++--- .../libgomp.c-c++-common/metadirective-6.c | 48 ++++++++++++++++++++++ 3 files changed, 73 insertions(+), 5 deletions(-) diff --git a/libgomp/ChangeLog.omp b/libgomp/ChangeLog.omp index 5a86870ad39..73618087b0d 100644 --- a/libgomp/ChangeLog.omp +++ b/libgomp/ChangeLog.omp @@ -1,3 +1,9 @@ +2022-12-06 Paul-Antoine Arras + + * config/gcn/selector.c (GOMP_evaluate_current_device): Recognise 'amdgcn' + as arch, and '-march' values (as well as 'gfx803') as isa traits. + * testsuite/libgomp.c-c++-common/metadirective-6.c: New test. + 2022-12-06 Tobias Burnus Backported from master: diff --git a/libgomp/config/gcn/selector.c b/libgomp/config/gcn/selector.c index 60793fc05d3..570bc1e8ae6 100644 --- a/libgomp/config/gcn/selector.c +++ b/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 __gfx803__ 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 __gfx900__ + if (strcmp (isa, "gfx900") == 0) + return true; +#endif + +#ifdef __gfx906__ + 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 a/libgomp/testsuite/libgomp.c-c++-common/metadirective-6.c b/libgomp/testsuite/libgomp.c-c++-common/metadirective-6.c new file mode 100644 index 00000000000..6d169001db1 --- /dev/null +++ b/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" } } } } */