From: Nathan Sidwell Date: Wed, 28 Oct 2015 14:27:30 +0000 (+0000) Subject: loop-g-1.c: New. X-Git-Url: https://git.libre-soc.org/?a=commitdiff_plain;h=e3091406955ec3f0642271143a5007f46ec8c876;p=gcc.git loop-g-1.c: New. * testsuite/libgomp.oacc-c-c++-common/loop-g-1.c: New. * testsuite/libgomp.oacc-c-c++-common/loop-g-2.c: New. * testsuite/libgomp.oacc-c-c++-common/loop-gwv-1.c: New. * testsuite/libgomp.oacc-c-c++-common/loop-v-1.c: New. * testsuite/libgomp.oacc-c-c++-common/loop-w-1.c: New. * testsuite/libgomp.oacc-c-c++-common/loop-wv-1.c: New. From-SVN: r229487 --- diff --git a/libgomp/ChangeLog b/libgomp/ChangeLog index 0f6fd91de1f..8f44af02606 100644 --- a/libgomp/ChangeLog +++ b/libgomp/ChangeLog @@ -1,3 +1,12 @@ +2015-10-28 Nathan Sidwell + + * testsuite/libgomp.oacc-c-c++-common/loop-g-1.c: New. + * testsuite/libgomp.oacc-c-c++-common/loop-g-2.c: New. + * testsuite/libgomp.oacc-c-c++-common/loop-gwv-1.c: New. + * testsuite/libgomp.oacc-c-c++-common/loop-v-1.c: New. + * testsuite/libgomp.oacc-c-c++-common/loop-w-1.c: New. + * testsuite/libgomp.oacc-c-c++-common/loop-wv-1.c: New. + 2015-10-27 Nathan Sidwell * plugin/plugin-nvptx.c (nvptx_exec): Remove check on compute diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-g-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-g-1.c new file mode 100644 index 00000000000..58545d04b26 --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-g-1.c @@ -0,0 +1,57 @@ +/* { dg-do run } */ +/* { dg-additional-options "-O2" */ + +#include + +#define N (32*32*32+17) +int main () +{ + int ary[N]; + int ix; + int exit = 0; + int ondev = 0; + + for (ix = 0; ix < N;ix++) + ary[ix] = -1; + +#pragma acc parallel num_gangs(32) vector_length(32) copy(ary) copy(ondev) + { +#pragma acc loop gang + for (unsigned ix = 0; ix < N; 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)); + ary[ix] = (g << 16) | (w << 8) | v; + ondev = 1; + } + else + ary[ix] = ix; + } + } + + for (ix = 0; ix < N; ix++) + { + int expected = ix; + if(ondev) + { + int g = ix / ((N + 31) / 32); + int w = 0; + int v = 0; + + expected = (g << 16) | (w << 8) | v; + } + + if (ary[ix] != expected) + { + exit = 1; + printf ("ary[%d]=%x expected %x\n", ix, ary[ix], expected); + } + } + + return exit; +} diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-g-2.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-g-2.c new file mode 100644 index 00000000000..c01c6fa66a6 --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-g-2.c @@ -0,0 +1,57 @@ +/* { dg-do run } */ +/* { dg-additional-options "-O2" */ + +#include + +#define N (32*32*32+17) +int main () +{ + int ary[N]; + int ix; + int exit = 0; + int ondev = 0; + + for (ix = 0; ix < N;ix++) + ary[ix] = -1; + +#pragma acc parallel num_gangs(32) vector_length(32) copy(ary) copy(ondev) + { +#pragma acc loop gang (static:1) + for (unsigned ix = 0; ix < N; 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)); + ary[ix] = (g << 16) | (w << 8) | v; + ondev = 1; + } + else + ary[ix] = ix; + } + } + + for (ix = 0; ix < N; ix++) + { + int expected = ix; + if(ondev) + { + int g = ix % 32; + int w = 0; + int v = 0; + + expected = (g << 16) | (w << 8) | v; + } + + if (ary[ix] != expected) + { + exit = 1; + printf ("ary[%d]=%x expected %x\n", ix, ary[ix], expected); + } + } + + return exit; +} diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-gwv-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-gwv-1.c new file mode 100644 index 00000000000..f23e2f379d2 --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-gwv-1.c @@ -0,0 +1,59 @@ +/* { dg-do run } */ +/* { dg-additional-options "-O2" */ + +#include + +#define N (32*32*32+17) +int main () +{ + int ary[N]; + int ix; + int exit = 0; + int ondev = 0; + + for (ix = 0; ix < N;ix++) + ary[ix] = -1; + +#pragma acc parallel num_gangs(32) num_workers(32) vector_length(32) copy(ary) copy(ondev) + { +#pragma acc loop gang worker vector + for (unsigned ix = 0; ix < N; 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)); + ary[ix] = (g << 16) | (w << 8) | v; + ondev = 1; + } + else + ary[ix] = ix; + } + } + + for (ix = 0; ix < N; ix++) + { + int expected = ix; + if(ondev) + { + int chunk_size = (N + 32*32*32 - 1) / (32*32*32); + + int g = ix / (chunk_size * 32 * 32); + int w = ix / 32 % 32; + int v = ix % 32; + + expected = (g << 16) | (w << 8) | v; + } + + if (ary[ix] != expected) + { + exit = 1; + printf ("ary[%d]=%x expected %x\n", ix, ary[ix], expected); + } + } + + return exit; +} diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-v-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-v-1.c new file mode 100644 index 00000000000..70c62922ec0 --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-v-1.c @@ -0,0 +1,57 @@ +/* { dg-do run } */ +/* { dg-additional-options "-O2" */ + +#include + +#define N (32*32*32+17) +int main () +{ + int ary[N]; + int ix; + int exit = 0; + int ondev = 0; + + for (ix = 0; ix < N;ix++) + ary[ix] = -1; + +#pragma acc parallel vector_length(32) copy(ary) copy(ondev) + { +#pragma acc loop vector + for (unsigned ix = 0; ix < N; 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)); + ary[ix] = (g << 16) | (w << 8) | v; + ondev = 1; + } + else + ary[ix] = ix; + } + } + + for (ix = 0; ix < N; ix++) + { + int expected = ix; + if(ondev) + { + int g = 0; + int w = 0; + int v = ix % 32; + + expected = (g << 16) | (w << 8) | v; + } + + if (ary[ix] != expected) + { + exit = 1; + printf ("ary[%d]=%x expected %x\n", ix, ary[ix], expected); + } + } + + return exit; +} diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-w-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-w-1.c new file mode 100644 index 00000000000..5473c2d609d --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-w-1.c @@ -0,0 +1,57 @@ +/* { dg-do run } */ +/* { dg-additional-options "-O2" */ + +#include + +#define N (32*32*32+17) +int main () +{ + int ary[N]; + int ix; + int exit = 0; + int ondev = 0; + + for (ix = 0; ix < N;ix++) + ary[ix] = -1; + +#pragma acc parallel num_workers(32) vector_length(32) copy(ary) copy(ondev) + { +#pragma acc loop worker + for (unsigned ix = 0; ix < N; 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)); + ary[ix] = (g << 16) | (w << 8) | v; + ondev = 1; + } + else + ary[ix] = ix; + } + } + + for (ix = 0; ix < N; ix++) + { + int expected = ix; + if(ondev) + { + int g = 0; + int w = ix % 32; + int v = 0; + + expected = (g << 16) | (w << 8) | v; + } + + if (ary[ix] != expected) + { + exit = 1; + printf ("ary[%d]=%x expected %x\n", ix, ary[ix], expected); + } + } + + return exit; +} diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-wv-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-wv-1.c new file mode 100644 index 00000000000..85e447652ad --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-wv-1.c @@ -0,0 +1,57 @@ +/* { dg-do run } */ +/* { dg-additional-options "-O2" */ + +#include + +#define N (32*32*32+17) +int main () +{ + int ary[N]; + int ix; + int exit = 0; + int ondev = 0; + + for (ix = 0; ix < N;ix++) + ary[ix] = -1; + +#pragma acc parallel num_workers(32) vector_length(32) copy(ary) copy(ondev) + { +#pragma acc loop worker vector + for (unsigned ix = 0; ix < N; 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)); + ary[ix] = (g << 16) | (w << 8) | v; + ondev = 1; + } + else + ary[ix] = ix; + } + } + + for (ix = 0; ix < N; ix++) + { + int expected = ix; + if(ondev) + { + int g = 0; + int w = (ix / 32) % 32; + int v = ix % 32; + + expected = (g << 16) | (w << 8) | v; + } + + if (ary[ix] != expected) + { + exit = 1; + printf ("ary[%d]=%x expected %x\n", ix, ary[ix], expected); + } + } + + return exit; +}