diff mbox

[pr/69916] ICE with empty loops

Message ID 56CE0D9D.8030704@acm.org
State New
Headers show

Commit Message

Nathan Sidwell Feb. 24, 2016, 8:07 p.m. UTC 
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

Comments

Bernd Schmidt March 4, 2016, 6:55 p.m. UTC | #1 
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
diff mbox

Patch

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