![Andrew Stubbs](/assets/img/avatar_default.png)
2020-01-20 Andrew Stubbs <ams@codesourcery.com> libgomp/ * testsuite/libgomp.oacc-c-c++-common/loop-auto-1.c: Skip test on gcn. * testsuite/libgomp.oacc-c-c++-common/loop-dim-default.c (main): Adjust test dimensions for amdgcn. * testsuite/libgomp.oacc-c-c++-common/loop-gwv-1.c (main): Adjust gang/worker/vector expectations dynamically. * testsuite/libgomp.oacc-c-c++-common/loop-red-gwv-1.c (main): Likewise. * testsuite/libgomp.oacc-c-c++-common/loop-red-v-1.c (main): Likewise. * testsuite/libgomp.oacc-c-c++-common/loop-red-v-2.c (main): Likewise. * testsuite/libgomp.oacc-c-c++-common/loop-red-w-1.c (main): Likewise. * testsuite/libgomp.oacc-c-c++-common/loop-red-w-2.c (main): Likewise. * testsuite/libgomp.oacc-c-c++-common/loop-red-wv-1.c (main): Likewise. * testsuite/libgomp.oacc-c-c++-common/loop-v-1.c (main): Likewise. * testsuite/libgomp.oacc-c-c++-common/loop-w-1.c (main): Likewise. * testsuite/libgomp.oacc-c-c++-common/loop-wv-1.c (main): Likewise. * testsuite/libgomp.oacc-c-c++-common/parallel-dims.c (acc_gang): Recognise acc_device_radeon. (acc_worker): Likewise. (acc_vector): Likewise. (main): Set expectations for amdgcn. * testsuite/libgomp.oacc-c-c++-common/routine-gwv-1.c (main): Adjust gang/worker/vector expectations dynamically. * testsuite/libgomp.oacc-c-c++-common/routine-v-1.c (main): Likewise. * testsuite/libgomp.oacc-c-c++-common/routine-w-1.c (main): Likewise. * testsuite/libgomp.oacc-c-c++-common/routine-wv-1.c (main): Likewise. * testsuite/libgomp.oacc-c-c++-common/routine-wv-2.c: Set expectations for amdgcn.
62 lines
1.4 KiB
C
62 lines
1.4 KiB
C
#include <stdio.h>
|
|
#include <openacc.h>
|
|
#include <gomp-constants.h>
|
|
|
|
#define N (32*32*32+17)
|
|
int main ()
|
|
{
|
|
int ix;
|
|
int ondev = 0;
|
|
int t = 0, h = 0;
|
|
int gangsize, workersize, vectorsize;
|
|
|
|
#pragma acc parallel num_gangs(32) num_workers(32) vector_length(32) \
|
|
copy(ondev) copyout(gangsize, workersize, vectorsize)
|
|
{
|
|
#pragma acc loop gang worker vector reduction(+:t)
|
|
for (unsigned ix = 0; ix < N; ix++)
|
|
{
|
|
int val = ix;
|
|
|
|
if (acc_on_device (acc_device_not_host))
|
|
{
|
|
int g, w, v;
|
|
|
|
g = __builtin_goacc_parlevel_id (GOMP_DIM_GANG);
|
|
w = __builtin_goacc_parlevel_id (GOMP_DIM_WORKER);
|
|
v = __builtin_goacc_parlevel_id (GOMP_DIM_VECTOR);
|
|
val = (g << 16) | (w << 8) | v;
|
|
ondev = 1;
|
|
}
|
|
t += val;
|
|
}
|
|
gangsize = __builtin_goacc_parlevel_size (GOMP_DIM_GANG);
|
|
workersize = __builtin_goacc_parlevel_size (GOMP_DIM_WORKER);
|
|
vectorsize = __builtin_goacc_parlevel_size (GOMP_DIM_VECTOR);
|
|
}
|
|
|
|
for (ix = 0; ix < N; ix++)
|
|
{
|
|
int val = ix;
|
|
if (ondev)
|
|
{
|
|
int chunk_size = (N + gangsize * workersize * vectorsize - 1)
|
|
/ (gangsize * workersize * vectorsize);
|
|
|
|
int g = ix / (chunk_size * vectorsize * workersize);
|
|
int w = ix / vectorsize % workersize;
|
|
int v = ix % vectorsize;
|
|
|
|
val = (g << 16) | (w << 8) | v;
|
|
}
|
|
h += val;
|
|
}
|
|
if (t != h)
|
|
{
|
|
printf ("t=%x expected %x\n", t, h);
|
|
return 1;
|
|
}
|
|
|
|
return 0;
|
|
}
|