Hi! On 2020-01-20T16:53:49+0000, Andrew Stubbs <a...@codesourcery.com> wrote: > I've committed this testsuite-only patch to fix some test cases that > need GCN-specific settings in order to pass.
> * testsuite/libgomp.oacc-c-c++-common/parallel-dims.c > (acc_gang): Recognise acc_device_radeon. > (acc_worker): Likewise. > (acc_vector): Likewise. > --- a/libgomp/testsuite/libgomp.oacc-c-c++-common/parallel-dims.c > +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/parallel-dims.c > @@ -14,7 +14,8 @@ static unsigned int __attribute__ ((optimize ("O2"))) > acc_gang () > { > if (acc_on_device ((int) acc_device_host)) > return 0; > - else if (acc_on_device ((int) acc_device_nvidia)) > + else if (acc_on_device ((int) acc_device_nvidia) > + || acc_on_device ((int) acc_device_radeon)) > return __builtin_goacc_parlevel_id (GOMP_DIM_GANG); > else > __builtin_abort (); > @@ -25,7 +26,8 @@ static unsigned int __attribute__ ((optimize ("O2"))) > acc_worker () > { > if (acc_on_device ((int) acc_device_host)) > return 0; > - else if (acc_on_device ((int) acc_device_nvidia)) > + else if (acc_on_device ((int) acc_device_nvidia) > + || acc_on_device ((int) acc_device_radeon)) > return __builtin_goacc_parlevel_id (GOMP_DIM_WORKER); > else > __builtin_abort (); > @@ -36,7 +38,8 @@ static unsigned int __attribute__ ((optimize ("O2"))) > acc_vector () > { > if (acc_on_device ((int) acc_device_host)) > return 0; > - else if (acc_on_device ((int) acc_device_nvidia)) > + else if (acc_on_device ((int) acc_device_nvidia) > + || acc_on_device ((int) acc_device_radeon)) > return __builtin_goacc_parlevel_id (GOMP_DIM_VECTOR); > else > __builtin_abort (); Similar changes are necessary for 'libgomp.oacc-fortran/parallel-dims.f90', too -- and actually all that can be simplified; pushed "Fix 'libgomp.oacc-fortran/parallel-dims.f90' for 'acc_device_radeon'" to master branch in commit 32099c0d24adb93a031e0301ffd77b065b6f5472, see attached. > (main): Set expectations for amdgcn. > @@ -282,6 +285,12 @@ int main () > /* The GCC nvptx back end enforces num_workers (32). */ > workers_actual = 32; > } > + else if (acc_on_device (acc_device_radeon)) > + { > + /* The GCC GCN back end is limited to num_workers (16). > + Temporarily set this to 1 until multiple workers are permitted. */ > + workers_actual = 1; // 16; > + } ACK; working on that. > @@ -328,6 +337,11 @@ int main () > /* We're actually executing with num_workers (32). */ > /* workers_actual = 32; */ > } > + else if (acc_on_device (acc_device_radeon)) > + { > + /* The GCC GCN back end is limited to num_workers (16). */ > + workers_actual = 16; > + } That's surprising however, that here we're not similarly having to set 'workers_actual = 1' -- find the explanation in a forthcoming email. ;-) > [...] Grüße Thomas ----------------- Mentor Graphics (Deutschland) GmbH, Arnulfstrasse 201, 80634 München Registergericht München HRB 106955, Geschäftsführer: Thomas Heurung, Frank Thürauf
>From 32099c0d24adb93a031e0301ffd77b065b6f5472 Mon Sep 17 00:00:00 2001 From: Thomas Schwinge <tho...@codesourcery.com> Date: Fri, 4 Jun 2021 15:31:53 +0200 Subject: [PATCH] Fix 'libgomp.oacc-fortran/parallel-dims.f90' for 'acc_device_radeon' ..., by simplifying 'libgomp.oacc-c-c++-common/parallel-dims.c', and updating the former correspondingly. '__builtin_goacc_parlevel_id' does the right thing for all 'acc_device_*'. Follow-up to commit 09e0ad6253f4330977e1b2f116b5e289dc2c2a02 "Update OpenACC tests for amdgcn". libgomp/ * testsuite/libgomp.oacc-c-c++-common/parallel-dims.c: Simplify. * testsuite/libgomp.oacc-fortran/parallel-dims-aux.c: Update. --- .../libgomp.oacc-c-c++-common/parallel-dims.c | 32 ++++--------------- .../libgomp.oacc-fortran/parallel-dims-aux.c | 31 ++++-------------- 2 files changed, 12 insertions(+), 51 deletions(-) diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/parallel-dims.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/parallel-dims.c index c7412c2ef3a..974e1504534 100644 --- a/libgomp/testsuite/libgomp.oacc-c-c++-common/parallel-dims.c +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/parallel-dims.c @@ -10,42 +10,22 @@ #include <openacc.h> #include <gomp-constants.h> -/* TODO: "(int) acc_device_*" casts because of the C++ acc_on_device wrapper - not behaving as expected for -O0. */ #pragma acc routine seq -static unsigned int __attribute__ ((optimize ("O2"))) acc_gang () +static int acc_gang () { - if (acc_on_device ((int) acc_device_host)) - return 0; - else if (acc_on_device ((int) acc_device_nvidia) - || acc_on_device ((int) acc_device_radeon)) - return __builtin_goacc_parlevel_id (GOMP_DIM_GANG); - else - __builtin_abort (); + return __builtin_goacc_parlevel_id (GOMP_DIM_GANG); } #pragma acc routine seq -static unsigned int __attribute__ ((optimize ("O2"))) acc_worker () +static int acc_worker () { - if (acc_on_device ((int) acc_device_host)) - return 0; - else if (acc_on_device ((int) acc_device_nvidia) - || acc_on_device ((int) acc_device_radeon)) - return __builtin_goacc_parlevel_id (GOMP_DIM_WORKER); - else - __builtin_abort (); + return __builtin_goacc_parlevel_id (GOMP_DIM_WORKER); } #pragma acc routine seq -static unsigned int __attribute__ ((optimize ("O2"))) acc_vector () +static int acc_vector () { - if (acc_on_device ((int) acc_device_host)) - return 0; - else if (acc_on_device ((int) acc_device_nvidia) - || acc_on_device ((int) acc_device_radeon)) - return __builtin_goacc_parlevel_id (GOMP_DIM_VECTOR); - else - __builtin_abort (); + return __builtin_goacc_parlevel_id (GOMP_DIM_VECTOR); } diff --git a/libgomp/testsuite/libgomp.oacc-fortran/parallel-dims-aux.c b/libgomp/testsuite/libgomp.oacc-fortran/parallel-dims-aux.c index b5986f4afef..cdece32d147 100644 --- a/libgomp/testsuite/libgomp.oacc-fortran/parallel-dims-aux.c +++ b/libgomp/testsuite/libgomp.oacc-fortran/parallel-dims-aux.c @@ -5,41 +5,22 @@ /* Used by 'parallel-dims.f90'. */ -#include <limits.h> -#include <openacc.h> #include <gomp-constants.h> -/* TODO: "(int) acc_device_*" casts because of the C++ acc_on_device wrapper - not behaving as expected for -O0. */ #pragma acc routine seq -/* static */ unsigned int __attribute__ ((optimize ("O2"))) acc_gang () +/* static */ int acc_gang () { - if (acc_on_device ((int) acc_device_host)) - return 0; - else if (acc_on_device ((int) acc_device_nvidia)) - return __builtin_goacc_parlevel_id (GOMP_DIM_GANG); - else - __builtin_abort (); + return __builtin_goacc_parlevel_id (GOMP_DIM_GANG); } #pragma acc routine seq -/* static */ unsigned int __attribute__ ((optimize ("O2"))) acc_worker () +/* static */ int acc_worker () { - if (acc_on_device ((int) acc_device_host)) - return 0; - else if (acc_on_device ((int) acc_device_nvidia)) - return __builtin_goacc_parlevel_id (GOMP_DIM_WORKER); - else - __builtin_abort (); + return __builtin_goacc_parlevel_id (GOMP_DIM_WORKER); } #pragma acc routine seq -/* static */ unsigned int __attribute__ ((optimize ("O2"))) acc_vector () +/* static */ int acc_vector () { - if (acc_on_device ((int) acc_device_host)) - return 0; - else if (acc_on_device ((int) acc_device_nvidia)) - return __builtin_goacc_parlevel_id (GOMP_DIM_VECTOR); - else - __builtin_abort (); + return __builtin_goacc_parlevel_id (GOMP_DIM_VECTOR); } -- 2.30.2