From patchwork Fri Sep 23 15:24:23 2022 Content-Type: text/plain; charset="utf-8" MIME-Version: 1.0 Content-Transfer-Encoding: 8bit X-Patchwork-Submitter: Tobias Burnus X-Patchwork-Id: 1422 Return-Path: Delivered-To: ouuuleilei@gmail.com Received: by 2002:a5d:5044:0:0:0:0:0 with SMTP id h4csp277046wrt; Fri, 23 Sep 2022 08:25:11 -0700 (PDT) X-Google-Smtp-Source: AMsMyM6W8VNQTh9pkhv1i0/Cps0zKvMiC66nY9an+lFmc1itZvLE2NkqCWONCR7kS/CtW1IuaEiC X-Received: by 2002:a17:907:7610:b0:77c:b7a:9fb2 with SMTP id jx16-20020a170907761000b0077c0b7a9fb2mr7277992ejc.468.1663946711076; Fri, 23 Sep 2022 08:25:11 -0700 (PDT) ARC-Seal: i=1; a=rsa-sha256; t=1663946711; cv=none; d=google.com; s=arc-20160816; b=wCppBfbxTn4IBh+2HbiE42krW5rbCU2KMoVAAkoMcpeCr2Wll2Nbxf9BwFYfGRZNWk uA9Mpcz9NVwhyPu0lTjVul2HLVlgnra7RKA5qD55dwF6Zob9JX2Sx1JVX0+TJpSkDHrg Wbue4pilndVSgKv5Gkhh0xVZ70uHOKi5Le1SDplQXFsLuKRsP29nnNfO1vZj5ULWyK4p cRNPfK3BDQfvuLKpu0rCZe3yCjd8Md6kUwy5ad9bRbExvoOhv0SrQr/yDD6n+xj0cexw t9NfxH8YmlIvOLYVqYcujQkQhAqlFa077m3TFtgOjIZZ+YD4dLAvoL6Nmha9cqNuTupK qpBA== ARC-Message-Signature: i=1; a=rsa-sha256; c=relaxed/relaxed; d=google.com; s=arc-20160816; h=sender:errors-to:list-subscribe:list-help:list-post:list-archive :list-unsubscribe:list-id:precedence:subject:from:to :content-language:user-agent:mime-version:date:message-id :ironport-sdr:dmarc-filter:delivered-to; bh=tXv4RV5VhtzbIRZlYXJrMXTZl6LTJPqc9iMhN9jYudQ=; b=XKUZ0WtCbdHf0a0gSUyONXCEVV7uRBAde4c6D7VJcI/D8M+vLRymhpwBOf35obsY35 qByQzDEcf8yD/hlglTnrTL0XOHC2G+1yJf1HNNLscDV0dMICm2Jj01V8x2np+5F04Y5R dkZj13asfcmHTOqvzF6Ce2l90FJU/jczdnW9siG3FKpHSaQX9qmKKrLEwNo4izy1cE7a mDrmeAwanSHe9QciW5XRsJzYVyZFc4dqkxDSk7wZapMoOZN99GLn14cvMWuWkhUgDJOQ QVKGOoXMpklY2zhUmX8JCPd4paMPRqZ6QA6VBy3t2zwCb1SiKiTNe14FAJ1RcZwqvNKB 46+g== ARC-Authentication-Results: i=1; mx.google.com; spf=pass (google.com: domain of gcc-patches-bounces+ouuuleilei=gmail.com@gcc.gnu.org designates 8.43.85.97 as permitted sender) smtp.mailfrom="gcc-patches-bounces+ouuuleilei=gmail.com@gcc.gnu.org" Received: from sourceware.org (ip-8-43-85-97.sourceware.org. [8.43.85.97]) by mx.google.com with ESMTPS id qf38-20020a1709077f2600b00781f51771b6si8750273ejc.900.2022.09.23.08.25.10 for (version=TLS1_3 cipher=TLS_AES_256_GCM_SHA384 bits=256/256); Fri, 23 Sep 2022 08:25:11 -0700 (PDT) Received-SPF: pass (google.com: domain of gcc-patches-bounces+ouuuleilei=gmail.com@gcc.gnu.org designates 8.43.85.97 as permitted sender) client-ip=8.43.85.97; Authentication-Results: mx.google.com; spf=pass (google.com: domain of gcc-patches-bounces+ouuuleilei=gmail.com@gcc.gnu.org designates 8.43.85.97 as permitted sender) smtp.mailfrom="gcc-patches-bounces+ouuuleilei=gmail.com@gcc.gnu.org" Received: from server2.sourceware.org (localhost [IPv6:::1]) by sourceware.org (Postfix) with ESMTP id 78A013856091 for ; Fri, 23 Sep 2022 15:24:55 +0000 (GMT) X-Original-To: gcc-patches@gcc.gnu.org Delivered-To: gcc-patches@gcc.gnu.org Received: from esa2.mentor.iphmx.com (esa2.mentor.iphmx.com [68.232.141.98]) by sourceware.org (Postfix) with ESMTPS id 5A0D83857351 for ; Fri, 23 Sep 2022 15:24:29 +0000 (GMT) DMARC-Filter: OpenDMARC Filter v1.4.1 sourceware.org 5A0D83857351 Authentication-Results: sourceware.org; dmarc=none (p=none dis=none) header.from=codesourcery.com Authentication-Results: sourceware.org; spf=pass smtp.mailfrom=mentor.com X-IronPort-AV: E=Sophos;i="5.93,339,1654588800"; d="diff'?scan'208,217";a="83588064" Received: from orw-gwy-02-in.mentorg.com ([192.94.38.167]) by esa2.mentor.iphmx.com with ESMTP; 23 Sep 2022 07:24:27 -0800 IronPort-SDR: 4kJ3hvmI38WIUqAW2aRGDeJGsyScssHNB0zM2ua/nBElZJLky73P6ffvBqa4AKPcNxLv0Xk6jo bmdsUUATMTkgdYfbvexIwt5txr8VEvcit1+cbT1Uj9J+lApLUpBh5oOOHUPUGPiSq3J1zICiLD 7nRebLKnEC8nhLQxCvFYVCMXoe9bqhQthSX8EgLQxDbfWv6eV23PUQaysTOig3v7qchVKLrm5h 8vzDu0ez7TVH8s5cCJxX0V+k51EuSSuG8jCZppp+Ycju/j8wGXoH3SIUz+BXiqb94gJ4hi9bti D/k= Message-ID: Date: Fri, 23 Sep 2022 17:24:23 +0200 MIME-Version: 1.0 User-Agent: Mozilla/5.0 (X11; Linux x86_64; rv:102.0) Gecko/20100101 Thunderbird/102.3.0 Content-Language: en-US To: gcc-patches , Thomas Schwinge From: Tobias Burnus Subject: [Patch] OpenACC: Fix reduction tree-sharing issue [PR106982] X-Originating-IP: [137.202.0.90] X-ClientProxiedBy: svr-ies-mbx-11.mgc.mentorg.com (139.181.222.11) To svr-ies-mbx-12.mgc.mentorg.com (139.181.222.12) X-Spam-Status: No, score=-11.4 required=5.0 tests=BAYES_00, GIT_PATCH_0, HEADER_FROM_DIFFERENT_DOMAINS, HTML_MESSAGE, KAM_DMARC_STATUS, RCVD_IN_MSPIKE_H2, SPF_HELO_PASS, SPF_PASS, TXREP autolearn=ham autolearn_force=no version=3.4.6 X-Spam-Checker-Version: SpamAssassin 3.4.6 (2021-04-09) on server2.sourceware.org X-Content-Filtered-By: Mailman/MimeDel 2.1.29 X-BeenThere: gcc-patches@gcc.gnu.org X-Mailman-Version: 2.1.29 Precedence: list List-Id: Gcc-patches mailing list List-Unsubscribe: , List-Archive: List-Post: List-Help: List-Subscribe: , Errors-To: gcc-patches-bounces+ouuuleilei=gmail.com@gcc.gnu.org Sender: "Gcc-patches" X-getmail-retrieved-from-mailbox: =?utf-8?q?INBOX?= X-GMAIL-THRID: =?utf-8?q?1744774587011189895?= X-GMAIL-MSGID: =?utf-8?q?1744774587011189895?= 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 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. 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) { } +}