From mboxrd@z Thu Jan 1 00:00:00 1970 Return-Path: Received: (qmail 24885 invoked by alias); 9 Oct 2019 11:35:32 -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 24780 invoked by uid 89); 9 Oct 2019 11:35:32 -0000 Authentication-Results: sourceware.org; auth=none X-Spam-SWARE-Status: No, score=-15.9 required=5.0 tests=AWL,BAYES_00,GIT_PATCH_0,GIT_PATCH_1,GIT_PATCH_2,GIT_PATCH_3,SPF_PASS autolearn=ham version=3.3.1 spammy= 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; Wed, 09 Oct 2019 11:35:31 +0000 IronPort-SDR: rX1c0qArwstP6S2RDiBG6Y3fU2RzXqN7PueiUkgqvCsuzkCc5xGBY+aXmWhXftrLIntdugVRKU s/9k0TroN4SLehbySQ0G+JodGgYphLmthRLm/qwm/A/fMVnPsdwZAXkSFcOpv0e+XCF95okaQ0 Y3XDLiYkho8Sj/pTIXXXOyouCirnsRPy7Z+ebUi3TDgdELs26FgPxPtouzAqg5I122Cq9Kd+w4 znvuKU4ZNTQulOkXOrFXhCVFZaaWyd1osVyxuAkiaKJMk/zYiMI04RRL0pHmMUhXvNhAB0TN9c 6Zk= Received: from orw-gwy-01-in.mentorg.com ([192.94.38.165]) by esa3.mentor.iphmx.com with ESMTP; 09 Oct 2019 03:35:29 -0800 IronPort-SDR: v9cosRdQT1+KVAL1f4WAJBFO6gh9SswdUV5EAzVZXJZa90eb9yX8R1zUHx+TRwrCDBkEYVttlH N2Ocm54sy96lhLv2PfbYKRlut5wqk5jysJaqDiiudWzgPTnK13qI2vHhNlWc/ICGqCmCdq8ohH N2frTW8sfpsYxPnB89wDe/TT+c8BVkPuhIY1HYxYR2zQGSlKljdUDs2mrEvr+kRAqa+5RvhCPu hc2bRNUBcqorlbNbIc+eNVyZzhnCHZvFFshRdwFeNu9iBzIbqh8SimszRzdweukJR0z2NuFPJU N3I= From: Thomas Schwinge To: CC: Julian Brown , Jakub Jelinek Subject: PR92036 "OpenACC 'firstprivate' clause: initial value" User-Agent: Notmuch/0.29.1+93~g67ed7df (https://notmuchmail.org) Emacs/26.1 (x86_64-pc-linux-gnu) Date: Wed, 09 Oct 2019 11:43:00 -0000 Message-ID: <87o8yqywwf.fsf@euler.schwinge.homeip.net> MIME-Version: 1.0 Content-Type: multipart/signed; boundary="==-=-="; micalg=pgp-sha512; protocol="application/pgp-signature" Return-Path: tschwing@mentor.com X-SW-Source: 2019-10/txt/msg00606.txt.bz2 --==-=-= Content-Type: multipart/mixed; boundary="=-=-=" --=-=-= Content-Type: text/plain; charset=utf-8 Content-Transfer-Encoding: quoted-printable Content-length: 267 Hi! In r276757 "[PR92036] Add 'libgomp.oacc-c-c++-common/data-firstprivate-1.c'", I committed the attached to document the status quo, and then what I think it instead should be per PR92036 "OpenACC 'firstprivate' clause: initial value". Gr=C3=BC=C3=9Fe Thomas --=-=-= Content-Type: text/x-diff Content-Disposition: inline; filename=0001-PR92036-Add-libgomp.oacc-c-c-common-data-first.trunk.patch Content-Transfer-Encoding: quoted-printable Content-length: 5309 =46rom 5fc105387fda327bf5e36ae6f997e415da6d1f37 Mon Sep 17 00:00:00 2001 From: tschwinge Date: Wed, 9 Oct 2019 11:31:14 +0000 Subject: [PATCH] [PR92036] Add 'libgomp.oacc-c-c++-common/data-firstprivate-1.c' libgomp/ PR middle-end/92036 * testsuite/libgomp.oacc-c-c++-common/data-firstprivate-1.c: New file. git-svn-id: svn+ssh://gcc.gnu.org/svn/gcc/trunk@276757 138bc75d-0d04-0410-9= 61f-82ee72b054a4 --- libgomp/ChangeLog | 6 + .../data-firstprivate-1.c | 165 ++++++++++++++++++ 2 files changed, 171 insertions(+) create mode 100644 libgomp/testsuite/libgomp.oacc-c-c++-common/data-firstp= rivate-1.c diff --git a/libgomp/ChangeLog b/libgomp/ChangeLog index 1b43c456741..319a1911882 100644 --- a/libgomp/ChangeLog +++ b/libgomp/ChangeLog @@ -1,3 +1,9 @@ +2019-10-09 Thomas Schwinge + + PR middle-end/92036 + * testsuite/libgomp.oacc-c-c++-common/data-firstprivate-1.c: New + file. + 2019-10-09 Tobias Burnus =20 PR testsuite/91884 diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/data-firstprivate-= 1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/data-firstprivate-1.c new file mode 100644 index 00000000000..8900a4e070d --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/data-firstprivate-1.c @@ -0,0 +1,165 @@ +/* Test behavior of 'firstprivate' lexically vs. dynamically nested inside= a + 'data' region. */ + +#include + + +#define VERIFY(x) \ + do { \ + if (!(x)) \ + abort (); \ + } while (0); + + +/* This is basically and extended version of 't2' from 'firstprivate-1.c'.= */ + +int lexically_nested_val =3D 2; + +static void +lexically_nested () +{ +#pragma acc data \ + copy (lexically_nested_val) + { + VERIFY (lexically_nested_val =3D=3D 2); + +#pragma acc parallel \ + present (lexically_nested_val) + { + VERIFY (lexically_nested_val =3D=3D 2); + + /* This updates the device copy, or shared variable. */ + lexically_nested_val =3D 7; + } + +#if ACC_MEM_SHARED + VERIFY (lexically_nested_val =3D=3D 7); +#else + VERIFY (lexically_nested_val =3D=3D 2); +#endif + + /* This only updates the local/shared variable, but not the device + copy. */ + lexically_nested_val =3D 5; + +#pragma acc parallel \ + firstprivate (lexically_nested_val) + { +#if 1 /* Current behavior. */ + /* The 'firstprivate' copy is initialized from the device copy, or + shared variable. */ +# if ACC_MEM_SHARED + VERIFY (lexically_nested_val =3D=3D 5); +# else + VERIFY (lexically_nested_val =3D=3D 7); +# endif +#else /* Expected behavior per PR92036. */ + /* The 'firstprivate' copy is initialized from the local thread. */ + VERIFY (lexically_nested_val =3D=3D 5); +#endif + + /* This updates the 'firstprivate' copy only, but not the shared + variable. */ + lexically_nested_val =3D 9; + } + + VERIFY (lexically_nested_val =3D=3D 5); + } + /* If not shared, the device copy has now been copied back. */ + +#if ACC_MEM_SHARED + VERIFY (lexically_nested_val =3D=3D 5); +#else + VERIFY (lexically_nested_val =3D=3D 7); +#endif +} + + +int dynamically_nested_val =3D 2; + +/* Same as above, but compute construct 1 broken out, so no longer lexical= ly + nested inside 'data' region. */ + +static void +dynamically_nested_compute_1 () +{ +#pragma acc parallel \ + present (dynamically_nested_val) + { + VERIFY (dynamically_nested_val =3D=3D 2); + + /* This updates the device copy, or shared variable. */ + dynamically_nested_val =3D 7; + } +} + +/* Same as above, but compute construct 2 broken out, so no longer lexical= ly + nested inside 'data' region. */ + +static void +dynamically_nested_compute_2 () +{ +#pragma acc parallel \ + firstprivate (dynamically_nested_val) + { +#if 1 /* Current behavior. */ + /* The 'firstprivate' copy is initialized from the device copy, or s= hared + variable. */ +# if ACC_MEM_SHARED + VERIFY (dynamically_nested_val =3D=3D 5); +# else + VERIFY (dynamically_nested_val =3D=3D 7); +# endif +#else /* Expected behavior per PR92036. */ + /* The 'firstprivate' copy is initialized from the local thread. */ + VERIFY (dynamically_nested_val =3D=3D 5); +#endif + + /* This updates the 'firstprivate' copy only, but not the shared + variable. */ + dynamically_nested_val =3D 9; + } +} + +static void +dynamically_nested () +{ +#pragma acc data \ + copy (dynamically_nested_val) + { + VERIFY (dynamically_nested_val =3D=3D 2); + + dynamically_nested_compute_1 (); + +#if ACC_MEM_SHARED + VERIFY (dynamically_nested_val =3D=3D 7); +#else + VERIFY (dynamically_nested_val =3D=3D 2); +#endif + + /* This only updates the local/shared variable, but not the device + copy. */ + dynamically_nested_val =3D 5; + + dynamically_nested_compute_2 (); + + VERIFY (dynamically_nested_val =3D=3D 5); + } + /* If not shared, the device copy has now been copied back. */ + +#if ACC_MEM_SHARED + VERIFY (dynamically_nested_val =3D=3D 5); +#else + VERIFY (dynamically_nested_val =3D=3D 7); +#endif +} + + +int +main() +{ + lexically_nested (); + dynamically_nested (); + + return 0; +} --=20 2.17.1 --=-=-=-- --==-=-= Content-Type: application/pgp-signature; name="signature.asc" Content-length: 658 -----BEGIN PGP SIGNATURE----- iQGzBAEBCgAdFiEEU9WEfWKGQazCmycCAKI7+41Q4XkFAl2dxfAACgkQAKI7+41Q 4XkGBQv+KqFfuKgouzjrEw4l2nPThPhPrRAqQn1298aFv5mKB+CcITFe4g0cNo3K z9TqjrUqIH1Hvi/qynMcqeAmZ/28ZJQ6SMDwN6trZ3S8pm2mFcbhVfKnu3CEbcGA oMajEJ9fMrhtUQbLJBBlpFDjTiXi/o/B0Qoq3xSqHvjQljPpZsux5bVeiwUiOn3P B/XO1TIixn6in6yQYwrFxChspzasSGjaGVbE2z/71RAnBysd+yg79aKC8s9lSKnV ZZX5IyCf9gYPPbgCAGbvP+QonSxs7bJ9OkLdm23JVu0CGnoAMnsLrorZiecP5WZQ fTTRtWnQmdICm7V4ckPmhE/0F4t+QwXDCv/KCx+Krm/fEGOtdUGkdexQ4hICvZzg CDmeHW9WszI0XC3jMjbZd5TKJwUgL7rEVLYaSDL7gQnJq/LuoIefEvbUkPepom8N CI1hNHyli9aOZd80ojlR4UX6Zu26rSQcmp+uH5YSm488TFoFldkGcreD3co7ppEb i/+DNeL9 =UGjz -----END PGP SIGNATURE----- --==-=-=--