From mboxrd@z Thu Jan 1 00:00:00 1970 Return-Path: Received: by sourceware.org (Postfix, from userid 1534) id 5BD4C3858D3C; Thu, 29 Sep 2022 11:48:28 +0000 (GMT) DKIM-Filter: OpenDKIM Filter v2.11.0 sourceware.org 5BD4C3858D3C DKIM-Signature: v=1; a=rsa-sha256; c=relaxed/relaxed; d=gcc.gnu.org; s=default; t=1664452108; bh=xpvdqZQI8KIDugJwPRLXm6esbLaYfqugjC9ekPbQd64=; h=From:To:Subject:Date:From; b=b9DoXmF5Tph/VzbWajztT+O2AGixKwKFLNbo8H5CqPT5S90+cFnR++lyAPdCqIOYX Cdt9cyT9vgTDLFaGOXBOFFrRURTGNBA+Ce19MAxSYpI25dpBeBsXlSU+Q3MHUALVvS gQjXd21qx7BstI0FWZjM4/3i01cdYkooA1guE41E= 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 r12-8793] OpenACC: Fix reduction tree-sharing issue [PR106982] X-Act-Checkin: gcc X-Git-Author: Tobias Burnus X-Git-Refname: refs/heads/releases/gcc-12 X-Git-Oldrev: 7f323c284f785397c9794c1936aacbcd7cc63ba8 X-Git-Newrev: afea1ae84f0c3e64137289c44c756195babe1845 Message-Id: <20220929114828.5BD4C3858D3C@sourceware.org> Date: Thu, 29 Sep 2022 11:48:28 +0000 (GMT) List-Id: https://gcc.gnu.org/g:afea1ae84f0c3e64137289c44c756195babe1845 commit r12-8793-gafea1ae84f0c3e64137289c44c756195babe1845 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. (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) { } +}