routine-1.c: New.

	gcc/testsuite/
	* c-c++-common/goacc/routine-1.c: New.
	* c-c++-common/goacc/routine-2.c: New.
	* c-c++-common/goacc/routine-3.c: New.
	* c-c++-common/goacc/routine-4.c: New.
	* c-c++-common/goacc/routine-5.c: New.

	libgomp/
	* testsuite/libgomp.oacc-c-c++-common/firstprivate-1.c: New.
	* testsuite/libgomp.oacc-c-c++-common/routine-g-1.c: New.
	* testsuite/libgomp.oacc-c-c++-common/routine-gwv-1.c: New.
	* testsuite/libgomp.oacc-c-c++-common/routine-v-1.c: New.
	* testsuite/libgomp.oacc-c-c++-common/routine-w-1.c: New.
	* testsuite/libgomp.oacc-c-c++-common/routine-wv-1.c: New.

From-SVN: r230080

7 files changed, 372 insertions, 0 deletions

diff --git a/libgomp/ChangeLog b/libgomp/ChangeLog
index ab2a25a..08d467b 100644
--- a/libgomp/ChangeLog
+++ b/libgomp/ChangeLog
@@ -1,3 +1,12 @@
+2015-11-09  Nathan Sidwell  <nathan@codesourcery.com>
+
+	* testsuite/libgomp.oacc-c-c++-common/firstprivate-1.c: New.
+	* testsuite/libgomp.oacc-c-c++-common/routine-g-1.c: New.
+	* testsuite/libgomp.oacc-c-c++-common/routine-gwv-1.c: New.
+	* testsuite/libgomp.oacc-c-c++-common/routine-v-1.c: New.
+	* testsuite/libgomp.oacc-c-c++-common/routine-w-1.c: New.
+	* testsuite/libgomp.oacc-c-c++-common/routine-wv-1.c: New.
+
 2015-11-06  Thomas Schwinge  <thomas@codesourcery.com>
 
 	* testsuite/libgomp.oacc-c-c++-common/loop-red-v-2.c: XFAIL.
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/firstprivate-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/firstprivate-1.c
new file mode 100644
index 0000000..7f5d3d3
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/firstprivate-1.c
@@ -0,0 +1,41 @@
+/* { dg-do run } */
+
+#include  <openacc.h>
+
+int main ()
+{
+  int ok = 1;
+  int val = 2;
+  int ary[32];
+  int ondev = 0;
+
+  for (int i = 0; i < 32; i++)
+    ary[i] = ~0;
+  
+#pragma acc parallel num_gangs (32) copy (ok) firstprivate (val) copy(ary, ondev)
+  {
+    ondev = acc_on_device (acc_device_not_host);
+#pragma acc loop gang(static:1)
+    for (unsigned i = 0; i < 32; i++)
+      {
+	if (val != 2)
+	  ok = 0;
+	val += i;
+	ary[i] = val;
+      }
+  }
+
+  if (ondev)
+    {
+      if (!ok)
+	return 1;
+      if (val != 2)
+	return 1;
+
+      for (int i = 0; i < 32; i++)
+	if (ary[i] != 2 + i)
+	  return 1;
+    }
+  
+  return 0;
+}
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/routine-g-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/routine-g-1.c
new file mode 100644
index 0000000..b1f238a
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/routine-g-1.c
@@ -0,0 +1,64 @@
+/* { dg-do run } */
+/* { dg-additional-options "-O2" */
+
+#include <stdio.h>
+
+#define N (32*32*32+17)
+
+#pragma acc routine gang
+void __attribute__ ((noinline)) gang (int ary[N])
+{
+#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;
+	  }
+	else
+	  ary[ix] = ix;
+      }
+}
+
+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)
+  {
+    ondev = __builtin_acc_on_device (5);
+    gang (ary);
+  }
+
+  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/routine-gwv-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/routine-gwv-1.c
new file mode 100644
index 0000000..5980747
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/routine-gwv-1.c
@@ -0,0 +1,66 @@
+/* { dg-do run } */
+/* { dg-additional-options "-O2" */
+
+#include <stdio.h>
+
+#define N (32*32*32+17)
+
+#pragma acc routine gang
+void __attribute__ ((noinline)) gang (int ary[N])
+{
+#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;
+	}
+      else
+	ary[ix] = ix;
+    }
+}
+
+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)
+  {
+    ondev = __builtin_acc_on_device (5);
+    gang (ary);
+  }
+
+  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/routine-v-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/routine-v-1.c
new file mode 100644
index 0000000..29e0f90
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/routine-v-1.c
@@ -0,0 +1,64 @@
+/* { dg-do run } */
+/* { dg-additional-options "-O2" */
+
+#include <stdio.h>
+
+#define N (32*32*32+17)
+
+#pragma acc routine vector
+void __attribute__ ((noinline)) vector (int ary[N])
+{
+#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;
+	}
+      else
+	ary[ix] = ix;
+    }
+}
+
+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)
+  {
+    ondev = __builtin_acc_on_device (5);
+    vector (ary);
+  }
+
+  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/routine-w-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/routine-w-1.c
new file mode 100644
index 0000000..95e7174
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/routine-w-1.c
@@ -0,0 +1,64 @@
+/* { dg-do run } */
+/* { dg-additional-options "-O2" */
+
+#include <stdio.h>
+
+#define N (32*32*32+17)
+
+#pragma acc routine worker
+void __attribute__ ((noinline)) worker (int ary[N])
+{
+#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;
+	}
+      else
+	ary[ix] = ix;
+    }
+}
+
+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)
+  {
+    ondev = __builtin_acc_on_device (5);
+    worker (ary);
+  }
+
+  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/routine-wv-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/routine-wv-1.c
new file mode 100644
index 0000000..3e8dcde
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/routine-wv-1.c
@@ -0,0 +1,64 @@
+/* { dg-do run } */
+/* { dg-additional-options "-O2" */
+
+#include <stdio.h>
+
+#define N (32*32*32+17)
+
+#pragma acc routine worker
+void __attribute__ ((noinline)) worker (int ary[N])
+{
+#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;
+	}
+      else
+	ary[ix] = ix;
+    }
+}
+
+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)
+  {
+    ondev = __builtin_acc_on_device (5);
+    worker (ary);
+  }
+
+  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;
+}