public inbox for gcc-cvs@sourceware.org
help / color / mirror / Atom feed
* [gcc r12-8793] OpenACC: Fix reduction tree-sharing issue [PR106982]
@ 2022-09-29 11:48 Tobias Burnus
0 siblings, 0 replies; only message in thread
From: Tobias Burnus @ 2022-09-29 11:48 UTC (permalink / raw)
To: gcc-cvs
https://gcc.gnu.org/g:afea1ae84f0c3e64137289c44c756195babe1845
commit r12-8793-gafea1ae84f0c3e64137289c44c756195babe1845
Author: Tobias Burnus <tobias@codesourcery.com>
Date: Mon Sep 26 12:45:28 2022 +0200
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.
(cherry picked from commit d3df98807b58df186061ad52ff87cc09ba593e9b)
Diff:
---
gcc/omp-low.cc | 19 +++++++++++--------
gcc/testsuite/c-c++-common/goacc/reduction-7.c | 22 ++++++++++++++++++++++
gcc/testsuite/c-c++-common/goacc/reduction-8.c | 12 ++++++++++++
3 files changed, 45 insertions(+), 8 deletions(-)
diff --git a/gcc/omp-low.cc b/gcc/omp-low.cc
index 584ec09d31b..e9482f2363b 100644
--- a/gcc/omp-low.cc
+++ b/gcc/omp-low.cc
@@ -7581,6 +7581,7 @@ lower_oacc_reductions (location_t loc, tree clauses, tree level, bool inner,
incoming = build_simple_mem_ref (incoming);
}
else
+ /* Note that 'var' might be a mem ref. */
v1 = v2 = v3 = var;
/* Determine position in reduction buffer, which may be used
@@ -7609,26 +7610,28 @@ 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);
- gimplify_assign (v3, fini_call, &before_join);
- gimplify_assign (outgoing, teardown_call, &after_join);
+ gimplify_assign (unshare_expr (v1), setup_call, &before_fork);
+ gimplify_assign (unshare_expr (v2), init_call, &after_fork);
+ gimplify_assign (unshare_expr (v3), fini_call, &before_join);
+ gimplify_assign (unshare_expr (outgoing), teardown_call, &after_join);
}
/* Now stitch things together. */
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) { }
+}
^ permalink raw reply [flat|nested] only message in thread
only message in thread, other threads:[~2022-09-29 11:48 UTC | newest]
Thread overview: (only message) (download: mbox.gz / follow: Atom feed)
-- links below jump to the message on this page --
2022-09-29 11:48 [gcc r12-8793] OpenACC: Fix reduction tree-sharing issue [PR106982] Tobias Burnus
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).