From patchwork Thu Sep 22 22:38:40 2016 Content-Type: text/plain; charset="utf-8" MIME-Version: 1.0 Content-Transfer-Encoding: 7bit X-Patchwork-Submitter: Cesar Philippidis X-Patchwork-Id: 673744 Return-Path: X-Original-To: incoming@patchwork.ozlabs.org Delivered-To: patchwork-incoming@bilbo.ozlabs.org Received: from sourceware.org (server1.sourceware.org [209.132.180.131]) (using TLSv1.2 with cipher ECDHE-RSA-AES256-GCM-SHA384 (256/256 bits)) (No client certificate requested) by ozlabs.org (Postfix) with ESMTPS id 3sgBHY30Tlz9t17 for ; Fri, 23 Sep 2016 08:39:03 +1000 (AEST) Authentication-Results: ozlabs.org; dkim=pass (1024-bit key; unprotected) header.d=gcc.gnu.org header.i=@gcc.gnu.org header.b=o5n1fzlK; dkim-atps=neutral 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=n+nX1xbIk6ZgWb/un6STCDUkfMapFf5Uerr3JvTaKz3zOWVl8D3ga ye/YP+n6q/8AdQysy0SyNCuNe6C4xqpAkaSfFdofswgYwrQtC+nxwdMUZIBfn0Yc uYbpEIqPWizKyogcX/wks92G658nM+qR0HtYTsGNmXGnn9dIcEmUxY= 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=cMAWa/+o8bwiQBu4/qqidtPN/Lk=; b=o5n1fzlK5ue4ISgpebg4 SoKldyFXkzl5nKpVMfcSqhqSwqlzebGWwyn3SLPQ8KfXrTjLy1Zcvsw4Hoo8L1MJ 7lDMcber5rWeJ3c5s6quMLqY3hvrEvIPATOTcW4X7iViIwHi1AloWrVkQ8vXUHBG bNkynSuhie2NZ3maE7/vEJI= Received: (qmail 89298 invoked by alias); 22 Sep 2016 22:38:54 -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 89283 invoked by uid 89); 22 Sep 2016 22:38:54 -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=cesar, 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; Thu, 22 Sep 2016 22:38:44 +0000 Received: from svr-orw-fem-03.mgc.mentorg.com ([147.34.97.39]) by relay1.mentorg.com with esmtp id 1bnCdc-0001zm-Vh from Cesar_Philippidis@mentor.com ; Thu, 22 Sep 2016 15:38:40 -0700 Received: from [127.0.0.1] (147.34.91.1) by svr-orw-fem-03.mgc.mentorg.com (147.34.97.39) with Microsoft SMTP Server id 14.3.224.2; Thu, 22 Sep 2016 15:38:40 -0700 From: Cesar Philippidis Subject: [openacc] fix an ICE with acc declared VLAs To: "gcc-patches@gcc.gnu.org" , Jakub Jelinek Message-ID: Date: Thu, 22 Sep 2016 15:38:40 -0700 User-Agent: Mozilla/5.0 (X11; Linux x86_64; rv:45.0) Gecko/20100101 Thunderbird/45.2.0 MIME-Version: 1.0 There's a bug with ACC DECLARE clauses involving VLA variables that causes lower_omp_target to thow an ICE on VLA decls. The problem occurred because those clauses are never gimplified. This patch resolves that ICE by teaching gimplify_oacc_declare how to gimplify them. It turns out that gimplify_adjust_omp_clauses removes the gimplify_omp_ctx created by gimplify_scan_omp_clauses. To workaround that problem, I just taught gimplify_oacc_declare to ensure that each gimplified decl has an "oacc declare target" attribute. This seems like a reasonable solution because acc declare is an executable directive, and hence really doesn't have a normal scope like ACC DATA. The declare code does install destructors to unmap those clauses (except for those vars which were ACC DECLARE CREATE'd -- but that's a problem for patch), but that doesn't depend on the gimplify omp ctx. The only other major thing that had to be updated was oacc_default_clause, because omp_notice_variable no longer has a gimple omp ctx to inspect for existing variables. But that's where the "oacc declare target" attribute comes into play. Is this patch OK for trunk? Cesar 2016-09-22 Cesar Philippidis 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 e378ed0..19095ff 100644 --- a/gcc/gimplify.c +++ b/gcc/gimplify.c @@ -6114,6 +6114,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. @@ -6214,6 +6224,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)) @@ -6244,7 +6255,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 @@ -6714,6 +6725,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: @@ -8589,20 +8601,28 @@ 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 (TREE_CODE (decl) == VAR_DECL + 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) && DECL_CONTEXT (decl) == current_function_decl) { @@ -8616,7 +8636,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; +}