OpenE2K/gcc
2
0
Fork 0
Code Issues Pull Requests Projects Releases Wiki Activity

Fix 'libgomp.oacc-fortran/parallel-dims.f90' for 'acc_device_radeon'

Browse Source 
..., 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.
This commit is contained in:
Thomas Schwinge 2021-06-04 15:31:53 +02:00
parent 984df1e163
commit 32099c0d24
2 changed files with 12 additions and 51 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

32
libgomp/testsuite/libgomp.oacc-c-c++-common/parallel-dims.c
View File

@ -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);
}

31
libgomp/testsuite/libgomp.oacc-fortran/parallel-dims-aux.c
View File

@ -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);
}