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

9 files changed, 408 insertions, 4 deletions

diff --git a/libgomp/ChangeLog b/libgomp/ChangeLog
index 0effe1c..e0897e8 100644
--- a/libgomp/ChangeLog
+++ b/libgomp/ChangeLog
@@ -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.
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-red-g-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-red-g-1.c
new file mode 100644
index 0000000..f315cdc
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-red-g-1.c
@@ -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;
+}
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-red-gwv-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-red-gwv-1.c
new file mode 100644
index 0000000..1663e65
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-red-gwv-1.c
@@ -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;
+}
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-red-v-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-red-v-1.c
new file mode 100644
index 0000000..be6916a
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-red-v-1.c
@@ -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;
+}
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-red-v-2.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-red-v-2.c
new file mode 100644
index 0000000..e66732d
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-red-v-2.c
@@ -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;
+}
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-red-w-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-red-w-1.c
new file mode 100644
index 0000000..85fe6eb
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-red-w-1.c
@@ -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;
+}
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 0000000..0059077
--- /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;
+}
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-red-wv-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-red-wv-1.c
new file mode 100644
index 0000000..3a5be5c
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-red-wv-1.c
@@ -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;
+}
diff --git a/libgomp/testsuite/libgomp.oacc-fortran/reduction-5.f90 b/libgomp/testsuite/libgomp.oacc-fortran/reduction-5.f90
index df44a7a..ec13e4e 100644
--- a/libgomp/testsuite/libgomp.oacc-fortran/reduction-5.f90
+++ b/libgomp/testsuite/libgomp.oacc-fortran/reduction-5.f90
@@ -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