From mboxrd@z Thu Jan 1 00:00:00 1970 Return-Path: Received: by sourceware.org (Postfix, from userid 7899) id B858A3858D28; Wed, 30 Nov 2022 11:38:06 +0000 (GMT) DKIM-Filter: OpenDKIM Filter v2.11.0 sourceware.org B858A3858D28 DKIM-Signature: v=1; a=rsa-sha256; c=relaxed/relaxed; d=gcc.gnu.org; s=default; t=1669808286; bh=L2gTXBl80Tu1kFYTnNIYetM7KS1YAofGHTV+rfJy4vc=; h=From:To:Subject:Date:From; b=ZJlJQwRw7yKjPckKTT/Z7bcdY8B+Z8VcHB8FUwpoaWalL8DwA0XXHx70LEDxUDrMf ozwR7FjXveOOPnU0SriFeezMlgKT32nnf+kb0WkPc3yBAvGhEO/uIU/1yirMFzsUx3 fLBPAEpTrGhIUxObueix4zXBiJo3CYndZtEAyoBs= 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' 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: ec5445902910ad709f33a0f13fcd314b3d4e3e87 X-Git-Newrev: b60c566317ae14916249544604bd7d6e46837be5 Message-Id: <20221130113806.B858A3858D28@sourceware.org> Date: Wed, 30 Nov 2022 11:38:06 +0000 (GMT) List-Id: https://gcc.gnu.org/g:b60c566317ae14916249544604bd7d6e46837be5 commit b60c566317ae14916249544604bd7d6e46837be5 Author: Paul-Antoine Arras Date: Tue Nov 29 16:22:07 2022 +0100 amdgcn: Support AMD-specific 'isa' traits in OpenMP context selectors Add support for gfx803 as an alias for fiji. Add test cases for all supported 'isa' values. gcc/ChangeLog: * config/gcn/gcn.cc (gcn_omp_device_kind_arch_isa): Add gfx803. * config/gcn/t-omp-device: Add gfx803. libgomp/ChangeLog: * testsuite/libgomp.c/declare-variant-4-fiji.c: New test. * testsuite/libgomp.c/declare-variant-4-gfx803.c: New test. * testsuite/libgomp.c/declare-variant-4-gfx900.c: New test. * testsuite/libgomp.c/declare-variant-4-gfx906.c: New test. * testsuite/libgomp.c/declare-variant-4-gfx908.c: New test. * testsuite/libgomp.c/declare-variant-4-gfx90a.c: New test. * testsuite/libgomp.c/declare-variant-4.h: New header file. (cherry picked from commit 1fd508744eccda9ad9c6d6fcce5b2ea9c568818d) Diff: --- gcc/ChangeLog.omp | 8 +++ gcc/config/gcn/gcn.cc | 2 +- gcc/config/gcn/t-omp-device | 2 +- libgomp/ChangeLog.omp | 13 +++++ .../testsuite/libgomp.c/declare-variant-4-fiji.c | 8 +++ .../testsuite/libgomp.c/declare-variant-4-gfx803.c | 7 +++ .../testsuite/libgomp.c/declare-variant-4-gfx900.c | 7 +++ .../testsuite/libgomp.c/declare-variant-4-gfx906.c | 7 +++ .../testsuite/libgomp.c/declare-variant-4-gfx908.c | 7 +++ .../testsuite/libgomp.c/declare-variant-4-gfx90a.c | 7 +++ libgomp/testsuite/libgomp.c/declare-variant-4.h | 63 ++++++++++++++++++++++ 11 files changed, 129 insertions(+), 2 deletions(-) diff --git a/gcc/ChangeLog.omp b/gcc/ChangeLog.omp index 48443438e39..b6204f7c760 100644 --- a/gcc/ChangeLog.omp +++ b/gcc/ChangeLog.omp @@ -1,3 +1,11 @@ +2022-11-30 Paul-Antoine Arras + + Backported from master: + 2022-11-30 Paul-Antoine Arras + + * config/gcn/gcn.cc (gcn_omp_device_kind_arch_isa): Add gfx803. + * config/gcn/t-omp-device: Add gfx803. + 2022-11-28 Tobias Burnus Backported from master: diff --git a/gcc/config/gcn/gcn.cc b/gcc/config/gcn/gcn.cc index 85f20a46d67..b9a18613e15 100644 --- a/gcc/config/gcn/gcn.cc +++ b/gcc/config/gcn/gcn.cc @@ -3016,7 +3016,7 @@ gcn_omp_device_kind_arch_isa (enum omp_device_kind_arch_isa trait, case omp_device_arch: return strcmp (name, "amdgcn") == 0 || strcmp (name, "gcn") == 0; case omp_device_isa: - if (strcmp (name, "fiji") == 0) + if (strcmp (name, "fiji") == 0 || strcmp (name, "gfx803") == 0) return gcn_arch == PROCESSOR_FIJI; if (strcmp (name, "gfx900") == 0) return gcn_arch == PROCESSOR_VEGA10; diff --git a/gcc/config/gcn/t-omp-device b/gcc/config/gcn/t-omp-device index 27d36db894b..538624f7ec7 100644 --- a/gcc/config/gcn/t-omp-device +++ b/gcc/config/gcn/t-omp-device @@ -1,4 +1,4 @@ omp-device-properties-gcn: $(srcdir)/config/gcn/gcn.cc echo kind: gpu > $@ echo arch: amdgcn gcn >> $@ - echo isa: fiji gfx900 gfx906 gfx908 gfx90a >> $@ + echo isa: fiji gfx803 gfx900 gfx906 gfx908 gfx90a >> $@ diff --git a/libgomp/ChangeLog.omp b/libgomp/ChangeLog.omp index 499f4ecd6f0..4985ae73858 100644 --- a/libgomp/ChangeLog.omp +++ b/libgomp/ChangeLog.omp @@ -1,3 +1,16 @@ +2022-11-30 Paul-Antoine Arras + + Backported from master: + 2022-11-30 Paul-Antoine Arras + + * testsuite/libgomp.c/declare-variant-4-fiji.c: New test. + * testsuite/libgomp.c/declare-variant-4-gfx803.c: New test. + * testsuite/libgomp.c/declare-variant-4-gfx900.c: New test. + * testsuite/libgomp.c/declare-variant-4-gfx906.c: New test. + * testsuite/libgomp.c/declare-variant-4-gfx908.c: New test. + * testsuite/libgomp.c/declare-variant-4-gfx90a.c: New test. + * testsuite/libgomp.c/declare-variant-4.h: New header file. + 2022-11-28 Tobias Burnus Backported from master: diff --git a/libgomp/testsuite/libgomp.c/declare-variant-4-fiji.c b/libgomp/testsuite/libgomp.c/declare-variant-4-fiji.c new file mode 100644 index 00000000000..ae2af1cc00c --- /dev/null +++ b/libgomp/testsuite/libgomp.c/declare-variant-4-fiji.c @@ -0,0 +1,8 @@ +/* { dg-do run { target { offload_target_amdgcn } } } */ +/* { dg-skip-if "fiji/gfx803 only" { ! amdgcn-*-* } { "*" } { "-foffload=-march=fiji" } } */ +/* { dg-additional-options "-foffload=-fdump-tree-optimized" } */ + +#define USE_FIJI_FOR_GFX803 +#include "declare-variant-4.h" + +/* { dg-final { scan-offload-tree-dump "= gfx803 \\(\\);" "optimized" } } */ diff --git a/libgomp/testsuite/libgomp.c/declare-variant-4-gfx803.c b/libgomp/testsuite/libgomp.c/declare-variant-4-gfx803.c new file mode 100644 index 00000000000..e0437a04d65 --- /dev/null +++ b/libgomp/testsuite/libgomp.c/declare-variant-4-gfx803.c @@ -0,0 +1,7 @@ +/* { dg-do run { target { offload_target_amdgcn } } } */ +/* { dg-skip-if "fiji/gfx803 only" { ! amdgcn-*-* } { "*" } { "-foffload=-march=fiji" } } */ +/* { dg-additional-options "-foffload=-fdump-tree-optimized" } */ + +#include "declare-variant-4.h" + +/* { dg-final { scan-offload-tree-dump "= gfx803 \\(\\);" "optimized" } } */ diff --git a/libgomp/testsuite/libgomp.c/declare-variant-4-gfx900.c b/libgomp/testsuite/libgomp.c/declare-variant-4-gfx900.c new file mode 100644 index 00000000000..8de03725dec --- /dev/null +++ b/libgomp/testsuite/libgomp.c/declare-variant-4-gfx900.c @@ -0,0 +1,7 @@ +/* { dg-do run { target { offload_target_amdgcn } } } */ +/* { dg-skip-if "gfx900 only" { ! amdgcn-*-* } { "*" } { "-foffload=-march=gfx900" } } */ +/* { dg-additional-options "-foffload=-fdump-tree-optimized" } */ + +#include "declare-variant-4.h" + +/* { dg-final { scan-offload-tree-dump "= gfx900 \\(\\);" "optimized" } } */ diff --git a/libgomp/testsuite/libgomp.c/declare-variant-4-gfx906.c b/libgomp/testsuite/libgomp.c/declare-variant-4-gfx906.c new file mode 100644 index 00000000000..be6f193ed3a --- /dev/null +++ b/libgomp/testsuite/libgomp.c/declare-variant-4-gfx906.c @@ -0,0 +1,7 @@ +/* { dg-do run { target { offload_target_amdgcn } } } */ +/* { dg-skip-if "gfx906 only" { ! amdgcn-*-* } { "*" } { "-foffload=-march=gfx906" } } */ +/* { dg-additional-options "-foffload=-fdump-tree-optimized" } */ + +#include "declare-variant-4.h" + +/* { dg-final { scan-offload-tree-dump "= gfx906 \\(\\);" "optimized" } } */ diff --git a/libgomp/testsuite/libgomp.c/declare-variant-4-gfx908.c b/libgomp/testsuite/libgomp.c/declare-variant-4-gfx908.c new file mode 100644 index 00000000000..311fad9074d --- /dev/null +++ b/libgomp/testsuite/libgomp.c/declare-variant-4-gfx908.c @@ -0,0 +1,7 @@ +/* { dg-do run { target { offload_target_amdgcn } } } */ +/* { dg-skip-if "gfx908 only" { ! amdgcn-*-* } { "*" } { "-foffload=-march=gfx908" } } */ +/* { dg-additional-options "-foffload=-fdump-tree-optimized" } */ + +#include "declare-variant-4.h" + +/* { dg-final { scan-offload-tree-dump "= gfx908 \\(\\);" "optimized" } } */ diff --git a/libgomp/testsuite/libgomp.c/declare-variant-4-gfx90a.c b/libgomp/testsuite/libgomp.c/declare-variant-4-gfx90a.c new file mode 100644 index 00000000000..96cc14ca0a3 --- /dev/null +++ b/libgomp/testsuite/libgomp.c/declare-variant-4-gfx90a.c @@ -0,0 +1,7 @@ +/* { dg-do run { target { offload_target_amdgcn } } } */ +/* { dg-skip-if "gfx90a only" { ! amdgcn-*-* } { "*" } { "-foffload=-march=gfx90a" } } */ +/* { dg-additional-options "-foffload=-fdump-tree-optimized" } */ + +#include "declare-variant-4.h" + +/* { dg-final { scan-offload-tree-dump "= gfx90a \\(\\);" "optimized" } } */ diff --git a/libgomp/testsuite/libgomp.c/declare-variant-4.h b/libgomp/testsuite/libgomp.c/declare-variant-4.h new file mode 100644 index 00000000000..2d7c1ef1a5a --- /dev/null +++ b/libgomp/testsuite/libgomp.c/declare-variant-4.h @@ -0,0 +1,63 @@ +#pragma omp declare target +int +gfx803 (void) +{ + return 0x803; +} + +int +gfx900 (void) +{ + return 0x900; +} + +int +gfx906 (void) +{ + return 0x906; +} + +int +gfx908 (void) +{ + return 0x908; +} + +int +gfx90a (void) +{ + return 0x90a; +} + +#ifdef USE_FIJI_FOR_GFX803 +#pragma omp declare variant(gfx803) match(device = {isa("fiji")}) +#else +#pragma omp declare variant(gfx803) match(device = {isa("gfx803")}) +#endif +#pragma omp declare variant(gfx900) match(device = {isa("gfx900")}) +#pragma omp declare variant(gfx906) match(device = {isa("gfx906")}) +#pragma omp declare variant(gfx908) match(device = {isa("gfx908")}) +#pragma omp declare variant(gfx90a) match(device = {isa("gfx90a")}) +int +f (void) +{ + return 0; +} + +#pragma omp end declare target + +int +main (void) +{ + int v = 0; + +#pragma omp target map(from : v) + v = f (); + + if (v == 0) + __builtin_abort (); + + __builtin_printf ("AMDGCN accelerator: gfx%x\n", v); + + return 0; +}