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 32902385702E for ; Tue, 30 Jun 2020 16:07:39 +0000 (GMT) DMARC-Filter: OpenDMARC Filter v1.3.2 sourceware.org 32902385702E Authentication-Results: sourceware.org; dmarc=none (p=none dis=none) header.from=codesourcery.com Authentication-Results: sourceware.org; spf=pass smtp.mailfrom=Thomas_Schwinge@mentor.com IronPort-SDR: LRopG+HdJpQWnv1QivORdb6qyaXmK6TfZiRtqki1T4RnNvGit86oHIKQNKej0GesYH05MFIXWe kgSAgVp/Dif8HDftInEQ6p/TD4eClF52Jw5Cpwvq5N3iV+kF0sUl55Y6qqPhwKh13jU797h061 xfHjLzf8ZEJGPokBv+4vN9L5FzMXgbzHyZlKTI48pTZuiJ9nGt+oGGofJB5db4Wy5Klg5OKWQZ 1T8TXy+XAZg6Pq2Ip1JJulscGzS7Aux32Cs0iqpyDnAIirDcdzm10BDW/0MVUDY/Oma0OdLB3e HdA= X-IronPort-AV: E=Sophos;i="5.75,297,1589270400"; d="scan'208,223";a="52599362" Received: from orw-gwy-01-in.mentorg.com ([192.94.38.165]) by esa1.mentor.iphmx.com with ESMTP; 30 Jun 2020 08:07:37 -0800 IronPort-SDR: g+5NpbuvkpmEf0SOR6Op0xA+/1gZFo/McAqdNeeYm/J89sut0kKKhnGbuMGDUZkAFYs5BRVMVJ PXUIjwpDm09xKpxH8KDdsGKCNTOD3i4nhfL7bXH+ZVKz0gu/ksBrWVRK8abTBGX7SQ7t3Fv4DX KvHg5xcg/ZvbWZbU1Y1iqOpzyQPXsQjVrn87AKnMMsw32TxG8GV+lQm6cAMGP+FpeHbElKIH9n uRKtkzed03P4C75hO1QrSbMCYMIqRR7rwpfm7OMNYk9sGT3/Xz+7aUNH0DCEfpH6oG+D6kQtbs e3U= From: Thomas Schwinge To: Alexandre Oliva , , Tom de Vries CC: Jakub Jelinek , , Richard Biener , , Tobias Burnus Subject: Re: drop -aux{dir,base}, revamp -dump{dir,base} In-Reply-To: References: <874krkqte6.fsf@euler.schwinge.homeip.net> User-Agent: Notmuch/0.29.1+93~g67ed7df (https://notmuchmail.org) Emacs/26.3 (x86_64-pc-linux-gnu) Date: Tue, 30 Jun 2020 18:07:26 +0200 Message-ID: <87eepwzge9.fsf@euler.schwinge.homeip.net> MIME-Version: 1.0 Content-Type: multipart/mixed; boundary="=-=-=" X-Spam-Status: No, score=-10.7 required=5.0 tests=BAYES_00, GIT_PATCH_0, HEADER_FROM_DIFFERENT_DOMAINS, KAM_DMARC_STATUS, SPF_HELO_PASS, SPF_PASS, TXREP autolearn=ham autolearn_force=no version=3.4.2 X-Spam-Checker-Version: SpamAssassin 3.4.2 (2018-09-13) 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: Tue, 30 Jun 2020 16:07:41 -0000 --=-=-= Content-Type: text/plain; charset="utf-8" Content-Transfer-Encoding: quoted-printable Hi! Many thanks, Alexandra and Tobias for working this out together! On 2020-06-23T06:50:26-0300, Alexandre Oliva wrote: > On Jun 9, 2020, Thomas Schwinge wrote: > >> Previously, for '-foffload=3Dnvptx-none -foffload=3D-fdump-rtl-mach >> -save-temps -o ./nvptx-merged-loop.exe', GCC produced the expected >> 'nvptx-merged-loop.o.307r.mach'. > > I believe the patch I've just installed fixes the UNRESOLVED results > caused by not finding dump files. Yes, confirmed, thanks! A few small issues remain, for those I'll respond elsewhere in this thread. >> Consider 'libgomp.oacc-c-c++-common/pr85381-2.c': > >> /* { dg-additional-options "-save-temps" } */ > >> /* { dg-final { scan-assembler-times "bar.sync" 2 } } */ > >> This expects to scan the PTX offloading compilation assembler code (not >> host code!), expecting that nvptx offloading code assembly is produced >> after the host code, and thus overwrites the latter file. (Yes, that's >> certainly ugly/fragile...) > > I'm afraid this will need further adjusting in the testsuite, as we'll > store the nvptx asm saved aux output in a separate file. > scan-assembler-times will no longer work for this purpose, we'll need > something that knows how to find the offloaded asm. So, that's (now?) easy enough to repair. I've pushed "[testsuite] Replace fragile 'scan-assembler' with 'scan-offload-rtl' in 'libgomp.oacc-c-c++-common/pr85381*.c'", see attached. Gr=C3=BC=C3=9Fe Thomas ----------------- Mentor Graphics (Deutschland) GmbH, Arnulfstra=C3=9Fe 201, 80634 M=C3=BCnch= en / Germany Registergericht M=C3=BCnchen HRB 106955, Gesch=C3=A4ftsf=C3=BChrer: Thomas = Heurung, Alexander Walter --=-=-= Content-Type: text/x-diff Content-Disposition: inline; filename="0001-testsuite-Replace-fragile-scan-assembler-with-scan-o.patch" >From 8a8efad09811b5c08cfd9d469c6f1b6ba0c848f1 Mon Sep 17 00:00:00 2001 From: Thomas Schwinge Date: Tue, 30 Jun 2020 05:24:17 +0200 Subject: [PATCH] [testsuite] Replace fragile 'scan-assembler' with 'scan-offload-rtl' in 'libgomp.oacc-c-c++-common/pr85381*.c' These test cases use directives similar to: /* { dg-additional-options "-save-temps" } */ /* { dg-final { scan-assembler-times "bar.sync" 2 } } */ This expects to scan the PTX offloading compilation assembler code (not host code!), expecting that nvptx offloading code assembly is produced after the host code, and thus overwrites the latter file. (Yes, that's certainly ugly/fragile...) ..., and this broke with recent commit 1dedc12d186a110854537e1279b4e6c29f2df35a "revamp dump and aux output names" plus fix-up commit commit efc16503ca10bc0e934e0bace5777500e4dc757a "handle dumpbase in offloading, adjust testsuite" (short summary: file names changed), so let's finally make that robust. libgomp/ * testsuite/libgomp.oacc-c-c++-common/pr85381-2.c: Replace fragile 'scan-assembler' with 'scan-offload-rtl'. * testsuite/libgomp.oacc-c-c++-common/pr85381-3.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/pr85381-4.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/pr85381-5.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/pr85381.c: Likewise. --- libgomp/testsuite/libgomp.oacc-c-c++-common/pr85381-2.c | 4 ++-- libgomp/testsuite/libgomp.oacc-c-c++-common/pr85381-3.c | 4 ++-- libgomp/testsuite/libgomp.oacc-c-c++-common/pr85381-4.c | 4 ++-- libgomp/testsuite/libgomp.oacc-c-c++-common/pr85381-5.c | 4 ++-- libgomp/testsuite/libgomp.oacc-c-c++-common/pr85381.c | 4 ++-- 5 files changed, 10 insertions(+), 10 deletions(-) diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/pr85381-2.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/pr85381-2.c index 6570c64afff5..84b9c01443e5 100644 --- a/libgomp/testsuite/libgomp.oacc-c-c++-common/pr85381-2.c +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/pr85381-2.c @@ -1,6 +1,6 @@ -/* { dg-additional-options "-save-temps" } */ /* { dg-do run { target openacc_nvidia_accel_selected } } { dg-skip-if "" { *-*-* } { "*" } { "-O2" } } */ +/* { dg-additional-options "-foffload=-fdump-rtl-mach" } */ int main (void) @@ -33,4 +33,4 @@ main (void) so the loop is not recognized as empty loop (which we detect by seeing if joining immediately follows forked). */ -/* { dg-final { scan-assembler-times "bar.sync" 2 } } */ +/* { dg-final { scan-offload-rtl-dump-times "nvptx_barsync" 2 "mach" } } */ diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/pr85381-3.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/pr85381-3.c index c5d1c5add68e..cddbf2719067 100644 --- a/libgomp/testsuite/libgomp.oacc-c-c++-common/pr85381-3.c +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/pr85381-3.c @@ -1,6 +1,6 @@ -/* { dg-additional-options "-save-temps -w" } */ /* { dg-do run { target openacc_nvidia_accel_selected } } { dg-skip-if "" { *-*-* } { "*" } { "-O2" } } */ +/* { dg-additional-options "-foffload=-fdump-rtl-mach" } */ int a; #pragma acc declare create(a) @@ -32,4 +32,4 @@ main (void) return 0; } -/* { dg-final { scan-assembler-not "bar.sync" } } */ +/* { dg-final { scan-offload-rtl-dump-not "nvptx_barsync" "mach" } } */ diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/pr85381-4.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/pr85381-4.c index d955d79718df..e1679444172c 100644 --- a/libgomp/testsuite/libgomp.oacc-c-c++-common/pr85381-4.c +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/pr85381-4.c @@ -1,6 +1,6 @@ -/* { dg-additional-options "-save-temps -w" } */ /* { dg-do run { target openacc_nvidia_accel_selected } } { dg-skip-if "" { *-*-* } { "*" } { "-O2" } } */ +/* { dg-additional-options "-foffload=-fdump-rtl-mach" } */ #define n 1024 @@ -24,4 +24,4 @@ main (void) /* Atm, %ntid.y is broadcast from one loop to the next, so there are 2 bar.syncs for that (the other two are there for the same reason as in pr85381-2.c). Todo: Recompute %ntid.y instead of broadcasting it. */ -/* { dg-final { scan-assembler-times "bar.sync" 4 } } */ +/* { dg-final { scan-offload-rtl-dump-times "nvptx_barsync" 4 "mach" } } */ diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/pr85381-5.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/pr85381-5.c index 61e7e48f0c93..26ca5093c47d 100644 --- a/libgomp/testsuite/libgomp.oacc-c-c++-common/pr85381-5.c +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/pr85381-5.c @@ -1,6 +1,6 @@ -/* { dg-additional-options "-save-temps" } */ /* { dg-do run { target openacc_nvidia_accel_selected } } { dg-skip-if "" { *-*-* } { "*" } { "-O2" } } */ +/* { dg-additional-options "-foffload=-fdump-rtl-mach" } */ #define n 1024 @@ -21,4 +21,4 @@ main (void) return 0; } -/* { dg-final { scan-assembler-not "bar.sync" } } */ +/* { dg-final { scan-offload-rtl-dump-not "nvptx_barsync" "mach" } } */ diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/pr85381.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/pr85381.c index 2864dfcf3cb1..eda87743625b 100644 --- a/libgomp/testsuite/libgomp.oacc-c-c++-common/pr85381.c +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/pr85381.c @@ -1,6 +1,6 @@ -/* { dg-additional-options "-save-temps" } */ /* { dg-do run { target openacc_nvidia_accel_selected } } { dg-skip-if "" { *-*-* } { "*" } { "-O2" } } */ +/* { dg-additional-options "-foffload=-fdump-rtl-mach" } */ int main (void) @@ -15,4 +15,4 @@ main (void) return 0; } -/* { dg-final { scan-assembler-not "bar.sync" } } */ +/* { dg-final { scan-offload-rtl-dump-not "nvptx_barsync" "mach" } } */ -- 2.27.0 --=-=-=--