public inbox for gcc-cvs@sourceware.org help / color / mirror / Atom feed
From: Paul-Antoine Arras <parras@gcc.gnu.org> To: gcc-cvs@gcc.gnu.org Subject: [gcc/devel/omp/gcc-12] amdgcn: Support AMD-specific 'isa' and 'arch' traits in OpenMP context selectors Date: Tue, 6 Dec 2022 13:57:48 +0000 (GMT) [thread overview] Message-ID: <20221206135748.6B9193864A3A@sourceware.org> (raw) https://gcc.gnu.org/g:de21480cbed844883a02f063759710594c87589e commit de21480cbed844883a02f063759710594c87589e Author: Paul-Antoine Arras <pa@codesourcery.com> 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 <pa@codesourcery.com> + + * 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 <tobias@codesourcery.com> 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" } } } } */
reply other threads:[~2022-12-06 13:57 UTC|newest] Thread overview: [no followups] expand[flat|nested] mbox.gz Atom feed
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=20221206135748.6B9193864A3A@sourceware.org \ --to=parras@gcc.gnu.org \ --cc=gcc-cvs@gcc.gnu.org \ /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: linkBe 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).