omp-low.c (scan_sharing_clauses): Accept INDEPENDENT, AUTO & SEQ.

	gcc/
	* gcc/omp-low.c (scan_sharing_clauses): Accept INDEPENDENT, AUTO &
	SEQ.
	(oacc_loop_fixed_partitions): Correct return type to bool.
	(oacc_loop_auto_partitions): New.
	(oacc_loop_partition): Take mask argument, call
	oacc_loop_auto_partitions.
	(execute_oacc_device_lower): Provide mask to oacc_loop_partition.

	gcc/testsuite/
	* c-c++-common/goacc/loop-auto-1.c: New.

	libgomp/
	* testsuite/libgomp.oacc-c-c++-common/loop-auto-1.c: New.

From-SVN: r230354

2 files changed, 227 insertions, 0 deletions

diff --git a/libgomp/ChangeLog b/libgomp/ChangeLog
index eb1db83..2f5a4d1 100644
--- a/libgomp/ChangeLog
+++ b/libgomp/ChangeLog
@@ -1,5 +1,7 @@
 2015-11-13  Nathan Sidwell  <nathan@codesourcery.com>
 
+	* testsuite/libgomp.oacc-c-c++-common/loop-auto-1.c: New.
+
 	* testsuite/libgomp.oacc-c-c++-common/collapse-2.c: Sequential
 	loop is sequential.
 
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-auto-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-auto-1.c
new file mode 100644
index 0000000..174af91
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-auto-1.c
@@ -0,0 +1,225 @@
+/* { dg-do run } */
+/* { dg-additional-options "-O2" */
+
+#include <stdio.h>
+#include <openacc.h>
+
+int check (const int *ary, int size, int gp, int wp, int vp)
+{
+  int exit = 0;
+  int ix;
+  int gangs[32], workers[32], vectors[32];
+
+  for (ix = 0; ix < 32; ix++)
+    gangs[ix] = workers[ix] = vectors[ix] = 0;
+  
+  for (ix = 0; ix < size; ix++)
+    {
+      vectors[ary[ix] & 0xff]++;
+      workers[(ary[ix] >> 8) & 0xff]++;
+      gangs[(ary[ix] >> 16) & 0xff]++;
+    }
+
+  for (ix = 0; ix < 32; ix++)
+    {
+      if (gp)
+	{
+	  int expect = gangs[0];
+	  if (gangs[ix] != expect)
+	    {
+	      exit = 1;
+	      printf ("gang %d not used %d times\n", ix, expect);
+	    }
+	}
+      else if (ix && gangs[ix])
+	{
+	  exit = 1;
+	  printf ("gang %d unexpectedly used\n", ix);
+	}
+
+      if (wp)
+	{
+	  int expect = workers[0];
+	  if (workers[ix] != expect)
+	    {
+	      exit = 1;
+	      printf ("worker %d not used %d times\n", ix, expect);
+	    }
+	}
+      else if (ix && workers[ix])
+	{
+	  exit = 1;
+	  printf ("worker %d unexpectedly used\n", ix);
+	}
+
+      if (vp)
+	{
+	  int expect = vectors[0];
+	  if (vectors[ix] != expect)
+	    {
+	      exit = 1;
+	      printf ("vector %d not used %d times\n", ix, expect);
+	    }
+	}
+      else if (ix && vectors[ix])
+	{
+	  exit = 1;
+	  printf ("vector %d unexpectedly used\n", ix);
+	}
+      
+    }
+  return exit;
+}
+
+#pragma acc routine seq
+static int __attribute__((noinline)) place ()
+{
+  int r = 0;
+
+  if (acc_on_device (acc_device_nvidia))
+    {
+      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));
+      r = (g << 16) | (w << 8) | v;
+    }
+  return r;
+}
+
+static void clear (int *ary, int size)
+{
+  int ix;
+
+  for (ix = 0; ix < size; ix++)
+    ary[ix] = -1;
+}
+
+int vector_1 (int *ary, int size)
+{
+  clear (ary, size);
+  
+#pragma acc parallel num_workers (32) vector_length(32) copy(ary[0:size]) firstprivate (size)
+  {
+#pragma acc loop auto
+    for (int ix = 0; ix < size; ix++)
+      ary[ix] = place ();
+  }
+
+  return check (ary, size, 0, 0, 1);
+}
+
+int vector_2 (int *ary, int size)
+{
+  clear (ary, size);
+  
+#pragma acc parallel num_workers (32) vector_length(32) copy(ary[0:size]) firstprivate (size)
+  {
+#pragma acc loop worker
+    for (int jx = 0; jx <  size  / 64; jx++)
+#pragma acc loop auto
+      for (int ix = 0; ix < 64; ix++)
+	ary[ix + jx * 64] = place ();
+  }
+
+  return check (ary, size, 0, 1, 1);
+}
+
+int worker_1 (int *ary, int size)
+{
+  clear (ary, size);
+  
+#pragma acc parallel num_workers (32) vector_length(32) copy(ary[0:size]) firstprivate (size)
+  {
+#pragma acc loop auto
+    for (int jx = 0; jx <  size  / 64; jx++)
+#pragma acc loop vector
+      for (int ix = 0; ix < 64; ix++)
+	ary[ix + jx * 64] = place ();
+  }
+
+  return check (ary, size, 0, 1, 1);
+}
+
+int worker_2 (int *ary, int size)
+{
+  clear (ary, size);
+  
+#pragma acc parallel num_workers (32) vector_length(32) copy(ary[0:size]) firstprivate (size)
+  {
+#pragma acc loop auto
+    for (int jx = 0; jx <  size  / 64; jx++)
+#pragma acc loop auto
+      for (int ix = 0; ix < 64; ix++)
+	ary[ix + jx * 64] = place ();
+  }
+
+  return check (ary, size, 0, 1, 1);
+}
+
+int gang_1 (int *ary, int size)
+{
+  clear (ary, size);
+  
+#pragma acc parallel num_gangs (32) num_workers (32) vector_length(32) copy(ary[0:size]) firstprivate (size)
+  {
+#pragma acc loop auto
+    for (int jx = 0; jx <  size  / 64; jx++)
+#pragma acc loop worker
+      for (int ix = 0; ix < 64; ix++)
+	ary[ix + jx * 64] = place ();
+  }
+
+  return check (ary, size, 1, 1, 0);
+}
+
+int gang_2 (int *ary, int size)
+{
+  clear (ary, size);
+  
+#pragma acc parallel num_gangs (32) num_workers (32) vector_length(32) copy(ary[0:size]) firstprivate (size)
+  {
+#pragma acc loop auto
+    for (int kx = 0; kx < size / (32 * 32); kx++)
+#pragma acc loop auto
+      for (int jx = 0; jx <  32; jx++)
+#pragma acc loop auto
+	for (int ix = 0; ix < 32; ix++)
+	  ary[ix + jx * 32 + kx * 32 * 32] = place ();
+  }
+
+  return check (ary, size, 1, 1, 1);
+}
+
+#define N (32*32*32)
+int main ()
+{
+  int ondev = 0;
+
+#pragma acc parallel copy(ondev)
+  {
+    ondev = acc_on_device (acc_device_not_host);
+  }
+  if (!ondev)
+    return 0;
+  
+  int ary[N];
+
+  if (vector_1 (ary,  N))
+    return 1;
+  if (vector_2 (ary,  N))
+    return 1;
+
+  if (worker_1 (ary,  N))
+    return 1;
+  if (worker_2 (ary,  N))
+    return 1;
+  
+  if (gang_1 (ary,  N))
+    return 1;
+  if (gang_2 (ary,  N))
+    return 1;
+
+  return 0;
+}