From mboxrd@z Thu Jan 1 00:00:00 1970 Return-Path: Received: from esa3.mentor.iphmx.com (esa3.mentor.iphmx.com [68.232.137.180]) by sourceware.org (Postfix) with ESMTPS id 8C31338618A5 for ; Tue, 6 Oct 2020 14:44:39 +0000 (GMT) DMARC-Filter: OpenDMARC Filter v1.3.2 sourceware.org 8C31338618A5 Authentication-Results: sourceware.org; dmarc=none (p=none dis=none) header.from=mentor.com Authentication-Results: sourceware.org; spf=pass smtp.mailfrom=Tobias_Burnus@mentor.com IronPort-SDR: SAywe6jJP4YTYIBl41aCWDRR6OX+9wo2cmYWN0+63UaEoASdr1S4E7Fj3wZUQsTDTsxv+FGG6r 4bArSInS/jDSS/2mKGiKBX45xepvH6mfm5F5KDEqfktfNM/oRLcLYSKgRHIaaaeVgE1oaNKayA 1/ik9yiXZSw72Le2GyFLD03s4lXGDKpFeUvykwYmW0WXQ75MSVPUsSH1x4m5bo9QEaZKYrFrId fi6dF9zASPCdryozXh2YFnTeuLXXpFDt4N/TpsUQYU5l6EYiZnNfp4sx8OvgIFkUbK6rhOQgMm plQ= X-IronPort-AV: E=Sophos;i="5.77,343,1596528000"; d="scan'208";a="53606422" Received: from orw-gwy-02-in.mentorg.com ([192.94.38.167]) by esa3.mentor.iphmx.com with ESMTP; 06 Oct 2020 06:44:38 -0800 IronPort-SDR: JmOIh6VhZjtm/Pimd7W9tC7RfpcjtH2KY6/xVYKotzvKhy1EeK0eA8XDfRF9deE9jk9JIfAAh1 UOUSOoYCUEtgCnUJHTJLkUPQwse6z4SthQPmGI7WzCz0qenHSBUTcoXtfDQyKBUUnC0jY16utF Rt9MfQSVDxOqOz9/ZjhVWOF9BCbC/weFcPgW9GF1Su7PFakNFbu89Nb/yj081E+wfaHxiu8TlJ AeJnwXxKb7c16+03qDNy2nGXpbjJMOaaVa7QcwJT/xCyIYqn4s12SKaHt8BUm2Rq63fZ7wth2K 1O4= Subject: Re: [PATCH][openacc] Fix acc declare for VLAs To: Tom de Vries , CC: Jakub Jelinek , Thomas Schwinge References: <20201006132849.GA16366@delia> From: Tobias Burnus Message-ID: Date: Tue, 6 Oct 2020 16:44:29 +0200 User-Agent: Mozilla/5.0 (X11; Linux x86_64; rv:68.0) Gecko/20100101 Thunderbird/68.12.0 MIME-Version: 1.0 In-Reply-To: <20201006132849.GA16366@delia> Content-Type: text/plain; charset="utf-8"; format=flowed Content-Transfer-Encoding: quoted-printable Content-Language: en-US X-Originating-IP: [137.202.0.90] X-ClientProxiedBy: SVR-IES-MBX-03.mgc.mentorg.com (139.181.222.3) To svr-ies-mbx-01.mgc.mentorg.com (139.181.222.1) X-Spam-Status: No, score=-11.8 required=5.0 tests=BAYES_00, GIT_PATCH_0, KAM_DMARC_STATUS, NICE_REPLY_A, SPF_HELO_PASS, SPF_PASS, TXREP autolearn=ham autolearn_force=no version=3.4.2 X-Spam-Checker-Version: SpamAssassin 3.4.2 (2018-09-13) on server2.sourceware.org X-BeenThere: gcc-patches@gcc.gnu.org X-Mailman-Version: 2.1.29 Precedence: list List-Id: Gcc-patches mailing list List-Unsubscribe: , List-Archive: List-Post: List-Help: List-Subscribe: , X-List-Received-Date: Tue, 06 Oct 2020 14:44:41 -0000 LGTM. Thanks, Tobias On 10/6/20 3:28 PM, Tom de Vries wrote: > Hi, > > Consider test-case test.c, with VLA A: > ... > int main (void) { > int N =3D 1000; > int A[N]; > #pragma acc declare copy(A) > return 0; > } > ... > compiled using: > ... > $ gcc test.c -fopenacc -S -fdump-tree-all > ... > > At original, we have: > ... > #pragma acc declare map(tofrom:A); > ... > but at gimple, we have a map (to:A.1), but not a map (from:A.1): > ... > int[0:D.2074] * A.1; > > { > int A[0:D.2074] [value-expr: *A.1]; > > saved_stack.2 =3D __builtin_stack_save (); > try > { > A.1 =3D __builtin_alloca_with_align (D.2078, 32); > #pragma omp target oacc_declare map(to:(*A.1) [len: D.2076]) > } > finally > { > __builtin_stack_restore (saved_stack.2); > } > } > ... > > This is caused by the following incompatibility. When storing the desire= d > from clause in oacc_declare_returns, we use 'A.1' as the key: > ... > 10898 oacc_declare_returns->put (decl, c); > (gdb) call debug_generic_expr (decl) > A.1 > (gdb) call debug_generic_expr (c) > map(from:(*A.1)) > ... > but when looking it up, we use 'A' as the key: > ... > (gdb) > 1471 tree *c =3D oacc_declare_returns->get (t); > (gdb) call debug_generic_expr (t) > A > ... > > Fix this by extracing the 'A.1' lookup key from 'A' using the decl-expr. > > In addition, unshare the looked up value, to fix avoid running into > an "incorrect sharing of tree nodes" error. > > Using these two fixes, we get our desired: > ... > finally > { > + #pragma omp target oacc_declare map(from:(*A.1)) > __builtin_stack_restore (saved_stack.2); > } > ... > > Build on x86_64-linux with nvptx accelerator, tested libgomp. > > OK for trunk? > > Thanks, > - Tom > > [openacc] Fix acc declare for VLAs > > gcc/ChangeLog: > > 2020-10-06 Tom de Vries > > PR middle-end/90861 > * gimplify.c (gimplify_bind_expr): Handle lookup in > oacc_declare_returns using key with decl-expr. > > libgomp/ChangeLog: > > 2020-10-06 Tom de Vries > > PR middle-end/90861 > * testsuite/libgomp.oacc-c-c++-common/declare-vla.c: Remove xfail. > > --- > gcc/gimplify.c | 13 ++++++++= ++--- > libgomp/testsuite/libgomp.oacc-c-c++-common/declare-vla.c | 5 ----- > 2 files changed, 10 insertions(+), 8 deletions(-) > > diff --git a/gcc/gimplify.c b/gcc/gimplify.c > index 2dea03cce3d..fa89e797940 100644 > --- a/gcc/gimplify.c > +++ b/gcc/gimplify.c > @@ -1468,15 +1468,22 @@ gimplify_bind_expr (tree *expr_p, gimple_seq *pre= _p) > > if (flag_openacc && oacc_declare_returns !=3D NULL) > { > - tree *c =3D oacc_declare_returns->get (t); > + tree key =3D t; > + if (DECL_HAS_VALUE_EXPR_P (key)) > + { > + key =3D DECL_VALUE_EXPR (key); > + if (TREE_CODE (key) =3D=3D INDIRECT_REF) > + key =3D TREE_OPERAND (key, 0); > + } > + tree *c =3D oacc_declare_returns->get (key); > if (c !=3D NULL) > { > if (ret_clauses) > OMP_CLAUSE_CHAIN (*c) =3D ret_clauses; > > - ret_clauses =3D *c; > + ret_clauses =3D unshare_expr (*c); > > - oacc_declare_returns->remove (t); > + oacc_declare_returns->remove (key); > > if (oacc_declare_returns->is_empty ()) > { > diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/declare-vla.c b/= libgomp/testsuite/libgomp.oacc-c-c++-common/declare-vla.c > index 0f51badca42..714935772c1 100644 > --- a/libgomp/testsuite/libgomp.oacc-c-c++-common/declare-vla.c > +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/declare-vla.c > @@ -59,8 +59,3 @@ main () > > return 0; > } > - > - > -/* { dg-xfail-run-if "TODO PR90861" { *-*-* } { "-DACC_MEM_SHARED=3D0" }= } > - This might XPASS if the compiler happens to put the two 'A' VLAs at t= he same > - address. */ ----------------- Mentor Graphics (Deutschland) GmbH, Arnulfstra=C3=9Fe 201, 80634 M=C3=BCnch= en / Germany Registergericht M=C3=BCnchen HRB 106955, Gesch=C3=A4ftsf=C3=BChrer: Thomas = Heurung, Alexander Walter