From patchwork Tue Dec 6 17:20:24 2016 Content-Type: text/plain; charset="utf-8" MIME-Version: 1.0 Content-Transfer-Encoding: 7bit X-Patchwork-Submitter: Cesar Philippidis X-Patchwork-Id: 86878 Delivered-To: patch@linaro.org Received: by 10.140.20.101 with SMTP id 92csp2130303qgi; Tue, 6 Dec 2016 09:20:56 -0800 (PST) X-Received: by 10.84.218.13 with SMTP id q13mr137738503pli.168.1481044856470; Tue, 06 Dec 2016 09:20:56 -0800 (PST) Return-Path: Received: from sourceware.org (server1.sourceware.org. [209.132.180.131]) by mx.google.com with ESMTPS id i124si20223328pgd.33.2016.12.06.09.20.56 for (version=TLS1_2 cipher=ECDHE-RSA-AES128-GCM-SHA256 bits=128/128); Tue, 06 Dec 2016 09:20:56 -0800 (PST) Received-SPF: pass (google.com: domain of gcc-patches-return-443618-patch=linaro.org@gcc.gnu.org designates 209.132.180.131 as permitted sender) client-ip=209.132.180.131; Authentication-Results: mx.google.com; dkim=pass header.i=@gcc.gnu.org; spf=pass (google.com: domain of gcc-patches-return-443618-patch=linaro.org@gcc.gnu.org designates 209.132.180.131 as permitted sender) smtp.mailfrom=gcc-patches-return-443618-patch=linaro.org@gcc.gnu.org DomainKey-Signature: a=rsa-sha1; c=nofws; d=gcc.gnu.org; h=list-id :list-unsubscribe:list-archive:list-post:list-help:sender:from :subject:to:message-id:date:mime-version:content-type; q=dns; s= default; b=q1KgnEatZSjYcZBbSpA2urlFMSmvieY94SFdFTQol5PgFWNkRqEjm d6gpcx8VzaDAq9emVdl4QiBQk5LgBFzgl34BLguqLjpeGPHCMX5TXjV6UZjfZSzV tfewK2MYrQ8wVQde3iVq3P+S2wNM67rPrqW6y51q577udR+URg5yvM= DKIM-Signature: v=1; a=rsa-sha1; c=relaxed; d=gcc.gnu.org; h=list-id :list-unsubscribe:list-archive:list-post:list-help:sender:from :subject:to:message-id:date:mime-version:content-type; s= default; bh=/Uqgy0Y87E2oBcWFjKVgO1ppPEI=; b=hkJmGhid8rqkZ5yFESvw wQ9vgmhOY23a1915L4gyRf9Iekx+FMEeTjs0rXne7CM5rtTRMQwDhqJow4ZDfEiB wwdRbcKKy+Q+ymIPsGbpeL4VucoKKcF1T3zDpaMnIreIMswp5y6ihbTA5kJOvVYG 1PsmpQz7gUO72GHH41xuMNo= Received: (qmail 122982 invoked by alias); 6 Dec 2016 17:20:42 -0000 Mailing-List: contact gcc-patches-help@gcc.gnu.org; run by ezmlm Precedence: bulk List-Id: List-Unsubscribe: List-Archive: List-Post: List-Help: Sender: gcc-patches-owner@gcc.gnu.org Delivered-To: mailing list gcc-patches@gcc.gnu.org Received: (qmail 122960 invoked by uid 89); 6 Dec 2016 17:20:41 -0000 Authentication-Results: sourceware.org; auth=none X-Virus-Found: No X-Spam-SWARE-Status: No, score=-1.9 required=5.0 tests=AWL, BAYES_00, RCVD_IN_DNSWL_NONE, SPF_PASS, URIBL_RED autolearn=ham version=3.3.2 spammy=Hx-languages-length:2431, sk:cesarc, decl_size, cesar@codesourcery.com X-HELO: relay1.mentorg.com Received: from relay1.mentorg.com (HELO relay1.mentorg.com) (192.94.38.131) by sourceware.org (qpsmtpd/0.93/v0.84-503-g423c35a) with ESMTP; Tue, 06 Dec 2016 17:20:29 +0000 Received: from svr-orw-mbx-01.mgc.mentorg.com ([147.34.90.201]) by relay1.mentorg.com with esmtp id 1cEJPn-000308-EP from Cesar_Philippidis@mentor.com for gcc-patches@gcc.gnu.org; Tue, 06 Dec 2016 09:20:27 -0800 Received: from [127.0.0.1] (147.34.91.1) by svr-orw-mbx-01.mgc.mentorg.com (147.34.90.201) with Microsoft SMTP Server (TLS) id 15.0.1210.3; Tue, 6 Dec 2016 09:20:24 -0800 From: Cesar Philippidis Subject: [gomp4] fix implicit data clause linker error To: "gcc-patches@gcc.gnu.org" Message-ID: <068721cc-1150-385e-7c34-229d50488ea2@codesourcery.com> Date: Tue, 6 Dec 2016 09:20:24 -0800 User-Agent: Mozilla/5.0 (X11; Linux x86_64; rv:45.0) Gecko/20100101 Thunderbird/45.5.1 MIME-Version: 1.0 X-ClientProxiedBy: svr-orw-mbx-04.mgc.mentorg.com (147.34.90.204) To svr-orw-mbx-01.mgc.mentorg.com (147.34.90.201) The patch I introduced here to fix PR70828 in gomp-4_0-branch has a bug involving the creation of new implicit data clauses for variables inside ACC PARALLEL or KERNELS regions that are present inside an enclosing ACC DATA region. Specifically, as gimplifiy_adjust_omp_clauses_1 introduced new present data clauses for those variables, it neglected to link those clauses with the existing clauses. This is bad in any case, but for data clauses involving static or global variables, this would result in an LTO error because those global variable definitions aren't included inside the offloaded LTO partitions. I've applied this fix to gomp-4_0-branch. Cesar 2016-12-06 Cesar Philippidis gcc/ * gimplify.c (gimplify_adjust_omp_clauses_1): Link ACC new clauses with the old ones. gcc/testsuite/ * c-c++-common/goacc/acc-data-chain.c: New test. diff --git a/gcc/gimplify.c b/gcc/gimplify.c index 9649fae..00b1b23 100644 --- a/gcc/gimplify.c +++ b/gcc/gimplify.c @@ -7992,6 +7992,7 @@ gimplify_adjust_omp_clauses_1 (splay_tree_node n, void *data) gimplify_expr (&OMP_CLAUSE_SIZE (nc), pre_p, NULL, is_gimple_val, fb_rvalue); gimplify_omp_ctxp = ctx; + OMP_CLAUSE_CHAIN (nc) = OMP_CLAUSE_CHAIN (clause); OMP_CLAUSE_CHAIN (clause) = nc; } else if (DECL_SIZE (decl) diff --git a/gcc/testsuite/c-c++-common/goacc/acc-data-chain.c b/gcc/testsuite/c-c++-common/goacc/acc-data-chain.c new file mode 100644 index 0000000..c2f33ff --- /dev/null +++ b/gcc/testsuite/c-c++-common/goacc/acc-data-chain.c @@ -0,0 +1,24 @@ +/* Ensure that the gimplifier does not remove any existing clauses as + it inserts new implicit data clauses. */ + +/* { dg-additional-options "-fdump-tree-gimple" } */ + +#define N 100 +static int a[N], b[N]; + +int main(int argc, char *argv[]) +{ + int i; + +#pragma acc data copyin(a[0:N]) copyout (b[0:N]) + { +#pragma acc parallel loop + for (i = 0; i < N; i++) + b[i] = a[i]; + } + + return 0; +} + +// { dg-final { scan-tree-dump-times "omp target oacc_data map.force_from:b.0. .len: 400.. map.force_to:a.0. .len: 400.." 1 "gimple" } } +// { dg-final { scan-tree-dump-times "omp target oacc_parallel map.force_present:b.0. .len: 400.. map.firstprivate:b .pointer assign, bias: 0.. map.force_present:a.0. .len: 400.. map.firstprivate:a .pointer assign, bias: 0.." 1 "gimple" } }