From mboxrd@z Thu Jan 1 00:00:00 1970 Return-Path: Received: by sourceware.org (Postfix, from userid 1643) id 99B5B3847830; Tue, 8 Jun 2021 09:45:44 +0000 (GMT) DKIM-Filter: OpenDKIM Filter v2.11.0 sourceware.org 99B5B3847830 MIME-Version: 1.0 Content-Transfer-Encoding: 7bit Content-Type: text/plain; charset="utf-8" From: Thomas Schwinge To: gcc-cvs@gcc.gnu.org Subject: [gcc r12-1290] Fix 'libgomp.oacc-fortran/parallel-dims.f90' for 'acc_device_radeon' X-Act-Checkin: gcc X-Git-Author: Thomas Schwinge X-Git-Refname: refs/heads/master X-Git-Oldrev: 984df1e1630f262d782c00cefad2643b8e8469f8 X-Git-Newrev: 32099c0d24adb93a031e0301ffd77b065b6f5472 Message-Id: <20210608094544.99B5B3847830@sourceware.org> Date: Tue, 8 Jun 2021 09:45:44 +0000 (GMT) X-BeenThere: gcc-cvs@gcc.gnu.org X-Mailman-Version: 2.1.29 Precedence: list List-Id: Gcc-cvs mailing list List-Unsubscribe: , List-Archive: List-Help: List-Subscribe: , X-List-Received-Date: Tue, 08 Jun 2021 09:45:44 -0000 https://gcc.gnu.org/g:32099c0d24adb93a031e0301ffd77b065b6f5472 commit r12-1290-g32099c0d24adb93a031e0301ffd77b065b6f5472 Author: Thomas Schwinge Date: Fri Jun 4 15:31:53 2021 +0200 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. Diff: --- .../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 #include -/* 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 -#include #include -/* 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); }