diff mbox

[OpenACC] Fix reduction lowering segfault in omp-low

Message ID 1658f83f-e37e-ca90-df15-d8046f597090@codesourcery.com
State New
Headers show

Commit Message

Chung-Lin Tang Aug. 17, 2016, 12:17 p.m. UTC 
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  <cltang@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
diff mbox

Patch

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