re PR middle-end/69916 ([openacc] ICE in single_succ_edge called from oacc_loop_xform...

	gcc/
	PR middle-end/69916
	* omp-low.c (struct oacc_loop): Add ifns.
	(new_oacc_loop_raw): Initialize it.
	(finish_oacc_loop): Clear mask & flags if no ifns.
	(oacc_loop_discover_walk): Count IFN_GOACC_LOOP calls.
	(oacc_loop_xform_loop): Add ifns arg & adjust.
	(oacc_loop_process): Adjust oacc_loop_xform_loop call.

	gcc/testsuite/
	PR middle-end/69916
	* c-c-++-common/goacc/pr69916.c: New.

From-SVN: r234026

diff --git a/gcc/ChangeLog b/gcc/ChangeLog
index 0fe90a6ac10d96dea21f23a48ac9a3f183463d65..5b14b176aea622236482a858e09e437ddf2d1edb 100644 (file)
--- a/gcc/ChangeLog
+++ b/gcc/ChangeLog
@@ -1,3 +1,13 @@
+2016-03-07  Nathan Sidwell  <nathan@codesourcery.com>
+
+       PR middle-end/69916
+       * omp-low.c (struct oacc_loop): Add ifns.
+       (new_oacc_loop_raw): Initialize it.
+       (finish_oacc_loop): Clear mask & flags if no ifns.
+       (oacc_loop_discover_walk): Count IFN_GOACC_LOOP calls.
+       (oacc_loop_xform_loop): Add ifns arg & adjust.
+       (oacc_loop_process): Adjust oacc_loop_xform_loop call.
+
 2016-03-07  Richard Henderson  <rth@redhat.com>
 
        PR rtl-opt/70061

diff --git a/gcc/omp-low.c b/gcc/omp-low.c
index c69fe448da62f320fc91e95ce171d80f437575ef..82dec9d09b4d00737968552eadfaa5fa7d172dce 100644 (file)
--- a/gcc/omp-low.c
+++ b/gcc/omp-low.c
@@ -241,8 +241,9 @@ struct oacc_loop
   tree routine;  /* Pseudo-loop enclosing a routine.  */
 
   unsigned mask;   /* Partitioning mask.  */
-  unsigned flags;   /* Partitioning flags.  */
-  tree chunk_size;   /* Chunk size.  */
+  unsigned flags;  /* Partitioning flags.  */
+  unsigned ifns;   /* Contained loop abstraction functions.  */
+  tree chunk_size; /* Chunk size.  */
   gcall *head_end; /* Final marker of head sequence.  */
 };
 
@@ -20442,6 +20443,7 @@ new_oacc_loop_raw (oacc_loop *parent, location_t loc)
   loop->routine = NULL_TREE;
 
   loop->mask = loop->flags = 0;
+  loop->ifns = 0;
   loop->chunk_size = 0;
   loop->head_end = NULL;
 
@@ -20503,6 +20505,9 @@ new_oacc_loop_routine (oacc_loop *parent, gcall *call, tree decl, tree attrs)
 static oacc_loop *
 finish_oacc_loop (oacc_loop *loop)
 {
+  /* If the loop has been collapsed, don't partition it.  */
+  if (!loop->ifns)
+    loop->mask = loop->flags = 0;
   return loop->parent;
 }
 
@@ -20633,43 +20638,54 @@ oacc_loop_discover_walk (oacc_loop *loop, basic_block bb)
       if (!gimple_call_internal_p (call))
        continue;
 
-      if (gimple_call_internal_fn (call) != IFN_UNIQUE)
-       continue;
-
-      enum ifn_unique_kind kind
-       = (enum ifn_unique_kind) TREE_INT_CST_LOW (gimple_call_arg (call, 0));
-      if (kind == IFN_UNIQUE_OACC_HEAD_MARK
-         || kind == IFN_UNIQUE_OACC_TAIL_MARK)
+      switch (gimple_call_internal_fn (call))
        {
-         if (gimple_call_num_args (call) == 2)
-           {
-             gcc_assert (marker && !remaining);
-             marker = 0;
-             if (kind == IFN_UNIQUE_OACC_TAIL_MARK)
-               loop = finish_oacc_loop (loop);
-             else
-               loop->head_end = call;
-           }
-         else
-           {
-             int count = TREE_INT_CST_LOW (gimple_call_arg (call, 2));
+       default:
+         break;
+
+       case IFN_GOACC_LOOP:
+         /* Count the goacc loop abstraction fns, to determine if the
+            loop was collapsed already.  */
+         loop->ifns++;
+         break;
 
-             if (!marker)
+       case IFN_UNIQUE:
+         enum ifn_unique_kind kind
+           = (enum ifn_unique_kind) (TREE_INT_CST_LOW
+                                     (gimple_call_arg (call, 0)));
+         if (kind == IFN_UNIQUE_OACC_HEAD_MARK
+             || kind == IFN_UNIQUE_OACC_TAIL_MARK)
+           {
+             if (gimple_call_num_args (call) == 2)
                {
-                 if (kind == IFN_UNIQUE_OACC_HEAD_MARK)
-                   loop = new_oacc_loop (loop, call);
-                 remaining = count;
+                 gcc_assert (marker && !remaining);
+                 marker = 0;
+                 if (kind == IFN_UNIQUE_OACC_TAIL_MARK)
+                   loop = finish_oacc_loop (loop);
+                 else
+                   loop->head_end = call;
                }
-             gcc_assert (count == remaining);
-             if (remaining)
+             else
                {
-                 remaining--;
-                 if (kind == IFN_UNIQUE_OACC_HEAD_MARK)
-                   loop->heads[marker] = call;
-                 else
-                   loop->tails[remaining] = call;
+                 int count = TREE_INT_CST_LOW (gimple_call_arg (call, 2));
+
+                 if (!marker)
+                   {
+                     if (kind == IFN_UNIQUE_OACC_HEAD_MARK)
+                       loop = new_oacc_loop (loop, call);
+                     remaining = count;
+                   }
+                 gcc_assert (count == remaining);
+                 if (remaining)
+                   {
+                     remaining--;
+                     if (kind == IFN_UNIQUE_OACC_HEAD_MARK)
+                       loop->heads[marker] = call;
+                     else
+                       loop->tails[remaining] = call;
+                   }
+                 marker++;
                }
-             marker++;
            }
        }
     }
