gimplify.c (gimplify_scan_omp_clauses): No special handling for OMP_CLAUSE_TILE.

2017-02-09  Nathan Sidwell  <nathan@codesourcery.com>
	    Cesar Philippidis  <cesar@codesourcery.com>
	    Joseph Myers  <joseph@codesourcery.com>
	    Chung-Lin Tang  <cltang@codesourcery.com>

	gcc/
	* gimplify.c (gimplify_scan_omp_clauses): No special handling for
	OMP_CLAUSE_TILE.
	(gimplify_adjust_omp_clauses): Don't delete TILE.
	(gimplify_omp_for): Deal with TILE.
	* internal-fn.c (expand_GOACC_TILE): New function.
	* internal-fn.def (GOACC_DIM_POS): Comment may be overly conservative.
	(GOACC_TILE): New.
	* omp-expand.c (struct oacc_collapse): Add tile and outer fields.
	(expand_oacc_collapse_init): Add LOC paramter.  Initialize tile
	element fields.
	(expand_oacc_collapse_vars): Add INNER parm, adjust for tiling,
	avoid DIV for outermost collapse var.
	(expand_oacc_for): Insert tile element loop as needed.  Adjust.
	Remove out of date comments, fix whitespace.
	* omp-general.c (omp_extract_for_data): Deal with tiling.
	* omp-general.h (enum oacc_loop_flags): Add OLF_TILE flag,
	adjust OLF_DIM_BASE value.
	(struct omp_for_data): Add tiling field.
	* omp-low.c (scan_sharing_clauses): Allow OMP_CLAUSE_TILE.
	(lower_oacc_head_mark): Add OLF_TILE as appropriate.  Ensure 2 levels
	for auto loops.  Remove default auto determining, moved to
	oacc_loop_fixed_partitions.
	* omp-offload.c (struct oacc_loop): Change 'ifns' to vector of call
	stmts, add e_mask field.
	(oacc_dim_call): New function, abstracted out from oacc_thread_numbers.
	(oacc_thread_numbers): Use oacc_dim_call.
	(oacc_xform_tile): New.
	(new_oacc_loop_raw): Initialize e_mask, adjust for ifns vector.
	(finish_oacc_loop): Adjust for ifns vector.
	(oacc_loop_discover_walk): Append loop abstraction sites to list,
	add case for GOACC_TILE fns.
	(oacc_loop_xform_loop): Delete.
	(oacc_loop_process): Iterate over call list directly, and add
	handling for GOACC_TILE fns.
	(oacc_loop_fixed_partitions): Determine default auto, deal with TILE,
	dump partitioning.
	(oacc_loop_auto_partitions): Add outer_assign parm. Assign all but
	vector partitioning to outer loops.  Assign 2 partitions to loops
	when available. Add TILE handling.
	(oacc_loop_partition): Adjust oacc_loop_auto_partitions call.
	(execite_oacc_device_lower): Process GOACC_TILE fns, ignore unknown specs.
	* tree-nested.c (convert_nonlocal_omp_clauses): Allow OMP_CLAUSE_TILE.
	* tree.c (omp_clause_num_ops): Adjust TILE ops.
	* tree.h (OMP_CLAUSE_TILE_ITERVAR, OMP_CLAUSE_TILE_COUNT): New.

	gcc/c/
	* c-parser.c (c_parser_omp_clause_collapse): Disallow tile.
	(c_parser_oacc_clause_tile): Disallow collapse. Fix parsing and
	semantic checking.
	* c-parser.c (c_parser_omp_for_loop): Accept tiling constructs.

	gcc/cp/
	* parser.c (cp_parser_oacc_clause_tile): Disallow collapse.  Fix
	parsing.  Parse constant expression. Remove semantic checking.
	(cp_parser_omp_clause_collapse): Disallow tile.
	(cp_parser_omp_for_loop): Deal with tile clause.  Don't emit a parse
	error about missing for after already emitting one.  Use more
	conventional for idiom for unbounded loop.
	* pt.c (tsubst_omp_clauses): Handle OMP_CLAUSE_TILE.
	* semantics.c (finish_omp_clauses): Correct TILE semantic check.
	(finish_omp_for): Deal with tile clause.

	gcc/fortran/
	* openmp.c (resolve_omp_clauses): Error on directives
	containing both tile and collapse clauses.
	(resolve_oacc_loop_blocks): Represent '*' tile arguments as zero.
	* trans-openmp.c (gfc_trans_omp_do): Lower tiled loops like
	collapsed loops.

	gcc/testsuite/
	* c-c++-common/goacc/combined-directives.c: Remove xfail.
	* c-c++-common/goacc/loop-auto-1.c: Adjust and add additional case.
	* c-c++-common/goacc/loop-auto-2.c: New.
	* c-c++-common/goacc/tile.c: Include stdbool, fix expected errors.
	* c-c++-common/goacc/tile-2.c: New.
	* g++.dg/goacc/template.C: Test tile subst.  Adjust erroneous uses.
	* g++.dg/goacc/tile-1.C: New, check tile subst.
	* gcc.dg/goacc/loop-processing-1.c: Adjust dg-final pattern.
	* gfortran.dg/goacc/combined-directives.f90: Remove xfail.
	* gfortran.dg/goacc/tile-1.f90: New test.
	* gfortran.dg/goacc/tile-2.f90: New test.
	* gfortran.dg/goacc/tile-lowering.f95: New test.

	libgomp/
	* testsuite/libgomp.oacc-c-c++-common/tile-1.c: New.
	* testsuite/libgomp.oacc-c-c++-common/loop-auto-1.c: Adjust and
	add additional case.
	* testsuite/libgomp.oacc-c-c++-common/vprop.c: XFAIL under
	"openacc_nvidia_accel_selected".
	* libgomp.oacc-fortran/nested-function-1.f90 (test2):
	Add num_workers(8) clause.

