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

loop-red-g-1.c: New. 

* libgomp.oacc-c-c++-common/loop-red-g-1.c: New.
	* libgomp.oacc-c-c++-common/loop-red-gwv-1.c: New.
	* libgomp.oacc-c-c++-common/loop-red-v-1.c: New.
	* libgomp.oacc-c-c++-common/loop-red-v-2.c: New.
	* libgomp.oacc-c-c++-common/loop-red-w-1.c: New.
	* libgomp.oacc-c-c++-common/loop-red-w-2.c: New.
	* libgomp.oacc-c-c++-common/loop-red-wv-1.c: New.
	* libgomp.oacc-fortran/reduction-5.f90: Avoid reference var.

From-SVN: r229770
This commit is contained in:
Nathan Sidwell 2015-11-04 17:01:23 +00:00 committed by Nathan Sidwell
parent 4d49f9642a
commit 5d6823a25e
9 changed files with 408 additions and 4 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

11
libgomp/ChangeLog
View File

@ -1,3 +1,14 @@
2015-11-04 Nathan Sidwell <nathan@codesourcery.com>
* libgomp.oacc-c-c++-common/loop-red-g-1.c: New.
* libgomp.oacc-c-c++-common/loop-red-gwv-1.c: New.
* libgomp.oacc-c-c++-common/loop-red-v-1.c: New.
* libgomp.oacc-c-c++-common/loop-red-v-2.c: New.
* libgomp.oacc-c-c++-common/loop-red-w-1.c: New.
* libgomp.oacc-c-c++-common/loop-red-w-2.c: New.
* libgomp.oacc-c-c++-common/loop-red-wv-1.c: New.
* libgomp.oacc-fortran/reduction-5.f90: Avoid reference var.
2015-11-03 Nathan Sidwell <nathan@codesourcery.com>
* libgomp.h (struct acc_dispatch_t): Remove args from exec_func.

54
libgomp/testsuite/libgomp.oacc-c-c++-common/loop-red-g-1.c Normal file
View File

@ -0,0 +1,54 @@
/* { dg-do run } */
/* { dg-additional-options "-O2" */
#include <stdio.h>
#define N (32*32*32+17)
int main ()
{
int ix;
int ondev = 0;
int t = 0, h = 0;
#pragma acc parallel num_gangs(32) vector_length(32) copy(t) copy(ondev)
{
#pragma acc loop gang 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;
}
}
for (ix = 0; ix < N; ix++)
{
int val = ix;
if(ondev)
{
int g = ix / ((N + 31) / 32);
int w = 0;
int v = 0;
val = (g << 16) | (w << 8) | v;
}
h += val;
}
if (t != h)
{
printf ("t=%x expected %x\n", t, h);
return 1;
}
return 0;
}

56
libgomp/testsuite/libgomp.oacc-c-c++-common/loop-red-gwv-1.c Normal file
View File

@ -0,0 +1,56 @@
/* { dg-do run } */
/* { dg-additional-options "-O2" */
#include <stdio.h>
#define N (32*32*32+17)
int main ()
{
int ix;
int ondev = 0;
int t = 0, h = 0;
#pragma acc parallel num_gangs(32) num_workers(32) vector_length(32) copy(t) copy(ondev)
{
#pragma acc loop gang worker vector 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;
}
}
for (ix = 0; ix < N; ix++)
{
int val = 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;
val = (g << 16) | (w << 8) | v;
}
h += val;
}
if (t != h)
{
printf ("t=%x expected %x\n", t, h);
return 1;
}
return 0;
}

56
libgomp/testsuite/libgomp.oacc-c-c++-common/loop-red-v-1.c Normal file
View File

@ -0,0 +1,56 @@
/* { dg-do run } */
/* { dg-additional-options "-O2" */
#include <stdio.h>
#define N (32*32*32+17)
int main ()
{
int ix;
int ondev = 0;
int t = 0, h = 0;
#pragma acc parallel vector_length(32) copy(t) copy(ondev)
{
#pragma acc loop vector 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;
}
}
for (ix = 0; ix < N; ix++)
{
int val = ix;
if (ondev)
{
int g = 0;
int w = 0;
int v = ix % 32;
val = (g << 16) | (w << 8) | v;
}
h += val;
}
if (t != h)
{
printf ("t=%x expected %x\n", t, h);
return 1;
}
return 0;
}

59
libgomp/testsuite/libgomp.oacc-c-c++-common/loop-red-v-2.c Normal file
View File

@ -0,0 +1,59 @@
/* { 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 vector_length(32) copy(q) copy(ondev)
{
int t = q;
#pragma acc loop vector 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 = 0;
int v = ix % 32;
val = (g << 16) | (w << 8) | v;
}
h += val;
}
if (q != h)
{
printf ("t=%x expected %x\n", q, h);
return 1;
}
return 0;
}

54
libgomp/testsuite/libgomp.oacc-c-c++-common/loop-red-w-1.c Normal file
View File

@ -0,0 +1,54 @@
/* { dg-do run } */
/* { dg-additional-options "-O2" */
#include <stdio.h>
#define N (32*32*32+17)
int main ()
{
int ix;
int ondev = 0;
int t = 0, h = 0;
#pragma acc parallel num_workers(32) vector_length(32) copy(t) copy(ondev)
{
#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;
}
}
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 (t != h)
{
printf ("t=%x expected %x\n", t, h);
return 1;
}
return 0;
}

57
libgomp/testsuite/libgomp.oacc-c-c++-common/loop-red-w-2.c Normal file
View File

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

54
libgomp/testsuite/libgomp.oacc-c-c++-common/loop-red-wv-1.c Normal file
View File

@ -0,0 +1,54 @@
/* { dg-do run } */
/* { dg-additional-options "-O2" */
#include <stdio.h>
#define N (32*32*32+17)
int main ()
{
int ix;
int ondev = 0;
int t = 0, h = 0;
#pragma acc parallel num_workers(32) vector_length(32) copy(t) copy(ondev)
{
#pragma acc loop worker vector 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;
}
}
for (ix = 0; ix < N; ix++)
{
int val = ix;
if(ondev)
{
int g = 0;
int w = (ix / 32) % 32;
int v = ix % 32;
val = (g << 16) | (w << 8) | v;
}
h += val;
}
if (t != h)
{
printf ("t=%x expected %x\n", t, h);
return 1;
}
return 0;
}

11
libgomp/testsuite/libgomp.oacc-fortran/reduction-5.f90
View File

@ -21,12 +21,15 @@ end program reduction
subroutine redsub(sum, n, c)
integer :: sum, n, c
sum = 0
integer :: s
s = 0
!$acc parallel vector_length(n) copyin (n, c) num_gangs(1)
!$acc loop reduction(+:sum)
!$acc parallel vector_length(32) copyin (n, c) copy (s) num_gangs(1)
!$acc loop reduction(+:s)
do i = 1, n
sum = sum + c
s = s + c
end do
!$acc end parallel
sum = s
end subroutine redsub