From mboxrd@z Thu Jan 1 00:00:00 1970 Return-Path: Received: (qmail 55293 invoked by alias); 22 Mar 2017 13:40:12 -0000 Mailing-List: contact gcc-patches-help@gcc.gnu.org; run by ezmlm Precedence: bulk List-Id: List-Archive: List-Post: List-Help: Sender: gcc-patches-owner@gcc.gnu.org Received: (qmail 55246 invoked by uid 89); 22 Mar 2017 13:40:10 -0000 Authentication-Results: sourceware.org; auth=none X-Virus-Found: No X-Spam-SWARE-Status: No, score=-24.5 required=5.0 tests=AWL,BAYES_00,GIT_PATCH_0,GIT_PATCH_1,GIT_PATCH_2,GIT_PATCH_3,RCVD_IN_DNSWL_NONE,SPF_PASS,URIBL_RED autolearn=ham version=3.3.2 spammy=Cesar, cesar, H*Ad:U*cesar 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; Wed, 22 Mar 2017 13:40:08 +0000 Received: from svr-orw-mbx-04.mgc.mentorg.com ([147.34.90.204]) by relay1.mentorg.com with esmtp id 1cqgUg-0000c0-CH from Cesar_Philippidis@mentor.com ; Wed, 22 Mar 2017 06:40:06 -0700 Received: from [127.0.0.1] (147.34.91.1) by SVR-ORW-MBX-04.mgc.mentorg.com (147.34.90.204) with Microsoft SMTP Server (TLS) id 15.0.1210.3; Wed, 22 Mar 2017 06:40:03 -0700 To: "gcc-patches@gcc.gnu.org" , Jakub Jelinek From: Cesar Philippidis Subject: [PATCH] Fix PR80029 Message-ID: <278ce4b7-cc02-453b-10fc-8094e9078d9a@codesourcery.com> Date: Wed, 22 Mar 2017 13:40:00 -0000 User-Agent: Mozilla/5.0 (X11; Linux x86_64; rv:45.0) Gecko/20100101 Thunderbird/45.7.0 MIME-Version: 1.0 Content-Type: multipart/mixed; boundary="------------8468025A2D097CD85CE0D7B5" X-ClientProxiedBy: SVR-ORW-MBX-09.mgc.mentorg.com (147.34.90.209) To SVR-ORW-MBX-04.mgc.mentorg.com (147.34.90.204) X-SW-Source: 2017-03/txt/msg01161.txt.bz2 --------------8468025A2D097CD85CE0D7B5 Content-Type: text/plain; charset="utf-8" Content-Transfer-Encoding: 7bit Content-length: 309 In addition to resolving the memory leak involving OpenACC delare clauses, this patch also corrects an ICE involving VLA variables as data clause arguments used in acc declare constructs. More details can be found here . Is this OK for trunk? Cesar --------------8468025A2D097CD85CE0D7B5 Content-Type: text/x-patch; name="pr80029.diff" Content-Transfer-Encoding: 7bit Content-Disposition: attachment; filename="pr80029.diff" Content-length: 3846 2017-03-22 Cesar Philippidis PR c++/80029 gcc/ * gimplify.c (is_oacc_declared): New function. (oacc_default_clause): Use it to set default flags for acc declared variables inside parallel regions. (gimplify_scan_omp_clauses): Strip firstprivate pointers for acc declared variables. (gimplify_oacc_declare): Gimplify the declare clauses. Add the declare attribute to any decl as necessary. libgomp/ * testsuite/libgomp.oacc-c-c++-common/declare-vla.c: New test. diff --git a/gcc/gimplify.c b/gcc/gimplify.c index fbf136f..26d35d1 100644 --- a/gcc/gimplify.c +++ b/gcc/gimplify.c @@ -6787,6 +6787,16 @@ device_resident_p (tree decl) return false; } +/* Return true if DECL has an ACC DECLARE attribute. */ + +static bool +is_oacc_declared (tree decl) +{ + tree t = TREE_CODE (decl) == MEM_REF ? TREE_OPERAND (decl, 0) : decl; + tree declared = lookup_attribute ("oacc declare target", DECL_ATTRIBUTES (t)); + return declared != NULL_TREE; +} + /* Determine outer default flags for DECL mentioned in an OMP region but not declared in an enclosing clause. @@ -6887,6 +6897,7 @@ oacc_default_clause (struct gimplify_omp_ctx *ctx, tree decl, unsigned flags) { const char *rkind; bool on_device = false; + bool declared = is_oacc_declared (decl); tree type = TREE_TYPE (decl); if (lang_hooks.decls.omp_privatize_by_reference (decl)) @@ -6917,7 +6928,7 @@ oacc_default_clause (struct gimplify_omp_ctx *ctx, tree decl, unsigned flags) case ORT_ACC_PARALLEL: { - if (on_device || AGGREGATE_TYPE_P (type)) + if (on_device || AGGREGATE_TYPE_P (type) || declared) /* Aggregates default to 'present_or_copy'. */ flags |= GOVD_MAP; else @@ -7346,6 +7357,7 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p, case OMP_TARGET_DATA: case OMP_TARGET_ENTER_DATA: case OMP_TARGET_EXIT_DATA: + case OACC_DECLARE: case OACC_HOST_DATA: ctx->target_firstprivatize_array_bases = true; default: @@ -9231,18 +9243,26 @@ gimplify_oacc_declare (tree *expr_p, gimple_seq *pre_p) { tree expr = *expr_p; gomp_target *stmt; - tree clauses, t; + tree clauses, t, decl; clauses = OACC_DECLARE_CLAUSES (expr); gimplify_scan_omp_clauses (&clauses, pre_p, ORT_TARGET_DATA, OACC_DECLARE); + gimplify_adjust_omp_clauses (pre_p, NULL, &clauses, OACC_DECLARE); for (t = clauses; t; t = OMP_CLAUSE_CHAIN (t)) { - tree decl = OMP_CLAUSE_DECL (t); + decl = OMP_CLAUSE_DECL (t); if (TREE_CODE (decl) == MEM_REF) - continue; + decl = TREE_OPERAND (decl, 0); + + if (VAR_P (decl) && !is_oacc_declared (decl)) + { + tree attr = get_identifier ("oacc declare target"); + DECL_ATTRIBUTES (decl) = tree_cons (attr, NULL_TREE, + DECL_ATTRIBUTES (decl)); + } if (VAR_P (decl) && !is_global_var (decl) @@ -9258,7 +9278,8 @@ gimplify_oacc_declare (tree *expr_p, gimple_seq *pre_p) } } - omp_add_variable (gimplify_omp_ctxp, decl, GOVD_SEEN); + if (gimplify_omp_ctxp) + omp_add_variable (gimplify_omp_ctxp, decl, GOVD_SEEN); } stmt = gimple_build_omp_target (NULL, GF_OMP_TARGET_KIND_OACC_DECLARE, diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/declare-vla.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/declare-vla.c new file mode 100644 index 0000000..3ea148e --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/declare-vla.c @@ -0,0 +1,25 @@ +/* Verify that acc declare accept VLA variables. */ + +#include + +int +main () +{ + int N = 1000; + int i, A[N]; +#pragma acc declare copy(A) + + for (i = 0; i < N; i++) + A[i] = -i; + +#pragma acc kernels + for (i = 0; i < N; i++) + A[i] = i; + +#pragma acc update host(A) + + for (i = 0; i < N; i++) + assert (A[i] == i); + + return 0; +} --------------8468025A2D097CD85CE0D7B5--