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

gcc/testsuite/ChangeLog

@@ -1,3 +1,11 @@
+2015-11-09  Nathan Sidwell  <nathan@codesourcery.com>
+
+	* 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.
+
 2015-11-09  Eric Botcazou  <ebotcazou@adacore.com>
 
 	* gcc.target/sparc/sparc-ret.c: Rename to...

gcc/testsuite/c-c++-common/goacc/routine-1.c

@@ -0,0 +1,34 @@
+
+#pragma acc routine gang
+void gang (void)
+{
+}
+
+#pragma acc routine worker
+void worker (void)
+{
+}
+
+#pragma acc routine vector
+void vector (void)
+{
+}
+
+#pragma acc routine seq
+void seq (void)
+{
+}
+
+int main ()
+{
+
+#pragma acc parallel num_gangs (32) num_workers (32) vector_length (32)
+  {
+    gang ();
+    worker ();
+    vector ();
+    seq ();
+  }
+
+  return 0;
+}

gcc/testsuite/c-c++-common/goacc/routine-2.c

@@ -0,0 +1,21 @@
+#pragma acc routine gang worker /* { dg-error "multiple loop axes" } */
+void gang (void)
+{
+}
+
+#pragma acc routine worker vector /* { dg-error "multiple loop axes" } */
+void worker (void)
+{
+}
+
+#pragma acc routine vector seq /* { dg-error "multiple loop axes" } */
+void vector (void)
+{
+}
+
+#pragma acc routine seq gang /* { dg-error "multiple loop axes" } */
+void seq (void)
+{
+}
+
+#pragma acc routine (nothing) gang /* { dg-error "not been declared" } */

gcc/testsuite/c-c++-common/goacc/routine-3.c

@@ -0,0 +1,53 @@
+#pragma acc routine gang
+void gang (void) /* { dg-message "declared here" 3 } */
+{
+}
+
+#pragma acc routine worker
+void worker (void) /* { dg-message "declared here" 2 } */
+{
+}
+
+#pragma acc routine vector
+void vector (void) /* { dg-message "declared here" 1 } */
+{
+}
+
+#pragma acc routine seq
+void seq (void)
+{
+}
+
+int main ()
+{
+
+#pragma acc parallel num_gangs (32) num_workers (32) vector_length (32)
+  {
+    #pragma acc loop gang /* { dg-message "loop here" 1 } */
+    for (int i = 0; i < 10; i++)
+      {
+	gang (); /*  { dg-error "routine call uses same" } */
+	worker ();
+	vector ();
+	seq ();
+      }
+    #pragma acc loop worker /* { dg-message "loop here" 2 } */
+    for (int i = 0; i < 10; i++)
+      {
+	gang (); /*  { dg-error "routine call uses same" } */
+	worker (); /*  { dg-error "routine call uses same" } */
+	vector ();
+	seq ();
+      }
+    #pragma acc loop vector /* { dg-message "loop here" 3 } */
+    for (int i = 0; i < 10; i++)
+      {
+	gang (); /*  { dg-error "routine call uses same" } */
+	worker (); /*  { dg-error "routine call uses same" } */
+	vector (); /*  { dg-error "routine call uses same" } */
+	seq ();
+      }
+  }
+
+  return 0;
+}

gcc/testsuite/c-c++-common/goacc/routine-4.c

@@ -0,0 +1,41 @@
+
+void gang (void);
+void worker (void);
+void vector (void);
+
+#pragma acc routine (gang) gang
+#pragma acc routine (worker) worker
+#pragma acc routine (vector) vector
+  
+#pragma acc routine seq
+void seq (void)
+{
+  gang ();  /* { dg-error "routine call uses" } */
+  worker ();  /* { dg-error "routine call uses" } */
+  vector ();  /* { dg-error "routine call uses" } */
+  seq ();
+}
+
+void vector (void) /* { dg-message "declared here" 1 } */
+{
+  gang ();  /* { dg-error "routine call uses" } */
+  worker ();  /* { dg-error "routine call uses" } */
+  vector ();
+  seq ();
+}
+
+void worker (void) /* { dg-message "declared here" 2 } */
+{
+  gang ();  /* { dg-error "routine call uses" } */
+  worker ();
+  vector ();
+  seq ();
+}
+
+void gang (void) /* { dg-message "declared here" 3 } */
+{
+  gang ();
+  worker ();
+  vector ();
+  seq ();
+}

gcc/testsuite/c-c++-common/goacc/routine-5.c

@@ -0,0 +1,51 @@
+/* { dg-do compile } */
+
+#pragma acc routine /* { dg-error "not followed by" } */
+int a;
+
+#if 0 /* Disable for the moment.  */
+#pragma acc routine /* dg-error "not followed by" */
+void fn1 (void), fn1b (void);
+#endif
+
+#pragma acc routine /* { dg-error "not followed by" } */
+int b, fn2 (void);
+
+#if 0 /* Disable for the moment.  */
+#pragma acc routine /*  dg-error "not followed by"  */
+int fn3 (void), b2;
+#endif
+
+#pragma acc routine /* { dg-error "not followed by" } */
+typedef struct c c;
+
+#pragma acc routine /* { dg-error "not followed by" } */
+struct d {} d;
+
+#pragma acc routine /* { dg-error "not followed by" } */
+#pragma acc routine
+int fn4 (void);
+
+int fn5a (void);
+
+#pragma acc routine /* { dg-error "not followed by" } */
+#pragma acc routine (fn5a)
+int fn5 (void);
+
+#ifdef __cplusplus
+
+#pragma acc routine /* { dg-error "not followed by" "" { target c++ } } */
+namespace f {}
+
+namespace g {}
+
+#pragma acc routine /* { dg-error "not followed by" "" { target c++ } } */
+using namespace g;
+
+#pragma acc routine (g) /* { dg-error "does not refer to" "" { target c++ } } */
+
+#endif
+
+#pragma acc routine (a) /* { dg-error "does not refer to" } */
+  
+#pragma acc routine (c) /* { dg-error "does not refer to" } */

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.

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

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

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

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

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

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