From patchwork Fri Dec 13 14:13:53 2019 Content-Type: text/plain; charset="utf-8" MIME-Version: 1.0 Content-Transfer-Encoding: 8bit X-Patchwork-Submitter: Thomas Schwinge X-Patchwork-Id: 1209158 Return-Path: X-Original-To: incoming@patchwork.ozlabs.org Delivered-To: patchwork-incoming@bilbo.ozlabs.org Authentication-Results: ozlabs.org; spf=pass (sender SPF authorized) smtp.mailfrom=gcc.gnu.org (client-ip=209.132.180.131; helo=sourceware.org; envelope-from=gcc-patches-return-515888-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="cNUQN6Hc"; 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 47ZCMz3ZTzz9s4Y for ; Sat, 14 Dec 2019 01:14:23 +1100 (AEDT) DomainKey-Signature: a=rsa-sha1; c=nofws; d=gcc.gnu.org; h=list-id :list-unsubscribe:list-archive:list-post:list-help:sender:from :to:cc:subject:in-reply-to:references:date:message-id :mime-version:content-type; q=dns; s=default; b=pWR1xIasyvGWFBmi LL0AmTXf6YLly87E5JC/dBuTh87T2iVRyyIcEP8b53HoTq7t9o4LXG6G4rv1HmUH mUgusaXCNOQ/vNIhj5c3kxT8uzFXlGOK1G7lk1DNr68MRUlDIi5n1qDrhUaXxoms 5p4+82SlyTSUG3sV9m6TEsR870g= 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 :to:cc:subject:in-reply-to:references:date:message-id :mime-version:content-type; s=default; bh=7IvV2rIXyRD0B6Q758Rm7v er4ew=; b=cNUQN6HcVgfJ7XDW6/5kA4al+xz3bZhvP2zGDfjnpgCwMe7H3Vm+fu rwmE22c8pLU217snqWew/QlEDoE54UiORqJNMR/7P78Na05WJdTXt7vXIJvgWKQR R0mMYg+0c7MlJxAfyH2tI0FPW4iq3F29Xa8z82Zm7MEA3u0aJw4B0= Received: (qmail 76777 invoked by alias); 13 Dec 2019 14:14:15 -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 75092 invoked by uid 89); 13 Dec 2019 14:14:15 -0000 Authentication-Results: sourceware.org; auth=none X-Spam-SWARE-Status: No, score=-19.1 required=5.0 tests=AWL, BAYES_00, GIT_PATCH_0, GIT_PATCH_1, GIT_PATCH_2, GIT_PATCH_3, KAM_SHORT, SPF_PASS autolearn=ham version=3.3.1 spammy=respond, acc_is_present, 385, 10, ACC X-HELO: esa3.mentor.iphmx.com Received: from esa3.mentor.iphmx.com (HELO esa3.mentor.iphmx.com) (68.232.137.180) by sourceware.org (qpsmtpd/0.93/v0.84-503-g423c35a) with ESMTP; Fri, 13 Dec 2019 14:14:12 +0000 IronPort-SDR: fx+pEfWsQjbwLBbAeB7w9SjDKs22kzamBO0TQIWogdWK1BbroFB3deq4f88EMa7gPQpce7oCTp SyUUbiKVlyYJJzwTYdZmsNVIszyUvJuBPl6+sW2ae51WAbTjWar75WSAFEd36w5HP5SZmP+SQ2 Tq4h6CI7ByektO8sAtMArNpbSXRoeH4a+QbHf71ckRj96Ln017vmGlGHiXxdOIJiMdA4P6aL9V IXdmF/0Y75yfOlMkPfFInP0Mk9RWXIG6K4ZQJKLswHPB/tCaU/JqbjodphhP76lHhyzFxbHpMP CZ4= Received: from orw-gwy-02-in.mentorg.com ([192.94.38.167]) by esa3.mentor.iphmx.com with ESMTP; 13 Dec 2019 06:14:10 -0800 IronPort-SDR: nKMBPePUSjoUoj/SsBhbspT/yRzF6/uCCxYpFb3ovT2JjgE4VzAo9LUXYOBbB6Hz0J2+NXFFoW 2oqWel9Vp9EnM3ZoB2lasJ1t0gJJsfEh0aa6Sn7sQCclUVoLDNH90zssFVuOotAHUc3Onyjxx/ LgVgS6N3sqJwaB03H207MDfOnQNUOgtc5Yi0H8k0mFXaIyiCMoajj/uuoVwIjMa1xeXMqWWNTH MZPZGrT1EgZ93AmU3ewF7wAGaFl2USrIpJith8NPRIE147Tb67DABXlVPzbqdQrbzj0Hmd/i7z xhE= From: Thomas Schwinge To: , Julian Brown , Tobias Burnus CC: Jakub Jelinek , Subject: [OpenACC] Elaborate/simplify 'exit data' 'finalize' handling (was: [OpenACC] Update OpenACC data clause semantics to the 2.5 behavior) In-Reply-To: <7fa7637f-e7f5-d43d-13f1-706c77e8e957@codesourcery.com> References: <7fa7637f-e7f5-d43d-13f1-706c77e8e957@codesourcery.com> User-Agent: Notmuch/0.29.1+93~g67ed7df (https://notmuchmail.org) Emacs/26.1 (x86_64-pc-linux-gnu) Date: Fri, 13 Dec 2019 15:13:53 +0100 Message-ID: <87y2vgz526.fsf@euler.schwinge.homeip.net> MIME-Version: 1.0 Hi! Julian, Tobias, regarding the following OpenACC 'exit data' 'finalize' handling: On 2018-05-25T13:01:58-0700, Cesar Philippidis wrote: > --- a/gcc/gimplify.c > +++ b/gcc/gimplify.c > @@ -10859,6 +10849,53 @@ gimplify_omp_target_update (tree *expr_p, gimple_seq *pre_p) > + else if (TREE_CODE (expr) == OACC_EXIT_DATA > + && omp_find_clause (OMP_STANDALONE_CLAUSES (expr), > + OMP_CLAUSE_FINALIZE)) > + { > + /* Use GOMP_MAP_DELETE/GOMP_MAP_FORCE_FROM to denote that "finalize" > + semantics apply to all mappings of this OpenACC directive. */ > + bool finalize_marked = false; > + for (tree c = OMP_STANDALONE_CLAUSES (expr); c; c = OMP_CLAUSE_CHAIN (c)) > + if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP) > + switch (OMP_CLAUSE_MAP_KIND (c)) > + { > + case GOMP_MAP_FROM: > + OMP_CLAUSE_SET_MAP_KIND (c, GOMP_MAP_FORCE_FROM); > + finalize_marked = true; > + break; > + case GOMP_MAP_RELEASE: > + OMP_CLAUSE_SET_MAP_KIND (c, GOMP_MAP_DELETE); > + finalize_marked = true; > + break; > + default: > + /* Check consistency: libgomp relies on the very first data > + mapping clause being marked, so make sure we did that before > + any other mapping clauses. */ > + gcc_assert (finalize_marked); > + break; > + } > + } > --- a/libgomp/oacc-parallel.c > +++ b/libgomp/oacc-parallel.c > @@ -286,6 +360,17 @@ GOACC_enter_exit_data (int device, size_t mapnum, > + /* Determine whether "finalize" semantics apply to all mappings of this > + OpenACC directive. */ > + bool finalize = false; > + if (mapnum > 0) > + { > + unsigned char kind = kinds[0] & 0xff; > + if (kind == GOMP_MAP_DELETE > + || kind == GOMP_MAP_FORCE_FROM) > + finalize = true; > + } > + > @@ -360,22 +458,28 @@ GOACC_enter_exit_data (int device, size_t mapnum, > switch (kind) > { > - case GOMP_MAP_POINTER: > - gomp_acc_remove_pointer (hostaddrs[i], (kinds[i] & 0xff) > - == GOMP_MAP_FORCE_FROM, > - async, 1); > - break; > + case GOMP_MAP_RELEASE: > case GOMP_MAP_DELETE: > - acc_delete (hostaddrs[i], sizes[i]); > + if (acc_is_present (hostaddrs[i], sizes[i])) > + { > + if (finalize) > + acc_delete_finalize (hostaddrs[i], sizes[i]); > + else > + acc_delete (hostaddrs[i], sizes[i]); > + } > break; > + case GOMP_MAP_FROM: > case GOMP_MAP_FORCE_FROM: > - acc_copyout (hostaddrs[i], sizes[i]); > + if (finalize) > + acc_copyout_finalize (hostaddrs[i], sizes[i]); > + else > + acc_copyout (hostaddrs[i], sizes[i]); > break; > default: > gomp_fatal (">>>> GOACC_enter_exit_data UNHANDLED kind 0x%.2x", > @@ -385,10 +489,12 @@ GOACC_enter_exit_data (int device, size_t mapnum, > } > else > { > - gomp_acc_remove_pointer (hostaddrs[i], (kinds[i] & 0xff) > - == GOMP_MAP_FORCE_FROM, async, 3); > +[...] > + gomp_acc_remove_pointer (hostaddrs[i], sizes[i], copyfrom, async, > + finalize, pointer); ... does the attached patch "[OpenACC] Elaborate/simplify 'exit data' 'finalize' handling" (with "No functional changes") match your understanding of what's going on? If approving this patch, please respond with "Reviewed-by: NAME " so that your effort will be recorded in the commit log, see . (It will be a separate discussion to change the 'GOMP_MAP_POINTER', 'GOMP_MAP_TO_PSET' stuff later on -- thinking about the changes from Julian's big "OpenACC reference count overhaul" as well as "OpenACC/OpenMP 'target' 'exit data'/'update' optimizations". That patch here is just meant to document what's going at present, and simplify things as a preparation for other changes.) Grüße Thomas From 283577c63b374c3e368e3c0b68b90e19085f193c Mon Sep 17 00:00:00 2001 From: Thomas Schwinge Date: Fri, 13 Dec 2019 13:56:51 +0100 Subject: [PATCH] [OpenACC] Elaborate/simplify 'exit data' 'finalize' handling No functional changes. --- gcc/gimplify.c | 23 +++++++++++-------- gcc/testsuite/c-c++-common/goacc/finalize-1.c | 11 ++++++++- gcc/testsuite/gfortran.dg/goacc/finalize-1.f | 10 ++++++++ libgomp/oacc-mem.c | 14 +++-------- 4 files changed, 36 insertions(+), 22 deletions(-) diff --git a/gcc/gimplify.c b/gcc/gimplify.c index 9073680cb31..60a80cb8098 100644 --- a/gcc/gimplify.c +++ b/gcc/gimplify.c @@ -12738,27 +12738,30 @@ gimplify_omp_target_update (tree *expr_p, gimple_seq *pre_p) && omp_find_clause (OMP_STANDALONE_CLAUSES (expr), OMP_CLAUSE_FINALIZE)) { - /* Use GOMP_MAP_DELETE/GOMP_MAP_FORCE_FROM to denote that "finalize" - semantics apply to all mappings of this OpenACC directive. */ - bool finalize_marked = false; + /* Use GOMP_MAP_DELETE/GOMP_MAP_FORCE_FROM to denote "finalize" + semantics. */ + bool have_clause = false; for (tree c = OMP_STANDALONE_CLAUSES (expr); c; c = OMP_CLAUSE_CHAIN (c)) if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP) switch (OMP_CLAUSE_MAP_KIND (c)) { case GOMP_MAP_FROM: OMP_CLAUSE_SET_MAP_KIND (c, GOMP_MAP_FORCE_FROM); - finalize_marked = true; + have_clause = true; break; case GOMP_MAP_RELEASE: OMP_CLAUSE_SET_MAP_KIND (c, GOMP_MAP_DELETE); - finalize_marked = true; + have_clause = true; break; - default: - /* Check consistency: libgomp relies on the very first data - mapping clause being marked, so make sure we did that before - any other mapping clauses. */ - gcc_assert (finalize_marked); + case GOMP_MAP_POINTER: + case GOMP_MAP_TO_PSET: + /* TODO PR92929: we may see these here, but they'll always follow + one of the clauses above, and will be handled by libgomp as + one group, so no handling required here. */ + gcc_assert (have_clause); break; + default: + gcc_unreachable (); } } stmt = gimple_build_omp_target (NULL, kind, OMP_STANDALONE_CLAUSES (expr)); diff --git a/gcc/testsuite/c-c++-common/goacc/finalize-1.c b/gcc/testsuite/c-c++-common/goacc/finalize-1.c index 94820290b94..3d64b2e7cb3 100644 --- a/gcc/testsuite/c-c++-common/goacc/finalize-1.c +++ b/gcc/testsuite/c-c++-common/goacc/finalize-1.c @@ -4,8 +4,10 @@ extern int del_r; extern float del_f[3]; +extern char *del_f_p; extern double cpo_r[8]; extern long cpo_f; +extern char *cpo_f_p; void f () { @@ -17,6 +19,10 @@ void f () /* { dg-final { scan-tree-dump-times "(?n)#pragma acc exit data map\\(release:del_f\\) finalize;$" 1 "original" } } { dg-final { scan-tree-dump-times "(?n)#pragma omp target oacc_enter_exit_data map\\(delete:del_f \\\[len: \[0-9\]+\\\]\\) finalize$" 1 "gimple" } } */ +#pragma acc exit data finalize delete (del_f_p[2:5]) +/* { dg-final { scan-tree-dump-times "(?n)#pragma acc exit data map\\(release:\\*\\(del_f_p \\+ 2\\) \\\[len: 5\\\]\\) map\\(firstprivate:del_f_p \\\[pointer assign, bias: 2\\\]\\) finalize;$" 1 "original" } } + { dg-final { scan-tree-dump-times "(?n)#pragma omp target oacc_enter_exit_data map\\(delete:\[^ \]+ \\\[len: 5\\\]\\) finalize$" 1 "gimple" } } */ + #pragma acc exit data copyout (cpo_r) /* { dg-final { scan-tree-dump-times "(?n)#pragma acc exit data map\\(from:cpo_r\\);$" 1 "original" } } { dg-final { scan-tree-dump-times "(?n)#pragma omp target oacc_enter_exit_data map\\(from:cpo_r \\\[len: \[0-9\]+\\\]\\)$" 1 "gimple" } } */ @@ -24,5 +30,8 @@ void f () #pragma acc exit data copyout (cpo_f) finalize /* { dg-final { scan-tree-dump-times "(?n)#pragma acc exit data finalize map\\(from:cpo_f\\);$" 1 "original" } } { dg-final { scan-tree-dump-times "(?n)#pragma omp target oacc_enter_exit_data finalize map\\(force_from:cpo_f \\\[len: \[0-9\]+\\\]\\)$" 1 "gimple" } } */ -} +#pragma acc exit data copyout (cpo_f_p[4:10]) finalize +/* { dg-final { scan-tree-dump-times "(?n)#pragma acc exit data finalize map\\(from:\\*\\(cpo_f_p \\+ 4\\) \\\[len: 10\\\]\\) map\\(firstprivate:cpo_f_p \\\[pointer assign, bias: 4\\\]\\);$" 1 "original" } } + { dg-final { scan-tree-dump-times "(?n)#pragma omp target oacc_enter_exit_data finalize map\\(force_from:\[^ \]+ \\\[len: 10\\\]\\)$" 1 "gimple" } } */ +} diff --git a/gcc/testsuite/gfortran.dg/goacc/finalize-1.f b/gcc/testsuite/gfortran.dg/goacc/finalize-1.f index 5c7a921a2e3..ca642156e9f 100644 --- a/gcc/testsuite/gfortran.dg/goacc/finalize-1.f +++ b/gcc/testsuite/gfortran.dg/goacc/finalize-1.f @@ -6,8 +6,10 @@ IMPLICIT NONE INTEGER :: del_r REAL, DIMENSION (3) :: del_f + INTEGER (1), DIMENSION (:), ALLOCATABLE :: del_f_p DOUBLE PRECISION, DIMENSION (8) :: cpo_r LOGICAL :: cpo_f + INTEGER (1), DIMENSION (:), ALLOCATABLE :: cpo_f_p !$ACC EXIT DATA DELETE (del_r) ! { dg-final { scan-tree-dump-times "(?n)#pragma acc exit data map\\(release:del_r\\);$" 1 "original" } } @@ -17,6 +19,10 @@ ! { dg-final { scan-tree-dump-times "(?n)#pragma acc exit data map\\(release:del_f\\) finalize;$" 1 "original" } } ! { dg-final { scan-tree-dump-times "(?n)#pragma omp target oacc_enter_exit_data map\\(delete:del_f \\\[len: \[0-9\]+\\\]\\) finalize$" 1 "gimple" } } +!$ACC EXIT DATA FINALIZE DELETE (del_f_p(2:5)) +! { dg-final { scan-tree-dump-times "(?n)#pragma acc exit data map\\(release:\\*\\(c_char \\*\\) parm\\.0\\.data \\\[len: \[^\\\]\]+\\\]\\) map\\(to:del_f_p \\\[pointer set, len: \[0-9\]+\\\]\\) map\\(alloc:\\(integer\\(kind=1\\)\\\[0:\\\] \\* restrict\\) del_f_p\\.data \\\[pointer assign, bias: \\(sizetype\\) parm\\.0\\.data - \\(sizetype\\) del_f_p\\.data\\\]\\) finalize;$" 1 "original" } } +! { dg-final { scan-tree-dump-times "(?n)#pragma omp target oacc_enter_exit_data map\\(delete:MEM\\\[\\(c_char \\*\\)\[^\\\]\]+\\\] \\\[len: \[^\\\]\]+\\\]\\) map\\(to:del_f_p \\\[pointer set, len: \[0-9\]+\\\]\\) map\\(alloc:del_f_p\\.data \\\[pointer assign, bias: \[^\\\]\]+\\\]\\) finalize$" 1 "gimple" } } + !$ACC EXIT DATA COPYOUT (cpo_r) ! { dg-final { scan-tree-dump-times "(?n)#pragma acc exit data map\\(from:cpo_r\\);$" 1 "original" } } ! { dg-final { scan-tree-dump-times "(?n)#pragma omp target oacc_enter_exit_data map\\(from:cpo_r \\\[len: \[0-9\]+\\\]\\)$" 1 "gimple" } } @@ -24,4 +30,8 @@ !$ACC EXIT DATA COPYOUT (cpo_f) FINALIZE ! { dg-final { scan-tree-dump-times "(?n)#pragma acc exit data map\\(from:cpo_f\\) finalize;$" 1 "original" } } ! { dg-final { scan-tree-dump-times "(?n)#pragma omp target oacc_enter_exit_data map\\(force_from:cpo_f \\\[len: \[0-9\]+\\\]\\) finalize$" 1 "gimple" } } + +!$ACC EXIT DATA COPYOUT (cpo_f_p(4:10)) FINALIZE +! { dg-final { scan-tree-dump-times "(?n)#pragma acc exit data map\\(from:\\*\\(c_char \\*\\) parm\\.1\\.data \\\[len: \[^\\\]\]+\\\]\\) map\\(to:cpo_f_p \\\[pointer set, len: \[0-9\]+\\\]\\) map\\(alloc:\\(integer\\(kind=1\\)\\\[0:\\\] \\* restrict\\) cpo_f_p\\.data \\\[pointer assign, bias: \\(sizetype\\) parm\\.1\\.data - \\(sizetype\\) cpo_f_p\\.data\\\]\\) finalize;$" 1 "original" } } +! { dg-final { scan-tree-dump-times "(?n)#pragma omp target oacc_enter_exit_data map\\(force_from:MEM\\\[\\(c_char \\*\\)\[^\\\]\]+\\\] \\\[len: \[^\\\]\]+\\\]\\) map\\(to:cpo_f_p \\\[pointer set, len: \[0-9\]+\\\]\\) map\\(alloc:cpo_f_p\\.data \\\[pointer assign, bias: \[^\\\]\]+\\\]\\) finalize$" 1 "gimple" } } END SUBROUTINE f diff --git a/libgomp/oacc-mem.c b/libgomp/oacc-mem.c index 311f9585f77..291ef9192b9 100644 --- a/libgomp/oacc-mem.c +++ b/libgomp/oacc-mem.c @@ -1058,17 +1058,6 @@ GOACC_enter_exit_data (int flags_m, size_t mapnum, void **hostaddrs, thr = goacc_thread (); acc_dev = thr->dev; - /* Determine whether "finalize" semantics apply to all mappings of this - OpenACC directive. */ - bool finalize = false; - if (mapnum > 0) - { - unsigned char kind = kinds[0] & 0xff; - if (kind == GOMP_MAP_DELETE - || kind == GOMP_MAP_FORCE_FROM) - finalize = true; - } - /* Determine if this is an "acc enter data". */ for (i = 0; i < mapnum; ++i) { @@ -1221,6 +1210,9 @@ GOACC_enter_exit_data (int flags_m, size_t mapnum, void **hostaddrs, { unsigned char kind = kinds[i] & 0xff; + bool finalize = (kind == GOMP_MAP_DELETE + || kind == GOMP_MAP_FORCE_FROM); + int pointer = find_pointer (i, mapnum, kinds); if (!pointer) -- 2.17.1