From patchwork Wed Jul 10 00:52:51 2019 Content-Type: text/plain; charset="utf-8" MIME-Version: 1.0 Content-Transfer-Encoding: 7bit X-Patchwork-Submitter: Julian Brown X-Patchwork-Id: 1130088 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-504787-incoming=patchwork.ozlabs.org@gcc.gnu.org; receiver=) Authentication-Results: ozlabs.org; dmarc=none (p=none dis=none) header.from=codesourcery.com Authentication-Results: ozlabs.org; dkim=pass (1024-bit key; unprotected) header.d=gcc.gnu.org header.i=@gcc.gnu.org header.b="ENbT2E8i"; 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 45k0zR2B51z9sN6 for ; Wed, 10 Jul 2019 10:53:07 +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:date :from:to:subject:message-id:mime-version:content-type; q=dns; s= default; b=HBE4V+3onDXy69ZcEScUtvtRuojidiKIA2DOjLIh+xYwobX6dHN2n /nQ3csT1c88DbrbRNSSGzwJfNf2zwFFHMRY55R7SfvmAAoAdLQ1t/BM1Br3n2YLp SqL4TOXGDWtKZHev6+ut5af+aQJqpK9wR5UFEwRAkO0K77hEPa/Dv4= 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:date :from:to:subject:message-id:mime-version:content-type; s= default; bh=PjJKW+sIjGu+pyTBgDZSwvKMaNs=; b=ENbT2E8iTlS51MxBWhz8 RlqyKjxH2kP0pMuvw941GHCcdXk8EuKS+pOJgYLO38zqJkfpwmAPo7nOAOfwNlG/ wy0WZ1fkYUk9FPStT9D5mK8LnymnLGA2166cyxqdyR2FrRbHH/7cVgi01hj51crS xkMaQB2arprIXCzJjMBTWNg= Received: (qmail 31603 invoked by alias); 10 Jul 2019 00:52:59 -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 31595 invoked by uid 89); 10 Jul 2019 00:52:59 -0000 Authentication-Results: sourceware.org; auth=none X-Spam-SWARE-Status: No, score=-15.6 required=5.0 tests=AWL, BAYES_00, GIT_PATCH_0, GIT_PATCH_1, GIT_PATCH_2, GIT_PATCH_3, KAM_SHORT, RCVD_IN_DNSWL_NONE autolearn=ham version=3.3.1 spammy=sk:thomas, lifted 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, 10 Jul 2019 00:52:58 +0000 Received: from nat-ies.mentorg.com ([192.94.31.2] helo=SVR-IES-MBX-04.mgc.mentorg.com) by relay1.mentorg.com with esmtps (TLSv1.2:ECDHE-RSA-AES256-SHA384:256) id 1hl0qu-0004le-CE from Julian_Brown@mentor.com for gcc-patches@gcc.gnu.org; Tue, 09 Jul 2019 17:52:56 -0700 Received: from squid.athome (137.202.0.90) by SVR-IES-MBX-04.mgc.mentorg.com (139.181.222.4) with Microsoft SMTP Server (TLS) id 15.0.1320.4; Wed, 10 Jul 2019 01:52:52 +0100 Date: Wed, 10 Jul 2019 01:52:51 +0100 From: Julian Brown To: Subject: [og9] OpenACC assumed-size arrays with non-lexical data mappings Message-ID: <20190710015251.118ea4cb@squid.athome> MIME-Version: 1.0 X-IsSubscribed: yes Hi, This patch provides support for implicit mapping of assumed-sized arrays for OpenACC, in cases where those arrays have previously been mapped using non-lexical data mappings (e.g. "#pragma acc enter data"). Previously posted here: https://gcc.gnu.org/ml/gcc-patches/2016-08/msg02090.html and then revised: https://gcc.gnu.org/ml/gcc-patches/2017-07/msg00069.html It's not clear if this is required behaviour for OpenACC, but at least one test program we are using relies on the semantics introduced by this patch. Tested with offloading to nvptx. I will apply to the openacc-gcc-9-branch shortly. Julian ChangeLog 2019-07-10 Cesar Philippidis Thomas Schwinge Julian Brown gcc/ * gimplify.c (gimplify_adjust_omp_clauses_1): Raise error for assumed-size arrays in map clauses for Fortran/OpenMP. * omp-low.c (lower_omp_target): Set the size of assumed-size Fortran arrays to one to allow use of data already mapped on the offload device. gcc/fortran/ * trans-openmp.c (gfc_omp_finish_clause): Change clauses mapping assumed-size arrays to use the GOMP_MAP_FORCE_PRESENT map type. From 2c5a7e445ebadc920730c732279732d2f9b40598 Mon Sep 17 00:00:00 2001 From: Julian Brown Date: Thu, 4 Jul 2019 18:14:41 -0700 Subject: [PATCH 2/3] Assumed-size arrays with non-lexical data mappings gcc/fortran/ * trans-openmp.c (gfc_omp_finish_clause): Change clauses mapping assumed-size arrays to use the GOMP_MAP_FORCE_PRESENT map type. * gimplify.c (gimplify_adjust_omp_clauses_1): Raise error for assumed-size arrays in map clauses for Fortran/OpenMP. * omp-low.c (lower_omp_target): Set the size of assumed-size Fortran arrays to one to allow use of data already mapped on the offload device. --- gcc/fortran/ChangeLog.openacc | 9 +++++++++ gcc/fortran/trans-openmp.c | 22 +++++++++++++--------- gcc/gimplify.c | 14 ++++++++++++++ gcc/omp-low.c | 5 +++++ 4 files changed, 41 insertions(+), 9 deletions(-) diff --git a/gcc/fortran/ChangeLog.openacc b/gcc/fortran/ChangeLog.openacc index c44a5ebdb3b..beba7d94ad2 100644 --- a/gcc/fortran/ChangeLog.openacc +++ b/gcc/fortran/ChangeLog.openacc @@ -1,3 +1,12 @@ +2019-07-10 Julian Brown + + * trans-openmp.c (gfc_omp_finish_clause): Change clauses mapping + assumed-size arrays to use the GOMP_MAP_FORCE_PRESENT map type. + * gimplify.c (gimplify_adjust_omp_clauses_1): Raise error for + assumed-size arrays in map clauses for Fortran/OpenMP. + * omp-low.c (lower_omp_target): Set the size of assumed-size Fortran + arrays to one to allow use of data already mapped on the offload device. + 2019-07-10 Julian Brown * openmp.c (resolve_oacc_data_clauses): Allow polymorphic allocatable diff --git a/gcc/fortran/trans-openmp.c b/gcc/fortran/trans-openmp.c index d5ae0b717df..db009130c85 100644 --- a/gcc/fortran/trans-openmp.c +++ b/gcc/fortran/trans-openmp.c @@ -1137,10 +1137,18 @@ gfc_omp_finish_clause (tree c, gimple_seq *pre_p) tree decl = OMP_CLAUSE_DECL (c); /* Assumed-size arrays can't be mapped implicitly, they have to be mapped - explicitly using array sections. An exception is if the array is - mapped explicitly in an enclosing data construct for OpenACC, in which - case we see GOMP_MAP_FORCE_PRESENT here and do not need to raise an - error. */ + explicitly using array sections. For OpenACC this restriction is lifted + if the array has already been mapped: + + - Using a lexically-enclosing data region: in that case we see the + GOMP_MAP_FORCE_PRESENT mapping kind here. + + - Using a non-lexical data mapping ("acc enter data"). + + In the latter case we change the mapping type to GOMP_MAP_FORCE_PRESENT. + This raises an error for OpenMP in our the caller + (gimplify.c:gimplify_adjust_omp_clauses_1). OpenACC will raise a runtime + error if the assumed-size array is not mapped. */ if (OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_FORCE_PRESENT && TREE_CODE (decl) == PARM_DECL && GFC_ARRAY_TYPE_P (TREE_TYPE (decl)) @@ -1148,11 +1156,7 @@ gfc_omp_finish_clause (tree c, gimple_seq *pre_p) && GFC_TYPE_ARRAY_UBOUND (TREE_TYPE (decl), GFC_TYPE_ARRAY_RANK (TREE_TYPE (decl)) - 1) == NULL) - { - error_at (OMP_CLAUSE_LOCATION (c), - "implicit mapping of assumed size array %qD", decl); - return; - } + OMP_CLAUSE_SET_MAP_KIND (c, GOMP_MAP_FORCE_PRESENT); tree c2 = NULL_TREE, c3 = NULL_TREE, c4 = NULL_TREE; if (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_FORCE_DEVICEPTR) diff --git a/gcc/gimplify.c b/gcc/gimplify.c index 60e04ff8353..58142c9eb90 100644 --- a/gcc/gimplify.c +++ b/gcc/gimplify.c @@ -10088,7 +10088,21 @@ gimplify_adjust_omp_clauses_1 (splay_tree_node n, void *data) *list_p = clause; struct gimplify_omp_ctx *ctx = gimplify_omp_ctxp; gimplify_omp_ctxp = ctx->outer_context; + gomp_map_kind kind = (code == OMP_CLAUSE_MAP) ? OMP_CLAUSE_MAP_KIND (clause) + : (gomp_map_kind) GOMP_MAP_LAST; lang_hooks.decls.omp_finish_clause (clause, pre_p); + /* Allow OpenACC to have implicit assumed-size arrays via FORCE_PRESENT, + which should work as long as the array has previously been mapped + explicitly on the target (e.g. by "enter data"). Raise an error for + OpenMP. */ + if (lang_GNU_Fortran () + && code == OMP_CLAUSE_MAP + && (ctx->region_type & ORT_ACC) == 0 + && kind == GOMP_MAP_TOFROM + && OMP_CLAUSE_MAP_KIND (clause) == GOMP_MAP_FORCE_PRESENT) + error_at (OMP_CLAUSE_LOCATION (clause), + "implicit mapping of assumed size array %qD", + OMP_CLAUSE_DECL (clause)); if (gimplify_omp_ctxp) for (; clause != chain; clause = OMP_CLAUSE_CHAIN (clause)) if (OMP_CLAUSE_CODE (clause) == OMP_CLAUSE_MAP diff --git a/gcc/omp-low.c b/gcc/omp-low.c index 249cc1aac53..97c00217d9f 100644 --- a/gcc/omp-low.c +++ b/gcc/omp-low.c @@ -10453,6 +10453,11 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx) s = OMP_CLAUSE_SIZE (c); if (s == NULL_TREE) s = TYPE_SIZE_UNIT (TREE_TYPE (ovar)); + /* Fortran assumed-size arrays have zero size because the type is + incomplete. Set the size to one to allow the runtime to remap + any existing data that is already present on the accelerator. */ + if (s == NULL_TREE && is_gimple_omp_oacc (ctx->stmt)) + s = integer_one_node; s = fold_convert (size_type_node, s); decl_args = append_decl_arg (ovar, decl_args, ctx); purpose = size_int (map_idx++); -- 2.22.0