@@ -20772,13 +20788,19 @@ oacc_loop_xform_head_tail (gcall *from, int level)
 }
 
 /* Transform the IFN_GOACC_LOOP internal functions by providing the
-   determined partitioning mask and chunking argument.  */
+   determined partitioning mask and chunking argument.  END_MARKER
+   points at the end IFN_HEAD_TAIL call intgroducing the loop.  IFNS
+   is the number of IFN_GOACC_LOOP calls for the loop.  MASK_ARG is
+   the replacement partitioning mask and CHUNK_ARG is the replacement
+   chunking arg.  */
 
 static void
-oacc_loop_xform_loop (gcall *end_marker, tree mask_arg, tree chunk_arg)
+oacc_loop_xform_loop (gcall *end_marker, unsigned ifns,
+                     tree mask_arg, tree chunk_arg)
 {
   gimple_stmt_iterator gsi = gsi_for_stmt (end_marker);
   
+  gcc_checking_assert (ifns);
   for (;;)
     {
       for (; !gsi_end_p (gsi); gsi_next (&gsi))
@@ -20798,13 +20820,13 @@ oacc_loop_xform_loop (gcall *end_marker, tree mask_arg, tree chunk_arg)
 
          *gimple_call_arg_ptr (call, 5) = mask_arg;
          *gimple_call_arg_ptr (call, 4) = chunk_arg;
-         if (TREE_INT_CST_LOW (gimple_call_arg (call, 0))
-             == IFN_GOACC_LOOP_BOUND)
+         ifns--;
+         if (!ifns)
            return;
        }
 
-      /* If we didn't see LOOP_BOUND, it should be in the single
-        successor block.  */
+      /* The LOOP_BOUND ifn could be in the single successor
+        block.  */
       basic_block bb = single_succ (gsi_bb (gsi));
       gsi = gsi_start_bb (bb);
     }
@@ -20827,7 +20849,7 @@ oacc_loop_process (oacc_loop *loop)
       tree mask_arg = build_int_cst (unsigned_type_node, mask);
       tree chunk_arg = loop->chunk_size;
 
-      oacc_loop_xform_loop (loop->head_end, mask_arg, chunk_arg);
+      oacc_loop_xform_loop (loop->head_end, loop->ifns, mask_arg, chunk_arg);
 
       for (ix = 0; ix != GOMP_DIM_MAX && loop->heads[ix]; ix++)
        {

diff --git a/gcc/testsuite/ChangeLog b/gcc/testsuite/ChangeLog
index 952bfeb700aa808527c026b56fa005f328410f59..7d4d34249903d6918bd6e9ab361ddc0911438f57 100644 (file)
--- a/gcc/testsuite/ChangeLog
+++ b/gcc/testsuite/ChangeLog
@@ -1,3 +1,8 @@
+2016-03-07  Nathan Sidwell  <nathan@codesourcery.com>
+
+       PR middle-end/69916
+       * c-c-++-common/goacc/pr69916.c: New.
+
 2016-03-07  Richard Henderson  <rth@redhat.com>
 
        * gcc.c-torture/compile/pr70061.c: New test.
@@ -711,7 +716,6 @@
        PR c/69900
        * gcc.dg/pr69900.c: New test.
 
->>>>>>> .r233653
 2016-02-23  Martin Jambor  <mjambor@suse.cz>
 
        PR tree-optimization/69666

diff --git a/gcc/testsuite/c-c++-common/goacc/pr69916.c b/gcc/testsuite/c-c++-common/goacc/pr69916.c
new file mode 100644 (file)
index 0000000..e037af3
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/goacc/pr69916.c
@@ -0,0 +1,20 @@
+/* {  dg-additional-options "-O2" } */
+
+/* PR 69916, an loop determined to be empty sometime after omp-lower
+   and before oacc-device-lower can evaporate leading to no GOACC_LOOP
+   internal functions existing.  */
+
+int
+main (void)
+{
+
+#pragma acc parallel
+  {
+    int j = 0;
+#pragma acc loop private (j)
+    for (int i = 0; i < 10; i++)
+      j++;
+  }
+
+  return 0;
+}