From mboxrd@z Thu Jan 1 00:00:00 1970 Return-Path: Received: from esa1.mentor.iphmx.com (esa1.mentor.iphmx.com [68.232.129.153]) by sourceware.org (Postfix) with ESMTPS id AC3B5385696F for ; Thu, 7 Jul 2022 10:42:29 +0000 (GMT) DMARC-Filter: OpenDMARC Filter v1.4.1 sourceware.org AC3B5385696F 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,223";a="81113056" Received: from orw-gwy-02-in.mentorg.com ([192.94.38.167]) by esa1.mentor.iphmx.com with ESMTP; 07 Jul 2022 02:42:28 -0800 IronPort-SDR: s+XVY5cSKeNm1uGmtSzwMet+o0Oh8t392YbYLpoEMSy+46Kenv14XtfbTkhogskeGxz7DIVOGl uBX0qmsrqxp4D3nb2m1qcRli1EpZJ1bgp74VN35WhxHzqfcGOXXjb1kFmQtHfrp1d3eDY6pvUo cE5F21ISylDqVmdURk7brdpsGr1VXGchS45w0JK8ZObD+LP5WvUXuNP6Ng8sVJk9MEjTaaDKLe 4dyeqyAIHnkCczF5/beL+nZWnzBhd7io0j7Q7tU5kV9g48RxjS1KGv2T5WppTjigG+iNdMDd4/ sEE= From: Thomas Schwinge To: Tobias Burnus , CC: Jakub Jelinek 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) In-Reply-To: <92609eba-2c1b-99b4-c44d-99c438fabb59@codesourcery.com> 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> <92609eba-2c1b-99b4-c44d-99c438fabb59@codesourcery.com> User-Agent: Notmuch/0.29.3+94~g74c3f1b (https://notmuchmail.org) Emacs/27.1 (x86_64-pc-linux-gnu) Date: Thu, 7 Jul 2022 12:42:20 +0200 Message-ID: <87wncp5hmb.fsf@euler.schwinge.homeip.net> MIME-Version: 1.0 Content-Type: multipart/mixed; boundary="=-=-=" X-Originating-IP: [137.202.0.90] X-ClientProxiedBy: SVR-IES-MBX-07.mgc.mentorg.com (139.181.222.7) To svr-ies-mbx-10.mgc.mentorg.com (139.181.222.10) X-Spam-Status: No, score=-12.0 required=5.0 tests=BAYES_00, GIT_PATCH_0, HEADER_FROM_DIFFERENT_DOMAINS, KAM_DMARC_STATUS, KAM_SHORT, SPF_HELO_PASS, SPF_PASS, TXREP, T_SCC_BODY_TEXT_LINE autolearn=ham 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 10:42:32 -0000 --=-=-= Content-Type: text/plain; charset="utf-8" Content-Transfer-Encoding: quoted-printable Hi Tobias! On 2022-07-07T11:36:34+0200, Tobias Burnus wrote: > On 07.07.22 10:42, Thomas Schwinge wrote: >> In preparation for other changes: > ... >> On 2022-06-29T16:33:02+0200, Tobias Burnus wro= te: >>> +/* { dg-output "devices present but 'omp requires unified_address, uni= fied_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. ACK. >> Other file uses: 'requires reverse_offload', but that's inactive as >> there are no declare target directives, device constructs nor devic= e routines */ >> >> +/* For actual offload execution, prints the following (only) if GOMP_DE= BUG=3D1: >> + "devices present but 'omp requires unified_address, unified_shared_m= emory, 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= device > is available for that/those device types. I think that's what you imply > by "For actual offload execution" ACK. > but it is a bit hidden. > > Maybe s/For actual offload execution, prints/It may print/ is clearer? I've settled on: /* Depending on offload device capabilities, it may print something lik= e the following (only) if GOMP_DEBUG=3D1: "devices present but 'omp requires unified_address, unified_shared_m= emory, reverse_offload' cannot be fulfilled" and in that case does host-fallback execution. */ > 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 Right, I'm aware of that issue with remote testing, and that's why I didn't propose such output verification. (In a few other test cases, we do have 'dg-set-target-env-var GOMP_DEBUG "1"', which then at present are UNSUPPORTED for remote testing.) > 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. Use 'target offload_device_nvptx', not 'target offload_target_nvptx', etc. ;-) >> + >> #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 wi= th > unified-shared memory =E2=80=93 the implementation my copy the data to th= e > 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. Right, this is meant to describe/test the current GCC master branch behavior, where USM isn't supported, so I didn't consider that. But I agree, a source code comment should be added: As no offload devices support USM at present, we may verify host-fal= lback execution by absence of separate memory spaces. */ > (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_DE= BUG=3D1: >> + "devices present but 'omp requires unified_address, unified_shared_m= emory, 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 r= emove > 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_o= ffload >> >> -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. Used the wording as in 'libgomp.c-c++-common/requires-4.c'. Thanks for the review! Pushed to master branch commit 5647e2c3853cbd51a6536a84b8eb0eb3c210dfbf "Enhance 'libgomp.c-c++-common/requires-4.c', 'libgomp.c-c++-common/requires-5.c' testing", see attached. Gr=C3=BC=C3=9Fe Thomas ----------------- 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 --=-=-= Content-Type: text/x-diff Content-Disposition: inline; filename="0001-Enhance-libgomp.c-c-common-requires-4.c-libgomp.c-c-.patch" >From 5647e2c3853cbd51a6536a84b8eb0eb3c210dfbf Mon Sep 17 00:00:00 2001 From: Thomas Schwinge Date: Thu, 7 Jul 2022 09:45:42 +0200 Subject: [PATCH] Enhance 'libgomp.c-c++-common/requires-4.c', 'libgomp.c-c++-common/requires-5.c' testing These should compile and link and execute in all configurations; host-fallback execution, which we may actually verify. Follow-up to recent commit 683f11843974f0bdf42f79cdcbb0c2b43c7b81b0 "OpenMP: Move omp requires checks to libgomp". libgomp/ * testsuite/libgomp.c-c++-common/requires-4.c: Enhance testing. * testsuite/libgomp.c-c++-common/requires-5.c: Likewise. --- .../libgomp.c-c++-common/requires-4.c | 21 +++++++++++++----- .../libgomp.c-c++-common/requires-5.c | 22 +++++++++++++------ 2 files changed, 31 insertions(+), 12 deletions(-) diff --git a/libgomp/testsuite/libgomp.c-c++-common/requires-4.c b/libgomp/testsuite/libgomp.c-c++-common/requires-4.c index 128fdbb8463..6ed5a5f647a 100644 --- a/libgomp/testsuite/libgomp.c-c++-common/requires-4.c +++ b/libgomp/testsuite/libgomp.c-c++-common/requires-4.c @@ -1,22 +1,33 @@ -/* { 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. Other file uses: 'requires reverse_offload', but that's inactive as there are no declare target directives, device constructs nor device routines */ +/* Depending on offload device capabilities, it may print something like the + following (only) if GOMP_DEBUG=1: + "devices present but 'omp requires unified_address, unified_shared_memory, reverse_offload' cannot be fulfilled" + and in that case does host-fallback execution. + + No offload devices support USM at present, so we may verify host-fallback + execution by presence of separate memory spaces. */ + #pragma omp requires unified_address,unified_shared_memory -int a[10]; +int a[10] = { 0 }; extern void foo (void); int main (void) { - #pragma omp target + #pragma omp target map(to: a) + for (int i = 0; i < 10; i++) + a[i] = i; + for (int i = 0; i < 10; i++) - a[i] = 0; + if (a[i] != i) + __builtin_abort (); foo (); return 0; diff --git a/libgomp/testsuite/libgomp.c-c++-common/requires-5.c b/libgomp/testsuite/libgomp.c-c++-common/requires-5.c index c1e5540cfc5..7fe0c735d27 100644 --- a/libgomp/testsuite/libgomp.c-c++-common/requires-5.c +++ b/libgomp/testsuite/libgomp.c-c++-common/requires-5.c @@ -1,21 +1,29 @@ -/* { dg-do run { target { offload_target_nvptx || offload_target_amdgcn } } } */ /* { dg-additional-sources requires-5-aux.c } */ +/* Depending on offload device capabilities, it may print something like the + following (only) if GOMP_DEBUG=1: + "devices present but 'omp requires unified_address, unified_shared_memory, reverse_offload' cannot be fulfilled" + and in that case does host-fallback execution. + + As no offload devices support USM at present, we may verify host-fallback + execution by absence of separate memory spaces. */ + #pragma omp requires unified_shared_memory, unified_address, reverse_offload -int a[10]; +int a[10] = { 0 }; extern void foo (void); int main (void) { - #pragma omp target + #pragma omp target map(to: a) for (int i = 0; i < 10; i++) - a[i] = 0; + a[i] = i; + + for (int i = 0; i < 10; i++) + if (a[i] != i) + __builtin_abort (); foo (); return 0; } - -/* (Only) if GOMP_DEBUG=1, should print at runtime the following: - "devices present but 'omp requires unified_address, unified_shared_memory, reverse_offload' cannot be fulfilled" */ -- 2.35.1 --=-=-=--