[pr/69916] ICE with empty loops

[pr/69916] ICE with empty loops

[Nathan Sidwell]

[-- Attachment #1: Type: text/plain, Size: 1475 bytes --]

Jakub,
this patch fixes the ICE reported in pr69916 
(https://gcc.gnu.org/bugzilla/show_bug.cgi?id=69916)  The loop is lowered at 
omp-lowering, but subsequently determined to be dead before we get  to 
oacc-target-lower.  The loop CF is removed along with the (pure) IFN_OACC_LOOP 
function calls inserted during lowering.  However the IFN_UNIQUE loop head & 
tail calls remain (because they are not pure).  Thus in  the oacc-target-lower 
pass we rediscover the loop structure.

Firstly we assign a specific axis for this loop -- as it's auto.  That's a 
pessimization, but not wrong.  However, we then scan the  loop to adjust the 
expected OACC_LOOP calls with the determined partitioning information.  As 
they're not there, we end up falling out of the function and die with a 
single_succ_edge assert.  (In general we  might end up finding OACC_LOOP  calls 
of an inner loop, or meeting a block with more than one successor.  Either would 
be bad.)

This patch changes the loop transformation to count OACC_LOOP calls it 
encounters when rediscovering the loops, and uses that count for the OACC_LOOP 
adjustment scan (rather than expect OACC_LOOP_BOUND to be the last one).  That 
fixes the ICE.

While there it  is trivial to mark the loop as not to be partitioned, if we 
discover no OACC_LOOP calls, which addresses the pessimization mentioned above.

As the loop is no longer partitioned, the fork and join  markers, end up being 
deleted.

ok for trunk?

nathan

[-- Attachment #2: 69916.patch --]
[-- Type: text/x-patch, Size: 5843 bytes --]

2016-02-24  Nathan Sidwell  <nathan@codesourcery.com>

	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.

Index: omp-low.c
===================================================================
--- omp-low.c	(revision 233663)
+++ omp-low.c	(working copy)
@@ -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.  */
 };
 
@@ -20393,6 +20394,7 @@ new_oacc_loop_raw (oacc_loop *parent, lo
   loop->routine = NULL_TREE;
 
   loop->mask = loop->flags = 0;
+  loop->ifns = 0;
   loop->chunk_size = 0;
   loop->head_end = NULL;
 
@@ -20454,6 +20456,9 @@ new_oacc_loop_routine (oacc_loop *parent
 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;
 }
 
@@ -20584,43 +20589,54 @@ oacc_loop_discover_walk (oacc_loop *loop
       if (!gimple_call_internal_p (call))
 	continue;
 
-      if (gimple_call_internal_fn (call) != IFN_UNIQUE)
-	continue;
+      switch (gimple_call_internal_fn (call))
+	{
+	default:
+	  break;
 
-      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)
-	    {
-	      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));
+	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++;
 	    }
 	}
     }
@@ -20726,10 +20742,12 @@ oacc_loop_xform_head_tail (gcall *from,
    determined partitioning mask and chunking argument.  */
 
 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))
@@ -20749,13 +20767,13 @@ oacc_loop_xform_loop (gcall *end_marker,
 
 	  *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);
     }
@@ -20778,7 +20796,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++)
 	{
Index: testsuite/c-c++-common/goacc/pr69916.c
===================================================================
--- testsuite/c-c++-common/goacc/pr69916.c	(nonexistent)
+++ testsuite/c-c++-common/goacc/pr69916.c	(working copy)
@@ -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;
+}

Re: [pr/69916] ICE with empty loops

[Bernd Schmidt]

On 02/24/2016 09:07 PM, Nathan Sidwell wrote:
> this patch fixes the ICE reported in pr69916
> (https://gcc.gnu.org/bugzilla/show_bug.cgi?id=69916)  The loop is
> lowered at omp-lowering, but subsequently determined to be dead before
> we get  to oacc-target-lower.  The loop CF is removed along with the
> (pure) IFN_OACC_LOOP function calls inserted during lowering.  However
> the IFN_UNIQUE loop head & tail calls remain (because they are not
> pure).  Thus in  the oacc-target-lower pass we rediscover the loop
> structure.

> @@ -20726,10 +20742,12 @@ oacc_loop_xform_head_tail (gcall *from,
>      determined partitioning mask and chunking argument.  */
>
>   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)

Document the new arg.

> +  gcc_checking_assert (ifns);

I prefer "ifns != 0" and "ifns == 0" for non-booleans. I don't think 
it's a requirement, so your call.

> -      /* 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.  */

Lose the comma?

Ok with these changes.


Bernd