Re: [patch, OpenACC] Fix reduction lowering segfault in omp-low

[Chung-Lin Tang]

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  <clt...@codesourcery.com>
>>
>>         * 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;
+}