On 2016/8/15 5:57 PM, Jakub Jelinek wrote:
> On Mon, Aug 15, 2016 at 05:52:29PM +0800, Chung-Lin Tang wrote:
>> Hi Jakub,
>> This patch fixes an OpenACC reduction lowering segfault which
>> triggers when nested acc loop directives are present.
>> Cesar has reviewed this patch internally (since he mostly wrote
>> the code originally)
>>
>> Patch has been tested and committed to gomp-4_0-branch,
>> is this also okay for trunk?
>>
>> Thanks,
>> Chung-Lin
>>
>> 2016-08-15 Chung-Lin Tang <[email protected]>
>>
>> * omp-low.c (lower_oacc_reductions): Adjust variable lookup to use
>> maybe_lookup_decl, to handle nested acc loop directives.
>
> Is this covered by an existing testcase in the testsuite?
> If not, can you please add a testcase for it.
> Otherwise LGTM (not extra happy about accepting any kinds of contexts,
> but I hope the nesting diagnostics error out on OpenMP contexts mixed with
> OpenACC ones and hope that there can't be some other OpenACC context around
> that you wouldn't want to handle).
Thanks, I've committed the patch along with a new testcase (full patch as
attached).
Testcase is also backported to gomp-4_0-branch.
Thanks,
Chung-Lin
Index: omp-low.c
===================================================================
--- omp-low.c (revision 239529)
+++ omp-low.c (working copy)
@@ -5660,10 +5660,19 @@ lower_oacc_reductions (location_t loc, tree clause
outgoing = var;
incoming = omp_reduction_init_op (loc, rcode, type);
}
- else if (ctx->outer)
- incoming = outgoing = lookup_decl (orig, ctx->outer);
else
- incoming = outgoing = orig;
+ {
+ /* Try to look at enclosing contexts for reduction var,
+ use original if no mapping found. */
+ tree t = NULL_TREE;
+ omp_context *c = ctx->outer;
+ while (c && !t)
+ {
+ t = maybe_lookup_decl (orig, c);
+ c = c->outer;
+ }
+ incoming = outgoing = (t ? t : orig);
+ }
has_outer_reduction:;
}
Index: testsuite/c-c++-common/goacc/reduction-6.c
===================================================================
--- testsuite/c-c++-common/goacc/reduction-6.c (revision 0)
+++ testsuite/c-c++-common/goacc/reduction-6.c (revision 0)
@@ -0,0 +1,58 @@
+/* Check if different occurences of the reduction clause on
+ OpenACC loop nests will compile. */
+
+int foo (int N)
+{
+ int a = 0, b = 0, c = 0, d = 0, e = 0;
+
+ #pragma acc parallel
+ {
+ #pragma acc loop
+ for (int i = 0; i < N; i++)
+ {
+ #pragma acc loop reduction(+:a)
+ for (int j = 0; j < N; j++)
+ a += 1;
+ }
+ }
+
+ #pragma acc parallel
+ {
+ #pragma acc loop reduction(+:b)
+ for (int i = 0; i < N; i++)
+ {
+ #pragma acc loop
+ for (int j = 0; j < N; j++)
+ b += 1;
+ }
+ }
+
+ #pragma acc parallel
+ {
+ #pragma acc loop reduction(+:c)
+ for (int i = 0; i < N; i++)
+ {
+ #pragma acc loop reduction(+:c)
+ for (int j = 0; j < N; j++)
+ c += 1;
+ }
+ }
+
+ #pragma acc parallel loop
+ for (int i = 0; i < N; i++)
+ {
+ #pragma acc loop reduction(+:d)
+ for (int j = 0; j < N; j++)
+ d += 1;
+ }
+
+ #pragma acc parallel loop reduction(+:e)
+ for (int i = 0; i < N; i++)
+ {
+ #pragma acc loop reduction(+:e)
+ for (int j = 0; j < N; j++)
+ e += 1;
+ }
+
+ return a + b + c + d + e;
+}
