From mboxrd@z Thu Jan 1 00:00:00 1970 Return-Path: Received: by sourceware.org (Postfix, from userid 1534) id 1D45E3858D32; Mon, 26 Sep 2022 10:45:46 +0000 (GMT) DKIM-Filter: OpenDKIM Filter v2.11.0 sourceware.org 1D45E3858D32 DKIM-Signature: v=1; a=rsa-sha256; c=relaxed/relaxed; d=gcc.gnu.org; s=default; t=1664189146; bh=wIQNY7mbmiPJHZ5WyIpbemrmPry7Ds8msVKPRyLvT98=; h=From:To:Subject:Date:From; b=VdAKnkrhiivmXkN8a0JOZlSDuBmKQCgSgD4eKJid+nvy/eR+ixkiCQGLHPqWpLwOA Ik8JoXQPA9lpFjKcs8tKTgm+dgmJpnLIQPKOcWvEzI9AeQJr/DpHThTFkoGzI15HPv XfgmfAByPQ9CYnBfO/H5Z2Y2/buV5VGxyLWlkmFs= MIME-Version: 1.0 Content-Transfer-Encoding: 7bit Content-Type: text/plain; charset="utf-8" From: Tobias Burnus To: gcc-cvs@gcc.gnu.org Subject: [gcc r13-2868] OpenACC: Fix reduction tree-sharing issue [PR106982] X-Act-Checkin: gcc X-Git-Author: Tobias Burnus X-Git-Refname: refs/heads/master X-Git-Oldrev: 2387cfc7f6e2065f5c040d62763f6f3a5997a444 X-Git-Newrev: d3df98807b58df186061ad52ff87cc09ba593e9b Message-Id: <20220926104546.1D45E3858D32@sourceware.org> Date: Mon, 26 Sep 2022 10:45:46 +0000 (GMT) List-Id: https://gcc.gnu.org/g:d3df98807b58df186061ad52ff87cc09ba593e9b commit r13-2868-gd3df98807b58df186061ad52ff87cc09ba593e9b Author: Tobias Burnus 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 *)&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. 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 d9f9aaebc0b..dc42c752017 100644 --- a/gcc/omp-low.cc +++ b/gcc/omp-low.cc @@ -7631,6 +7631,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 @@ -7659,26 +7660,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) { } +}