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
2016-02-24 Nathan Sidwell <nat...@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;
+}
