From mboxrd@z Thu Jan 1 00:00:00 1970 Return-Path: Received: from esa4.mentor.iphmx.com (esa4.mentor.iphmx.com [68.232.137.252]) by sourceware.org (Postfix) with ESMTPS id 935FB3858D32 for ; Thu, 7 Jul 2022 09:36:48 +0000 (GMT) DMARC-Filter: OpenDMARC Filter v1.4.1 sourceware.org 935FB3858D32 Authentication-Results: sourceware.org; dmarc=none (p=none dis=none) header.from=codesourcery.com Authentication-Results: sourceware.org; spf=pass smtp.mailfrom=mentor.com X-IronPort-AV: E=Sophos;i="5.92,252,1650960000"; d="scan'208";a="78428226" Received: from orw-gwy-02-in.mentorg.com ([192.94.38.167]) by esa4.mentor.iphmx.com with ESMTP; 07 Jul 2022 01:36:41 -0800 IronPort-SDR: yzVw2DE6oTJplJ0OF9PsmLS9YYq0ZMbMw2/SAzffD6nZFaLczIdjmStd470WgXfDAm5bjsfpVs HcFmstAOJA9Eo2hpiDqw1U1zyi/pPOn2iVlGSNuUwhdJ6jYDET774P98NF/PvJ/eACQtjvCaJq p3KHddI9XY7rCmmlWsq3eBrCSfDTLJHYmAUv8ZFMyXzJXCIXEoqPaBVYNdgw62YwqCFIveYE0I 2di70ycBFWfbAme2AawxL6EIsIIyzMsl2vN+gQ0Ew0xFBqI3gAJqTOAl1MGWQN4XTtLdfI8ziB +1U= Message-ID: <92609eba-2c1b-99b4-c44d-99c438fabb59@codesourcery.com> Date: Thu, 7 Jul 2022 11:36:34 +0200 MIME-Version: 1.0 User-Agent: Mozilla/5.0 (X11; Linux x86_64; rv:91.0) Gecko/20100101 Thunderbird/91.11.0 Subject: Re: Enhance 'libgomp.c-c++-common/requires-4.c', 'libgomp.c-c++-common/requires-5.c' testing (was: [Patch][v4] OpenMP: Move omp requires checks to libgomp) Content-Language: en-US To: Thomas Schwinge , CC: Jakub Jelinek References: <07fec82a-41cf-fdc5-6307-c068dd95ef1a@mentor.com> <7f9c91c1-a479-f94f-ac14-1d6827ce671b@codesourcery.com> <5576fa00-0ddd-8046-17c1-d1cea82bdcf5@codesourcery.com> <87k08p48mb.fsf@euler.schwinge.homeip.net> From: Tobias Burnus In-Reply-To: <87k08p48mb.fsf@euler.schwinge.homeip.net> Content-Type: text/plain; charset="UTF-8"; format=flowed Content-Transfer-Encoding: quoted-printable X-Originating-IP: [137.202.0.90] X-ClientProxiedBy: SVR-IES-MBX-07.mgc.mentorg.com (139.181.222.7) To svr-ies-mbx-12.mgc.mentorg.com (139.181.222.12) X-Spam-Status: No, score=-5.4 required=5.0 tests=BAYES_00, HEADER_FROM_DIFFERENT_DOMAINS, KAM_DMARC_STATUS, KAM_SHORT, NICE_REPLY_A, SPF_HELO_PASS, SPF_PASS, TXREP, T_SCC_BODY_TEXT_LINE autolearn=no autolearn_force=no version=3.4.6 X-Spam-Checker-Version: SpamAssassin 3.4.6 (2021-04-09) 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: Thu, 07 Jul 2022 09:36:50 -0000 On 07.07.22 10:42, Thomas Schwinge wrote: > In preparation for other changes: ... > On 2022-06-29T16:33:02+0200, Tobias Burnus wrot= e: >> +/* { dg-output "devices present but 'omp requires unified_address, unif= ied_shared_memory, reverse_offload' cannot be fulfilled" } */ > (The latter diagnostic later got conditionalized by 'GOMP_DEBUG=3D1'.) > OK to push the attached "Enhance 'libgomp.c-c++-common/requires-4.c', > 'libgomp.c-c++-common/requires-5.c' testing"? ... > libgomp/ > * testsuite/libgomp.c-c++-common/requires-4.c: Enhance testing. > * testsuite/libgomp.c-c++-common/requires-5.c: Likewise. ... > --- a/libgomp/testsuite/libgomp.c-c++-common/requires-4.c > +++ b/libgomp/testsuite/libgomp.c-c++-common/requires-4.c > @@ -1,22 +1,29 @@ > -/* { dg-do link { target offloading_enabled } } */ > /* { dg-additional-options "-flto" } */ > /* { dg-additional-sources requires-4-aux.c } */ > > -/* Check diagnostic by device-compiler's or host compiler's lto1. > +/* Check no diagnostic by device-compiler's or host compiler's lto1. I note that without ENABLE_OFFLOADING that there is never any lto1 diagnostic. However, given that no diagnostic is expected, it also works for "! offloading_enabled". Thus, the change fine. > Other file uses: 'requires reverse_offload', but that's inactive as > there are no declare target directives, device constructs nor device= routines */ > > +/* For actual offload execution, prints the following (only) if GOMP_DEB= UG=3D1: > + "devices present but 'omp requires unified_address, unified_shared_me= mory, reverse_offload' cannot be fulfilled" > + and does host-fallback execution. */ The latter is only true when also device code is produced =E2=80=93 and a d= evice is available for that/those device types. I think that's what you imply by "For actual offload execution", but it is a bit hidden. Maybe s/For actual offload execution, prints/It may print/ is clearer? In principle, it would be nice if we could test for the output, but currently setting an env var for remote execution does not work, yet. Cf. https://gcc.gnu.org/pipermail/gcc-patches/2022-July/597773.html - When set, we could use offload_target_nvptx etc. (..._amdgcn, ..._any) to test =E2=80=93 as this guarantees that it is compiled for that device + = the device is available. > + > #pragma omp requires unified_address,unified_shared_memory > > -int a[10]; > +int a[10] =3D { 0 }; > extern void foo (void); > > int > main (void) > { > - #pragma omp target > + #pragma omp target map(to: a) Hmm, I wonder whether I like it or not. Without, there is an implicit "map(tofrom:a)". On the other hand, OpenMP permits that =E2=80=93 even with unified-shared memory =E2=80=93 the implementation my copy the data to the device. (For instance, to permit faster access to "a".) Thus, ... > + for (int i =3D 0; i < 10; i++) > + a[i] =3D i; > + > for (int i =3D 0; i < 10; i++) > - a[i] =3D 0; > + if (a[i] !=3D i) > + __builtin_abort (); ... this condition (back on the host) could also fail with USM. However, given that to my knowledge no USM implementation actually copies the data, I believe it is fine. (Disclaimer: I have not checked what OG12, but I guess it also does not copy it.) > --- a/libgomp/testsuite/libgomp.c-c++-common/requires-5.c > +++ b/libgomp/testsuite/libgomp.c-c++-common/requires-5.c > @@ -1,21 +1,25 @@ > -/* { dg-do run { target { offload_target_nvptx || offload_target_amdgcn = } } } */ > /* { dg-additional-sources requires-5-aux.c } */ > > +/* For actual offload execution, prints the following (only) if GOMP_DEB= UG=3D1: > + "devices present but 'omp requires unified_address, unified_shared_me= mory, reverse_offload' cannot be fulfilled" > + and does host-fallback execution. */ > + This wording is correct with the now-removed check =E2=80=93 but if you rem= ove the offload_target..., it only "might" print it, depending, well, on the conditions set by offload_target... > #pragma omp requires unified_shared_memory, unified_address, reverse_of= fload > > -int a[10]; > +int a[10] =3D { 0 }; > extern void foo (void); > > int > main (void) > { > - #pragma omp target > + #pragma omp target map(to: a) > + for (int i =3D 0; i < 10; i++) > + a[i] =3D i; > + > for (int i =3D 0; i < 10; i++) > - a[i] =3D 0; > + if (a[i] !=3D i) > + __builtin_abort (); > > foo (); > return 0; > } Thus: LGTM, if you update the GOMP_DEBUG=3D... wording, either using "might" (etc.) or by being more explicit. Once we have remote setenv, we probably want to add another testcase to check for the GOMP_DEBUG=3D1, copying an existing one, adding dg-output and restricting it to target offload_target_... Tobias ----------------- Siemens Electronic Design Automation GmbH; Anschrift: Arnulfstra=C3=9Fe 201= , 80634 M=C3=BCnchen; Gesellschaft mit beschr=C3=A4nkter Haftung; Gesch=C3= =A4ftsf=C3=BChrer: Thomas Heurung, Frank Th=C3=BCrauf; Sitz der Gesellschaf= t: M=C3=BCnchen; Registergericht M=C3=BCnchen, HRB 106955