From-SVN: r245300

4 files changed, 303 insertions, 3 deletions

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
index 3ca9388..863b6b3 100644
--- a/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-auto-1.c
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-auto-1.c
@@ -112,7 +112,7 @@ int vector_1 (int *ary, int size)
 	ary[ix] = place ();
   }
 
-  return check (ary, size, 0, 0, 1);
+  return check (ary, size, 0, 1, 1);
 }
 
 int vector_2 (int *ary, int size)
@@ -196,10 +196,24 @@ int gang_3 (int *ary, int size)
 	ary[ix + jx * 64] = place ();
   }
 
+  return check (ary, size, 1, 1, 1);
+}
+
+int gang_4 (int *ary, int size)
+{
+  clear (ary, size);
+  
+#pragma acc parallel vector_length(32) copy(ary[0:size]) firstprivate (size)
+  {
+#pragma acc loop auto
+    for (int jx = 0; jx <  size; jx++)
+      ary[jx] = place ();
+  }
+
   return check (ary, size, 1, 0, 1);
 }
 
-#define N (32*32*32)
+#define N (32*32*32*2)
 int main ()
 {
   int ondev = 0;
@@ -227,6 +241,8 @@ int main ()
     return 1;
   if (gang_3 (ary,  N))
     return 1;
+  if (gang_4 (ary,  N))
+    return 1;
 
   return 0;
 }
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/tile-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/tile-1.c
new file mode 100644
index 0000000..8dcb956
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/tile-1.c
@@ -0,0 +1,281 @@
+/* This code uses nvptx inline assembly guarded with acc_on_device, which is
+   not optimized away at -O0, and then confuses the target assembler.
+   { dg-skip-if "" { *-*-* } { "-O0" } { "" } } */
+
+/* { dg-additional-options "-fopenacc-dim=32" } */
+
+#include <stdio.h>
+#include <openacc.h>
+
+static 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 gang_vector_1 (int *ary, int size)
+{
+  clear (ary, size);
+#pragma acc parallel vector_length(32) num_gangs (32) copy (ary[0:size]) firstprivate (size)
+  {
+#pragma acc loop tile(128) gang vector
+    for (int jx = 0; jx < size; jx++)
+      ary[jx] = place ();
+  }
+
+  return check (ary, size, 1, 0, 1);
+}
+
+int gang_vector_2a (int *ary, int size)
+{
+  if (size % 256)
+    return 1;
+  
+  clear (ary, size);
+#pragma acc parallel vector_length(32) num_gangs (32) copy (ary[0:size]) firstprivate (size)
+  {
+#pragma acc loop tile(64, 64) gang vector
+    for (int jx = 0; jx < size / 256; jx++)
+      for (int ix = 0; ix < 256; ix++)
+	ary[jx * 256 + ix] = place ();
+  }
+
+  return check (ary, size, 1, 0, 1);
+}
+
+int gang_vector_2b (int *ary, int size)
+{
+  if (size % 256)
+    return 1;
+  
+  clear (ary, size);
+#pragma acc parallel vector_length(32) num_gangs (32) copy (ary[0:size]) firstprivate (size)
+  {
+#pragma acc loop tile(64, 64) gang vector
+    for (int jx = 0; jx < size; jx += 256)
+      for (int ix = 0; ix < 256; ix++)
+	ary[jx + ix] = place ();
+  }
+
+  return check (ary, size, 1, 0, 1);
+}
+
+int worker_vector_2a (int *ary, int size)
+{
+  if (size % 256)
+    return 1;
+
+  clear (ary, size);
+#pragma acc parallel vector_length(32) num_workers (32) copy (ary[0:size]) firstprivate (size)
+  {
+#pragma acc loop tile(64, 64) worker vector
+    for (int jx = 0; jx < size / 256; jx++)
+      for (int ix = 0; ix < 256; ix++)
+	ary[jx * 256 + ix] = place ();
+  }
+
+  return check (ary, size, 0, 1, 1);
+}
+
+int worker_vector_2b (int *ary, int size)
+{
+  if (size % 256)
+    return 1;
+
+  clear (ary, size);
+#pragma acc parallel vector_length(32) num_workers (32) copy (ary[0:size]) firstprivate (size)
+  {
+#pragma acc loop tile(64, 64) worker vector
+    for (int jx = 0; jx < size; jx += 256)
+      for (int ix = 0; ix < 256; ix++)
+	ary[jx + ix] = place ();
+  }
+
+  return check (ary, size, 0, 1, 1);
+}
+
+int gang_worker_vector_2a (int *ary, int size)
+{
+  if (size % 256)
+    return 1;
+  clear (ary, size);
+#pragma acc parallel vector_length(32) num_workers (32) num_gangs(32) copy (ary[0:size]) firstprivate (size)
+  {
+#pragma acc loop tile(32, 32)
+    for (int jx = 0; jx < size / 256; jx++)
+      for (int ix = 0; ix < 256; ix++)
+	ary[jx * 256 + ix] = place ();
+  }
+
+  return check (ary, size, 1, 1, 1);
+}
+
+int gang_worker_vector_2b (int *ary, int size)
+{
+  if (size % 256)
+    return 1;
+  clear (ary, size);
+#pragma acc parallel vector_length(32) num_workers (32) num_gangs(32) copy (ary[0:size]) firstprivate (size)
+  {
+#pragma acc loop tile(32, 32)
+    for (int jx = 0; jx < size; jx += 256)
+      for (int ix = 0; ix < 256; ix++)
+	ary[jx + ix] = place ();
+  }
+
+  return check (ary, size, 1, 1, 1);
+}
+
+int gang_worker_vector_star_2a (int *ary, int size)
+{
+  if (size % 256)
+    return 1;
+
+  clear (ary, size);
+#pragma acc parallel vector_length(32) num_workers (32) num_gangs(32) copy (ary[0:size]) firstprivate (size)
+  {
+#pragma acc loop tile(*, *)
+    for (int jx = 0; jx < size / 256; jx++)
+      for (int ix = 0; ix < 256; ix++)
+	ary[jx * 256 + ix] = place ();
+  }
+
+  return check (ary, size, 1, 1, 1);
+}
+
+int gang_worker_vector_star_2b (int *ary, int size)
+{
+  if (size % 256)
+    return 1;
+
+  clear (ary, size);
+#pragma acc parallel vector_length(32) num_workers (32) num_gangs(32) copy (ary[0:size]) firstprivate (size)
+  {
+#pragma acc loop tile(*, *)
+    for (int jx = 0; jx < size; jx +=256)
+      for (int ix = 0; ix < 256; ix++)
+	ary[jx + ix] = place ();
+  }
+
+  return check (ary, size, 1, 1, 1);
+}
+
+#define N (32*32*32*8)
+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 (gang_vector_1 (ary, N))
+    return 1;
+  if (gang_vector_2a (ary, N))
+    return 1;
+  if (worker_vector_2a (ary, N))
+    return 1;
+  if (gang_worker_vector_2a (ary, N))
+    return 1;
+  if (gang_worker_vector_star_2a (ary, N))
+    return 1;
+  if (gang_vector_2b (ary, N))
+    return 1;
+  if (worker_vector_2b (ary, N))
+    return 1;
+  if (gang_worker_vector_2b (ary, N))
+    return 1;
+  if (gang_worker_vector_star_2b (ary, N))
+    return 1;
+  return 0;
+}
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/vprop.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/vprop.c
index 17b9568..c2bce82 100644
--- a/libgomp/testsuite/libgomp.oacc-c-c++-common/vprop.c
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/vprop.c
@@ -1,3 +1,6 @@
+/* { dg-do run } */
+/* { dg-xfail-run-if "PR78266" { openacc_nvidia_accel_selected } { "*" } { "" } } */
+
 #include <assert.h>
 
 #define test(type)				\
diff --git a/libgomp/testsuite/libgomp.oacc-fortran/nested-function-1.f90 b/libgomp/testsuite/libgomp.oacc-fortran/nested-function-1.f90
index fdbca44..c4af199 100644
--- a/libgomp/testsuite/libgomp.oacc-fortran/nested-function-1.f90
+++ b/libgomp/testsuite/libgomp.oacc-fortran/nested-function-1.f90
@@ -33,7 +33,7 @@ firstdo: do i = 1, 3
   subroutine test2
     integer :: a(3,3,3), k, kk, kkk, l, ll, lll
     a = 0
-    !$acc parallel
+    !$acc parallel num_workers(8)
     ! Use "gang(static:1)" here and below to effectively turn gang-redundant
     ! execution mode into something like gang-single.
     !$acc loop gang(static:1) collapse(1)