public inbox for gcc-patches@gcc.gnu.org
 help / color / mirror / Atom feed
From: Tobias Burnus <tobias@codesourcery.com>
To: gcc-patches <gcc-patches@gcc.gnu.org>,
	Thomas Schwinge <thomas@codesourcery.com>
Subject: [Patch] OpenACC: Fix reduction tree-sharing issue [PR106982]
Date: Fri, 23 Sep 2022 17:24:23 +0200	[thread overview]
Message-ID: <b694809c-c969-1d8f-196b-589194312c02@codesourcery.com> (raw)


[-- Attachment #1.1: Type: text/plain, Size: 837 bytes --]

This fixes a tree-sharing ICE. It seems as if all unshare_expr
I added were required in this case. The first long testcase is
based on the real testcase from the OpenACC testsuite, the second
one is what reduction produced - but I thought some nested reduction
might be interesting as well; hence, I included both tests.


Bootstrapped and regtested on x86-64-gnu-linux w/o offloading.
OK for mainline and GCC 12?

(It gives an ICE with GCC 10 but not with GCC 9; thus,
more regression-fix backporting would be possible,
if someone cares.)

Tobias


-----------------
Siemens Electronic Design Automation GmbH; Anschrift: Arnulfstraße 201, 80634 München; Gesellschaft mit beschränkter Haftung; Geschäftsführer: Thomas Heurung, Frank Thürauf; Sitz der Gesellschaft: München; Registergericht München, HRB 106955

[-- Attachment #2: oacc-reduction.diff --]
[-- Type: text/x-patch, Size: 3744 bytes --]

OpenACC: Fix reduction tree-sharing issue [PR106982]

The tree for var == incoming == outgound was
'MEM <double[5]> [(double *)&reduced]' which caused the ICE
"incorrect sharing of tree nodes".

	PR middle-end/106982

gcc/ChangeLog:

	* omp-low.cc (lower_oacc_reductions): Add some unshare_expr.

gcc/testsuite/ChangeLog:

	* c-c++-common/goacc/reduction-7.c: New test.
	* c-c++-common/goacc/reduction-8.c: New test.

 gcc/omp-low.cc                                 | 17 ++++++++++++-----
 gcc/testsuite/c-c++-common/goacc/reduction-7.c | 22 ++++++++++++++++++++++
 gcc/testsuite/c-c++-common/goacc/reduction-8.c | 12 ++++++++++++
 3 files changed, 46 insertions(+), 5 deletions(-)

diff --git a/gcc/omp-low.cc b/gcc/omp-low.cc
index f0469d20b3d..8e07fb5d8a8 100644
--- a/gcc/omp-low.cc
+++ b/gcc/omp-low.cc
@@ -7631,7 +7631,12 @@ lower_oacc_reductions (location_t loc, tree clauses, tree level, bool inner,
 	      incoming = build_simple_mem_ref (incoming);
 	  }
 	else
-	  v1 = v2 = v3 = var;
+	  {
+	    v1 = unshare_expr (var);
+	    v2 = unshare_expr (var);
+	    v3 = unshare_expr (var);
+	    outgoing = unshare_expr (outgoing);
+	  }
 
 	/* Determine position in reduction buffer, which may be used
 	   by target.  The parser has ensured that this is not a
@@ -7659,21 +7664,23 @@ lower_oacc_reductions (location_t loc, tree clauses, tree level, bool inner,
 	  = build_call_expr_internal_loc (loc, IFN_GOACC_REDUCTION,
 					  TREE_TYPE (var), 6, setup_code,
 					  unshare_expr (ref_to_res),
-					  incoming, level, op, off);
+					  unshare_expr (incoming), level,
+					  op, off);
 	tree init_call
 	  = build_call_expr_internal_loc (loc, IFN_GOACC_REDUCTION,
 					  TREE_TYPE (var), 6, init_code,
 					  unshare_expr (ref_to_res),
-					  v1, level, op, off);
+					  unshare_expr (v1), level, op, off);
 	tree fini_call
 	  = build_call_expr_internal_loc (loc, IFN_GOACC_REDUCTION,
 					  TREE_TYPE (var), 6, fini_code,
 					  unshare_expr (ref_to_res),
-					  v2, level, op, off);
+					  unshare_expr (v2), level, op, off);
 	tree teardown_call
 	  = build_call_expr_internal_loc (loc, IFN_GOACC_REDUCTION,
 					  TREE_TYPE (var), 6, teardown_code,
-					  ref_to_res, v3, level, op, off);
+					  ref_to_res, unshare_expr (v3),
+					  level, op, off);
 
 	gimplify_assign (v1, setup_call, &before_fork);
 	gimplify_assign (v2, init_call, &after_fork);
diff --git a/gcc/testsuite/c-c++-common/goacc/reduction-7.c b/gcc/testsuite/c-c++-common/goacc/reduction-7.c
new file mode 100644
index 00000000000..482b0ab1984
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/goacc/reduction-7.c
@@ -0,0 +1,22 @@
+/* { dg-do compile } */
+
+/* PR middle-end/106982 */
+
+long long n = 100;
+int multiplicitive_n = 128;
+
+void test1(double *rand, double *a, double *b, double *c)
+{
+#pragma acc data copyin(a[0:10*multiplicitive_n], b[0:10*multiplicitive_n]) copyout(c[0:10])
+    {
+#pragma acc parallel loop
+        for (int i = 0; i < 10; ++i)
+        {
+        double temp = 1.0;
+#pragma acc loop vector reduction(*:temp)
+        for (int j = 0; j < multiplicitive_n; ++j)
+          temp *= a[(i * multiplicitive_n) + j] + b[(i * multiplicitive_n) + j];
+        c[i] = temp;
+        }
+    }
+}
diff --git a/gcc/testsuite/c-c++-common/goacc/reduction-8.c b/gcc/testsuite/c-c++-common/goacc/reduction-8.c
new file mode 100644
index 00000000000..2c3ed499d5b
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/goacc/reduction-8.c
@@ -0,0 +1,12 @@
+/* { dg-do compile } */
+
+/* PR middle-end/106982 */
+
+void test1(double *c)
+{
+    double reduced[5];
+#pragma acc parallel loop gang private(reduced)
+    for (int x = 0; x < 5; ++x)
+#pragma acc loop worker reduction(*:reduced)
+      for (int y = 0; y < 5; ++y) { }
+}

             reply	other threads:[~2022-09-23 15:24 UTC|newest]

Thread overview: 5+ messages / expand[flat|nested]  mbox.gz  Atom feed  top
2022-09-23 15:24 Tobias Burnus [this message]
2022-09-26  8:32 ` Richard Biener
2022-09-26  9:27   ` Tobias Burnus
2022-09-26  9:34     ` Richard Biener
2022-09-26 11:16       ` Thomas Schwinge

Reply instructions:

You may reply publicly to this message via plain-text email
using any one of the following methods:

* Save the following mbox file, import it into your mail client,
  and reply-to-all from there: mbox

  Avoid top-posting and favor interleaved quoting:
  https://en.wikipedia.org/wiki/Posting_style#Interleaved_style

* Reply using the --to, --cc, and --in-reply-to
  switches of git-send-email(1):

  git send-email \
    --in-reply-to=b694809c-c969-1d8f-196b-589194312c02@codesourcery.com \
    --to=tobias@codesourcery.com \
    --cc=gcc-patches@gcc.gnu.org \
    --cc=thomas@codesourcery.com \
    /path/to/YOUR_REPLY

  https://kernel.org/pub/software/scm/git/docs/git-send-email.html

* If your mail client supports setting the In-Reply-To header
  via mailto: links, try the mailto: link
Be sure your reply has a Subject: header at the top and a blank line before the message body.
This is a public inbox, see mirroring instructions
for how to clone and mirror all data and code used for this inbox;
as well as URLs for read-only IMAP folder(s) and NNTP newsgroup(s).