From patchwork Tue Aug 7 22:09:38 2018 Content-Type: text/plain; charset="utf-8" MIME-Version: 1.0 Content-Transfer-Encoding: 7bit X-Patchwork-Submitter: Cesar Philippidis X-Patchwork-Id: 954699 Return-Path: X-Original-To: incoming@patchwork.ozlabs.org Delivered-To: patchwork-incoming@bilbo.ozlabs.org Authentication-Results: ozlabs.org; spf=pass (mailfrom) smtp.mailfrom=gcc.gnu.org (client-ip=209.132.180.131; helo=sourceware.org; envelope-from=gcc-patches-return-483359-incoming=patchwork.ozlabs.org@gcc.gnu.org; receiver=) Authentication-Results: ozlabs.org; dmarc=none (p=none dis=none) header.from=mentor.com Authentication-Results: ozlabs.org; dkim=pass (1024-bit key; unprotected) header.d=gcc.gnu.org header.i=@gcc.gnu.org header.b="Ns+KB4wA"; dkim-atps=neutral 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 41lTGM1f5Pz9s4c for ; Wed, 8 Aug 2018 08:10:01 +1000 (AEST) DomainKey-Signature: a=rsa-sha1; c=nofws; d=gcc.gnu.org; h=list-id :list-unsubscribe:list-archive:list-post:list-help:sender:to :from:subject:message-id:date:mime-version:content-type; q=dns; s=default; b=qOTWyIq+TaXSl9IHHWkwrTIVSeZey9fLP0L5IyVZcFh36bDUYv gjvHQS4tI87XtR6QxPhM12IRFdKaZDqYN6CtMmAtrJgIRaferj9lexFFZC/qPmV5 mfhAY1tDVAc80aZ+weHw5rXRQhdXURw+XmI8lItMzCwHIpIe2lUzEEHPo= 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:to :from:subject:message-id:date:mime-version:content-type; s= default; bh=dQaznL/ojeDCIWbUKvC85iRBq/g=; b=Ns+KB4wAPlbumEIyYk8J saAbMgq7UA0yyrCJ4Dse+4TCiWLStIX2rfYRRjOrbLz7ARwjqrxaGooZqy4dXg5b lmTo3jIkUdV3nCR1cBJANbnHgqniFaNlA0s+mJ8an7m/7WHInBKcYTog0ycUC+Ny l/9afTuqTo3+cAPsx6oCz9M= Received: (qmail 58139 invoked by alias); 7 Aug 2018 22:09: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 58129 invoked by uid 89); 7 Aug 2018 22:09:54 -0000 Authentication-Results: sourceware.org; auth=none X-Spam-SWARE-Status: No, score=-24.9 required=5.0 tests=AWL, BAYES_00, GIT_PATCH_0, GIT_PATCH_1, GIT_PATCH_2, GIT_PATCH_3, KAM_SHORT, SPF_PASS, URIBL_RED autolearn=ham version=3.3.2 spammy=HX-detected-operating-system:Windows, 84, cleanly, sk:finaliz X-HELO: fencepost.gnu.org Received: from fencepost.gnu.org (HELO fencepost.gnu.org) (208.118.235.10) by sourceware.org (qpsmtpd/0.93/v0.84-503-g423c35a) with ESMTP; Tue, 07 Aug 2018 22:09:52 +0000 Received: from eggs.gnu.org ([2001:4830:134:3::10]:51049) by fencepost.gnu.org with esmtps (TLS1.0:RSA_AES_256_CBC_SHA1:256) (Exim 4.82) (envelope-from ) id 1fnAAo-0004YL-LF for gcc-patches@gnu.org; Tue, 07 Aug 2018 18:09:50 -0400 Received: from Debian-exim by eggs.gnu.org with spam-scanned (Exim 4.71) (envelope-from ) id 1fnAAl-0001Ix-AV for gcc-patches@gnu.org; Tue, 07 Aug 2018 18:09:50 -0400 Received: from relay1.mentorg.com ([192.94.38.131]:54431) by eggs.gnu.org with esmtps (TLS1.0:DHE_RSA_AES_256_CBC_SHA1:32) (Exim 4.71) (envelope-from ) id 1fnAAl-0001Dj-1p for gcc-patches@gnu.org; Tue, 07 Aug 2018 18:09:47 -0400 Received: from svr-orw-mbx-01.mgc.mentorg.com ([147.34.90.201]) by relay1.mentorg.com with esmtps (TLSv1.2:ECDHE-RSA-AES256-SHA384:256) id 1fnAAg-0006Fu-It from Cesar_Philippidis@mentor.com for gcc-patches@gnu.org; Tue, 07 Aug 2018 15:09:42 -0700 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.1320.4; Tue, 7 Aug 2018 15:09:40 -0700 To: "gcc-patches@gnu.org" From: Cesar Philippidis Subject: [PATCH][OpenACC] Update deviceptr handling during gimplification Message-ID: Date: Tue, 7 Aug 2018 15:09:38 -0700 User-Agent: Mozilla/5.0 (X11; Linux x86_64; rv:52.0) Gecko/20100101 Thunderbird/52.9.1 MIME-Version: 1.0 X-detected-operating-system: by eggs.gnu.org: Windows NT kernel [generic] [fuzzy] X-Received-From: 192.94.38.131 I had previously posted this patch as part of a monster deviceptr patch here . This patch breaks out the generic gimplifier changes. Essentially, with this patch, the gimplifier will now transfer deviceptr data clauses using GOMP_MAP_FORCE_DEVICEPTR. Is this patch OK for trunk? It bootstrapped / regression tested cleanly for x86_64 with nvptx offloading. Thanks, Cesar From b5cf37b795ce78c78f3f434ac6999f7094bd86aa Mon Sep 17 00:00:00 2001 From: Cesar Philippidis Date: Mon, 7 May 2018 08:23:48 -0700 Subject: [PATCH] [OpenACC] Update deviceptr handling 2018-XX-YY Cesar Philippidis gcc/fortran/ * trans-openmp.c (gfc_omp_finish_clause): Don't create pointer data mappings for deviceptr clauses. (gfc_trans_omp_clauses): Likewise. gcc/ * gimplify.c (enum gimplify_omp_var_data): Add GOVD_DEVICETPR. (omp_notice_variable): Add GOVD_DEVICEPTR attribute when appropriate. (gimplify_scan_omp_clauses): Likewise. (gimplify_adjust_omp_clauses_1): Set GOMP_MAP_FORCE_DEVICEPTR for implicit deviceptr mappings. gcc/testsuite/ * c-c++-common/goacc/deviceptr-4.c: Update expected data mapping. (cherry picked from openacc-gcc-7-branch commit d3de16b461545aac1925f0d7c2851c8c49a07d06 and commit f0514fe1899666bb5b8ee52601f5d4263d4c4646) --- gcc/fortran/trans-openmp.c | 9 +++++++++ gcc/gimplify.c | 12 +++++++++++- gcc/testsuite/c-c++-common/goacc/deviceptr-4.c | 2 +- 3 files changed, 21 insertions(+), 2 deletions(-) diff --git a/gcc/fortran/trans-openmp.c b/gcc/fortran/trans-openmp.c index f038f4c..ca31c88 100644 --- a/gcc/fortran/trans-openmp.c +++ b/gcc/fortran/trans-openmp.c @@ -1060,6 +1060,8 @@ gfc_omp_finish_clause (tree c, gimple_seq *pre_p) } tree c2 = NULL_TREE, c3 = NULL_TREE, c4 = NULL_TREE; + if (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_FORCE_DEVICEPTR) + return; if (POINTER_TYPE_P (TREE_TYPE (decl))) { if (!gfc_omp_privatize_by_reference (decl) @@ -2111,6 +2113,12 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses, if (n->expr == NULL || n->expr->ref->u.ar.type == AR_FULL) { if (POINTER_TYPE_P (TREE_TYPE (decl)) + && n->u.map_op == OMP_MAP_FORCE_DEVICEPTR) + { + OMP_CLAUSE_DECL (node) = decl; + goto finalize_map_clause; + } + else if (POINTER_TYPE_P (TREE_TYPE (decl)) && (gfc_omp_privatize_by_reference (decl) || GFC_DECL_GET_SCALAR_POINTER (decl) || GFC_DECL_GET_SCALAR_ALLOCATABLE (decl) @@ -2282,6 +2290,7 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses, ptr2 = fold_convert (sizetype, ptr2); OMP_CLAUSE_SIZE (node3) = fold_build2 (MINUS_EXPR, sizetype, ptr, ptr2); + finalize_map_clause:; } switch (n->u.map_op) { diff --git a/gcc/gimplify.c b/gcc/gimplify.c index 4a109ae..bcf862f 100644 --- a/gcc/gimplify.c +++ b/gcc/gimplify.c @@ -105,6 +105,9 @@ enum gimplify_omp_var_data /* Flag for GOVD_MAP: must be present already. */ GOVD_MAP_FORCE_PRESENT = 524288, + /* Flag for OpenACC deviceptrs. */ + GOVD_DEVICEPTR = (1<<21), + GOVD_DATA_SHARE_CLASS = (GOVD_SHARED | GOVD_PRIVATE | GOVD_FIRSTPRIVATE | GOVD_LASTPRIVATE | GOVD_REDUCTION | GOVD_LINEAR | GOVD_LOCAL) @@ -7232,6 +7235,7 @@ omp_notice_variable (struct gimplify_omp_ctx *ctx, tree decl, bool in_code) error ("variable %qE declared in enclosing " "% region", DECL_NAME (decl)); nflags |= GOVD_MAP; + nflags |= (n2->value & GOVD_DEVICEPTR); if (octx->region_type == ORT_ACC_DATA && (n2->value & GOVD_MAP_0LEN_ARRAY)) nflags |= GOVD_MAP_0LEN_ARRAY; @@ -8213,6 +8217,8 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p, if (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ALWAYS_TO || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ALWAYS_TOFROM) flags |= GOVD_MAP_ALWAYS_TO; + else if (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_FORCE_DEVICEPTR) + flags |= GOVD_DEVICEPTR; goto do_add; case OMP_CLAUSE_DEPEND: @@ -8828,7 +8834,8 @@ gimplify_adjust_omp_clauses_1 (splay_tree_node n, void *data) /* Not all combinations of these GOVD_MAP flags are actually valid. */ switch (flags & (GOVD_MAP_TO_ONLY | GOVD_MAP_FORCE - | GOVD_MAP_FORCE_PRESENT)) + | GOVD_MAP_FORCE_PRESENT + | GOVD_DEVICEPTR)) { case 0: kind = GOMP_MAP_TOFROM; @@ -8845,6 +8852,9 @@ gimplify_adjust_omp_clauses_1 (splay_tree_node n, void *data) case GOVD_MAP_FORCE_PRESENT: kind = GOMP_MAP_FORCE_PRESENT; break; + case GOVD_DEVICEPTR: + kind = GOMP_MAP_FORCE_DEVICEPTR; + break; default: gcc_unreachable (); } diff --git a/gcc/testsuite/c-c++-common/goacc/deviceptr-4.c b/gcc/testsuite/c-c++-common/goacc/deviceptr-4.c index db1b916..79a5162 100644 --- a/gcc/testsuite/c-c++-common/goacc/deviceptr-4.c +++ b/gcc/testsuite/c-c++-common/goacc/deviceptr-4.c @@ -8,4 +8,4 @@ subr (int *a) a[0] += 1.0; } -/* { dg-final { scan-tree-dump-times "#pragma omp target oacc_parallel.*map\\(tofrom:a" 1 "gimple" } } */ +/* { dg-final { scan-tree-dump-times "#pragma omp target oacc_parallel.*map\\(force_deviceptr:a" 1 "gimple" } } */ -- 2.7.4