09e0ad6253
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.
61 lines
1.0 KiB
C
61 lines
1.0 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 q = 0, h = 0;
|
|
int vectorsize;
|
|
|
|
#pragma acc parallel vector_length(32) copy(q) copy(ondev) copyout(vectorsize)
|
|
{
|
|
int t = q;
|
|
|
|
#pragma acc loop 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;
|
|
}
|
|
q = t;
|
|
vectorsize = __builtin_goacc_parlevel_size (GOMP_DIM_VECTOR);
|
|
}
|
|
|
|
for (ix = 0; ix < N; ix++)
|
|
{
|
|
int val = ix;
|
|
if (ondev)
|
|
{
|
|
int g = 0;
|
|
int w = 0;
|
|
int v = ix % vectorsize;
|
|
|
|
val = (g << 16) | (w << 8) | v;
|
|
}
|
|
h += val;
|
|
}
|
|
|
|
if (q != h)
|
|
{
|
|
printf ("t=%x expected %x\n", q, h);
|
|
return 1;
|
|
}
|
|
|
|
return 0;
|
|
}
|