From mboxrd@z Thu Jan 1 00:00:00 1970 Return-Path: Received: (qmail 33876 invoked by alias); 31 May 2016 07:28:58 -0000 Mailing-List: contact gcc-patches-help@gcc.gnu.org; run by ezmlm Precedence: bulk List-Id: List-Archive: List-Post: List-Help: Sender: gcc-patches-owner@gcc.gnu.org Received: (qmail 33855 invoked by uid 89); 31 May 2016 07:28:57 -0000 Authentication-Results: sourceware.org; auth=none X-Virus-Found: No X-Spam-SWARE-Status: No, score=-1.8 required=5.0 tests=AWL,BAYES_00,RCVD_IN_DNSWL_NONE,SPF_PASS autolearn=ham version=3.3.2 spammy=UD:gmane.org, saved, sk:tree_ad, sk:TREE_AD 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 (AES256-GCM-SHA384 encrypted) ESMTPS; Tue, 31 May 2016 07:28:47 +0000 Received: from svr-orw-fem-03.mgc.mentorg.com ([147.34.97.39]) by relay1.mentorg.com with esmtp id 1b7e6V-0002lc-Ae from Thomas_Schwinge@mentor.com ; Tue, 31 May 2016 00:28:43 -0700 Received: from tftp-cs (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; Tue, 31 May 2016 00:28:42 -0700 Received: by tftp-cs (Postfix, from userid 49978) id 56A42C22DF; Tue, 31 May 2016 00:28:42 -0700 (PDT) From: Thomas Schwinge To: Jakub Jelinek , Chung-Lin Tang , Cesar Philippidis CC: gcc-patches Subject: Re: [PATCH, OpenACC] Make reduction arguments addressable In-Reply-To: <20160530165341.GS28550@tucnak.redhat.com> References: <20160530165341.GS28550@tucnak.redhat.com> User-Agent: Notmuch/0.9-125-g4686d11 (http://notmuchmail.org) Emacs/24.5.1 (i586-pc-linux-gnu) Date: Tue, 31 May 2016 09:59:00 -0000 Message-ID: <878tyqsn32.fsf@kepler.schwinge.homeip.net> MIME-Version: 1.0 Content-Type: multipart/signed; boundary="=-=-="; micalg=pgp-sha1; protocol="application/pgp-signature" X-SW-Source: 2016-05/txt/msg02380.txt.bz2 --=-=-= Content-Type: text/plain; charset=utf-8 Content-Transfer-Encoding: quoted-printable Content-length: 3103 Hi! On Mon, 30 May 2016 18:53:41 +0200, Jakub Jelinek wrote: > On Mon, May 30, 2016 at 10:38:59PM +0800, Chung-Lin Tang wrote: > > Hi, a previous patch of Cesar's has made the middle-end omp-lowering > > automatically create and insert a tofrom (i.e. present_or_copy) map for > > parallel reductions. This allowed the user to not need explicit > > clauses to copy out the reduction result, but because reduction argumen= ts > > are not marked addressable, async does not work as expected, > > i.e. the asynchronous copy-out results are not used in the compiler gen= erated code. >=20 > If you need it only for async parallel/kernels? regions, can't you do that > only for those and not for others? Also, please add comments to the source code to document the need for such special handling. > > This patch fixes this in the front-ends, I've tested this patch without > > new regressions, and fixes some C++ OpenACC tests that regressed after > > my last OpenACC async patch. Is this okay for trunk? >=20 > Testcases in the testsuite or others? If the latter, we should add them. The r236772 commit "[PATCH, libgomp] Rewire OpenACC async", regressed (or, triggered/exposed the existing wrong behavior?) libgomp.oacc-c++/template-reduction.C execution testing for nvptx offloading. (Please always send email about such known regressions, and XFAIL them with your commit -- that would have saved me an hour yesterday, when I bisected recent changes to figure out why that test suddenly fails.) For reference, here is a test case, a reduced C version of libgomp/testsuite/libgomp.oacc-c++/template-reduction.C. This test case works (without Chung-Lin's "[PATCH, OpenACC] Make reduction arguments addressable" patch) if I enable "POCs", which surprises me a bit, because I thought after Cesar's recent changes, the gimplifier is doing the same thing of adding a data clause next to the reduction clause. Probably it's not doing the exactly same thing, though. Should it? Cesar, do you have any comments on this? For example (just guessing), should TREE_ADDRESSABLE be set where the gimplifier does its work, instead of in the three front ends? // Reduced C version of libgomp/testsuite/libgomp.oacc-c++/template-red= uction.C. =20=20=20=20 const int n =3D 100; =20=20=20=20 // Check present and async and an explicit firstprivate =20=20=20=20 int async_sum (int c) { int s =3D 0; =20=20=20=20 #define POCs //present_or_copy(s) #pragma acc parallel loop num_gangs (10) gang reduction (+:s) POCs firs= tprivate (c) async for (int i =3D 0; i < n; i++) s +=3D i+c; =20=20=20=20 #pragma acc wait =20=20=20=20 return s; } =20=20=20=20 int main() { int result =3D 0; =20=20=20=20 for (int i =3D 0; i < n; i++) { result +=3D i+1; } =20=20=20=20 if (async_sum (1) !=3D result) __builtin_abort (); =20=20=20=20 return 0; } Gr=C3=BC=C3=9Fe Thomas --=-=-= Content-Type: application/pgp-signature; name="signature.asc" Content-length: 472 -----BEGIN PGP SIGNATURE----- Version: GnuPG v1 iQEcBAEBAgAGBQJXTT0iAAoJEK3/DN1sMFFtS74H/2oPrFJp+y/qxS4sbojkPlEQ YDfER3mwHfLAXW4ht3cEi+ALT+I5/C8+ZqF2LWrPW75VnHYHlx0sCT9nXN8a+vWy gQaWF/PywbVJObstjlIMKQbZvIwl3AgKLCBznJLSZ6ssebdKgLtlffi61SjZCNB9 54s7A1Eml2gvAuUfbYfXcOC4H+MhzDP0l/AzhKxMquHXEA2seBELo7sRdOd3AIwZ GB2HyCAZ5MKHndnn6JyNwt2/5ETmO+RqNKCAQ5ne7ERA2CUnxT7kybSEomI+Gwsk 03dokOTca6+XOb+Z6JE/38x+2JuTM9V+3enqGD4dqOd26QcsUZm5/1M6xzdybZM= =nLeK -----END PGP SIGNATURE----- --=-=-=--