From mboxrd@z Thu Jan 1 00:00:00 1970 Return-Path: Received: (qmail 20623 invoked by alias); 26 Jun 2019 16:23:30 -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 20608 invoked by uid 89); 26 Jun 2019 16:23:30 -0000 Authentication-Results: sourceware.org; auth=none X-Spam-SWARE-Status: No, score=-9.3 required=5.0 tests=AWL,BAYES_00,GIT_PATCH_2,GIT_PATCH_3,RCVD_IN_DNSWL_NONE,SPF_PASS autolearn=ham version=3.3.1 spammy=2015, Ilya, ilya, *p2 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, 26 Jun 2019 16:23:28 +0000 Received: from svr-orw-mbx-06.mgc.mentorg.com ([147.34.90.206]) by relay1.mentorg.com with esmtps (TLSv1.2:ECDHE-RSA-AES256-SHA384:256) id 1hgAhb-0003Mg-9C from Thomas_Schwinge@mentor.com ; Wed, 26 Jun 2019 09:23:19 -0700 Received: from svr-orw-mbx-08.mgc.mentorg.com (147.34.90.208) by SVR-ORW-MBX-06.mgc.mentorg.com (147.34.90.206) with Microsoft SMTP Server (TLS) id 15.0.1320.4; Wed, 26 Jun 2019 09:23:17 -0700 Received: from tftp-cs (147.34.91.1) by svr-orw-mbx-08.mgc.mentorg.com (147.34.90.208) with Microsoft SMTP Server id 15.0.1320.4 via Frontend Transport; Wed, 26 Jun 2019 09:23:17 -0700 Received: by tftp-cs (Postfix, from userid 49978) id AA39FC2212; Wed, 26 Jun 2019 09:23:16 -0700 (PDT) From: Thomas Schwinge To: Ilya Verbin , Jakub Jelinek , Tom de Vries , Alexander Monakov CC: , Kirill Yukhin Subject: Re: [gomp4.5] Handle #pragma omp declare target link In-Reply-To: <20151214171733.GB63018@msticlxl57.ims.intel.com> References: <20151026183552.GD35077@msticlxl57.ims.intel.com> <20151026190539.GX478@tucnak.redhat.com> <20151026193904.GE35077@msticlxl57.ims.intel.com> <20151026194940.GZ478@tucnak.redhat.com> <20151116154043.GA18854@msticlxl57.ims.intel.com> <20151119153115.GZ5675@tucnak.redhat.com> <20151127165009.GA24771@msticlxl57.ims.intel.com> <20151130120459.GB5675@tucnak.redhat.com> <20151130202934.GA22962@msticlxl57.ims.intel.com> <20151130204902.GJ5675@tucnak.redhat.com> <20151214171733.GB63018@msticlxl57.ims.intel.com> User-Agent: Notmuch/0.9-125-g4686d11 (http://notmuchmail.org) Emacs/26.1 (x86_64-pc-linux-gnu) Date: Wed, 26 Jun 2019 16:23:00 -0000 Message-ID: <87woh848pi.fsf@euler.schwinge.homeip.net> MIME-Version: 1.0 Content-Type: multipart/signed; boundary="=-=-="; micalg=pgp-sha512; protocol="application/pgp-signature" X-SW-Source: 2019-06/txt/msg01675.txt.bz2 --=-=-= Content-Type: text/plain; charset=utf-8 Content-Transfer-Encoding: quoted-printable Content-length: 1965 Hi! On Mon, 14 Dec 2015 20:17:33 +0300, Ilya Verbin wrote: > Here is an updated patch [for "#pragma omp declare target link"] ..., that got committed long ago (trunk r231655), with additional changes later on. As has later been filed in PR81689, the test case added "libgomp.c/target-link-1.c fails for nvptx: #pragma omp target link not implemented". Curious, has anybody ever looked into what's going on/wrong? Gr=C3=BC=C3=9Fe Thomas > --- /dev/null > +++ b/libgomp/testsuite/libgomp.c/target-link-1.c > @@ -0,0 +1,63 @@ > +struct S { int s, t; }; > + > +int a =3D 1, b =3D 1; > +double c[27]; > +struct S d =3D { 8888, 8888 }; > +#pragma omp declare target link (a) to (b) link (c, d) > + > +int > +foo (void) > +{ > + return a++ + b++; > +} > + > +int > +bar (int n) > +{ > + int *p1 =3D &a; > + int *p2 =3D &b; > + c[n] +=3D 2.0; > + d.s -=3D 2; > + d.t -=3D 2; > + return *p1 + *p2 + d.s + d.t; > +} > + > +#pragma omp declare target (foo, bar) > + > +int > +main () > +{ > + a =3D b =3D 2; > + d.s =3D 17; > + d.t =3D 18; > + > + int res, n =3D 10; > + #pragma omp target map (to: a, b, c, d) map (from: res) > + { > + res =3D foo () + foo (); > + c[n] =3D 3.0; > + res +=3D bar (n); > + } > + > + int shared_mem =3D 0; > + #pragma omp target map (alloc: shared_mem) > + shared_mem =3D 1; > + > + if ((shared_mem && res !=3D (2 + 2) + (3 + 3) + (4 + 4 + 15 + 16)) > + || (!shared_mem && res !=3D (2 + 1) + (3 + 2) + (4 + 3 + 15 + 16))) > + __builtin_abort (); > + > + #pragma omp target enter data map (to: c) > + #pragma omp target update from (c) > + res =3D (int) (c[n] + 0.5); > + if ((shared_mem && res !=3D 5) || (!shared_mem && res !=3D 0)) > + __builtin_abort (); > + > + #pragma omp target map (to: a, b) map (from: res) > + res =3D foo (); > + > + if ((shared_mem && res !=3D 4 + 4) || (!shared_mem && res !=3D 2 + 3)) > + __builtin_abort (); > + > + return 0; > +} --=-=-= Content-Type: application/pgp-signature; name="signature.asc" Content-length: 658 -----BEGIN PGP SIGNATURE----- iQGzBAEBCgAdFiEEU9WEfWKGQazCmycCAKI7+41Q4XkFAl0Tm+sACgkQAKI7+41Q 4XmWCgv/SBVBi9KW+qTOBy3MEF4XGkhcFt5puZD8XBSoUEkuH9h8GWNKpTt4mbFB NURyCNsgzW+qpijGqpnjUMK/08Ll0JS5YuDS5kMsw9Nke2lYt3zFBXrGY0HvNq+k /CB/x3BFav/Q67SjLuuSXIpOq4Iz3TLa5DmypSL9w4HzSmRYiWLlf+F1flVSfycD CR8yPxPBowZa50RgF08weWOQVIOHo82dS1JbtEs06GcGnykx4Z2bs5u8JftwG/xO Mj5xRnCqjxxuD1A9CuyTjCl1Qbko4XyxQLv6gi0f1QXoYiyKulXqbP4Vvse2aA/S E2XQrG39oJob60Dwxd7B9jumB38H4Iaan5ajs8OUkXLDKsAnirRN/yWFJxUQeUV3 QS7JwjMvGLWCtKKNw5nt1CNN1IZWp6bA7DhCn2Fh4DpquSRx9o5dALypGzytWUle ss4SH+dheag+S/2NIlnUxBLhqLXl4Tfn3sNTGghOIIKInnFi8tRCrEBlzryXLA4L qhKctJ8G =8+0l -----END PGP SIGNATURE----- --=-=-=--