diff options
Diffstat (limited to 'libgomp/testsuite/libgomp.oacc-c-c++-common/loop-red-w-2.c')
-rw-r--r-- | libgomp/testsuite/libgomp.oacc-c-c++-common/loop-red-w-2.c | 57 |
1 files changed, 57 insertions, 0 deletions
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-red-w-2.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-red-w-2.c new file mode 100644 index 00000000000..0059077b685 --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-red-w-2.c @@ -0,0 +1,57 @@ +/* { dg-do run } */ +/* { dg-additional-options "-O2" */ + +#include <stdio.h> + +#define N (32*32*32+17) +int main () +{ + int ix; + int ondev = 0; + int q = 0, h = 0; + +#pragma acc parallel num_workers(32) vector_length(32) copy(q) copy(ondev) + { + int t = q; + +#pragma acc loop worker reduction(+:t) + for (unsigned ix = 0; ix < N; ix++) + { + int val = ix; + + if (__builtin_acc_on_device (5)) + { + int g = 0, w = 0, v = 0; + + __asm__ volatile ("mov.u32 %0,%%ctaid.x;" : "=r" (g)); + __asm__ volatile ("mov.u32 %0,%%tid.y;" : "=r" (w)); + __asm__ volatile ("mov.u32 %0,%%tid.x;" : "=r" (v)); + val = (g << 16) | (w << 8) | v; + ondev = 1; + } + t += val; + } + q = t; + } + + for (ix = 0; ix < N; ix++) + { + int val = ix; + if(ondev) + { + int g = 0; + int w = ix % 32; + int v = 0; + + val = (g << 16) | (w << 8) | v; + } + h += val; + } + if (q != h) + { + printf ("t=%x expected %x\n", q, h); + return 1; + } + + return 0; +} |