From patchwork Fri Apr 3 08:29:27 2020 Content-Type: text/plain; charset="utf-8" MIME-Version: 1.0 Content-Transfer-Encoding: 8bit X-Patchwork-Submitter: Thomas Schwinge X-Patchwork-Id: 1265955 Return-Path: X-Original-To: incoming@patchwork.ozlabs.org Delivered-To: patchwork-incoming@bilbo.ozlabs.org Authentication-Results: ozlabs.org; spf=pass (sender SPF authorized) smtp.mailfrom=gcc.gnu.org (client-ip=8.43.85.97; helo=sourceware.org; envelope-from=gcc-patches-bounces@gcc.gnu.org; receiver=) Authentication-Results: ozlabs.org; dmarc=none (p=none dis=none) header.from=codesourcery.com Received: from sourceware.org (server2.sourceware.org [8.43.85.97]) (using TLSv1.3 with cipher TLS_AES_256_GCM_SHA384 (256/256 bits) key-exchange X25519 server-signature RSA-PSS (4096 bits) server-digest SHA256) (No client certificate requested) by ozlabs.org (Postfix) with ESMTPS id 48ttQh4vSxz9sP7 for ; Fri, 3 Apr 2020 19:29:47 +1100 (AEDT) Received: from server2.sourceware.org (localhost [IPv6:::1]) by sourceware.org (Postfix) with ESMTP id B2953385BF92; Fri, 3 Apr 2020 08:29:42 +0000 (GMT) X-Original-To: gcc-patches@gcc.gnu.org Delivered-To: gcc-patches@gcc.gnu.org Received: from esa2.mentor.iphmx.com (esa2.mentor.iphmx.com [68.232.141.98]) by sourceware.org (Postfix) with ESMTPS id E6FCD385B833 for ; Fri, 3 Apr 2020 08:29:38 +0000 (GMT) DMARC-Filter: OpenDMARC Filter v1.3.2 sourceware.org E6FCD385B833 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: dZobI8n33M+/FdSipJwj+UJxF6buxfUs6cSjw0vacNm7jTA8mADTVwZ9XIZrTqkuxnIswY9xzR wjq4Si3Wg5BIJMVAkROJo0gyZZCRHf1txVGvkzmO2JRVUV085hQFW5HCkpPNJSr9M49iyI+KOP uNAEnMfvYLpSAb+5qZgS+740iOkHykz9+h4yORPi4FtZAgPpz80g6eYb5U4/b580e8WKqUNOVi 9379IbuB3Oybggsqfce/YYtpedlH4o74MdcRjzfnAgYSRpyk6JrFcL6B1vw+nHIiL7wcDYBs9e bFY= X-IronPort-AV: E=Sophos; i="5.72,339,1580803200"; d="scan'208,223"; a="47324683" Received: from orw-gwy-02-in.mentorg.com ([192.94.38.167]) by esa2.mentor.iphmx.com with ESMTP; 03 Apr 2020 00:29:37 -0800 IronPort-SDR: fW5LbMwOm9FjNMz5AJBm87Gf0kZJfkeJFAqNWeSq9F4LmsfkCmdAYKwNp/E2nhcSlBl+lDsh9R OMjexJMcepJQ7I5RLnXFp3HECch675jtb8uPW+jAbQiL/rDwlVxPkMwsxbNL4spf+Xq9JtdYV/ lMVPNbzE7RLya3Zgjmhv5RP+VagStG6OUTtlrpVWDd2VkhFrpp6lVQ7/WTZqqubydDcvlBgP5g +J4NaEcNI9vwRhG4iD+Dkp6kNEz/rBC3nWv7P8yIhxMb/1CzVGmZwwlDvkpw67S8biqfihi946 IAg= From: Thomas Schwinge To: Richard Biener , Tom de Vries , Subject: Revert "[nvptx, libgomp] Update pr85381-{2, 4}.c test-cases" [PR89713, PR94392] (was: [PATCH][RFC] c/94392 - only enable -ffinite-loops for C++) In-Reply-To: References: <0f331251-85c6-2cfe-e7aa-54bc7966ce57@redhat.com> User-Agent: Notmuch/0.29.1+93~g67ed7df (https://notmuchmail.org) Emacs/26.1 (x86_64-pc-linux-gnu) Date: Fri, 3 Apr 2020 10:29:27 +0200 Message-ID: <877dyxnfa0.fsf@euler.schwinge.homeip.net> MIME-Version: 1.0 X-Spam-Status: No, score=-21.0 required=5.0 tests=BAYES_00, GIT_PATCH_0, GIT_PATCH_1, GIT_PATCH_2, GIT_PATCH_3, HEADER_FROM_DIFFERENT_DOMAINS, KAM_DMARC_STATUS, KAM_LOTSOFHASH, SPF_HELO_NONE, 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: , Cc: "Joseph S. Myers" Errors-To: gcc-patches-bounces@gcc.gnu.org Sender: "Gcc-patches" Hi! On 2020-04-02T11:12:48+0200, Richard Biener wrote: > On Wed, 1 Apr 2020, Jason Merrill wrote: > >> On 4/1/20 9:36 AM, Richard Biener wrote: >> > This does away with enabling -ffinite-loops at -O2+ for all languages >> > and instead enables it selectively for C++ only. > I'm retesting the following [...] ..., which got pushed in commit 75efe9cb1f8938a713ce540dc3b27bc2afcd3fae "c/94392 - only enable -ffinite-loops for C++". I pushed the attached in commit 4f6a0888de52a2e523a6fd4235fe7f8193819c3b 'Revert "[nvptx, libgomp] Update pr85381-{2,4}.c test-cases" [PR89713, PR94392]'. As can be observed in two nvptx offloading test cases regressing, 'apparently now again "empty oacc loops are" no longer "removed before expand"' (quoting myself from the commit log, adapting Tom's commit log snippet from the reverted commit). It's not obvious to me how the "finite loop" property discussed/changed in Richard's commit 75efe9cb1f8938a713ce540dc3b27bc2afcd3fae "c/94392 - only enable -ffinite-loops for C++" relates to the previously observed optimization of removing "empty oacc loops [...] before expand" (after PR89713 commit c29c92c789d93848cc1c929838771bfc68cb272c "PR tree-optimization/89713 - Assume loop with an exit is finite"), but examining that in detail is for another day. Grüße Thomas ----------------- Mentor Graphics (Deutschland) GmbH, Arnulfstraße 201, 80634 München / Germany Registergericht München HRB 106955, Geschäftsführer: Thomas Heurung, Alexander Walter From 4f6a0888de52a2e523a6fd4235fe7f8193819c3b Mon Sep 17 00:00:00 2001 From: Thomas Schwinge Date: Fri, 3 Apr 2020 10:07:16 +0200 Subject: [PATCH] Revert "[nvptx, libgomp] Update pr85381-{2,4}.c test-cases" [PR89713, PR94392] In response to PR94392 commit 75efe9cb1f8938a713ce540dc3b27bc2afcd3fae "c/94392 - only enable -ffinite-loops for C++", this reverts PR89713 commit 00908992f2a78f213d227aea8dbab014a1361df0, as apparently now again "empty oacc loops are" no longer "removed before expand". libgomp/ PR tree-optimization/89713 PR c/94392 * testsuite/libgomp.oacc-c-c++-common/pr85381-2.c: Again expect 'bar.sync'. * testsuite/libgomp.oacc-c-c++-common/pr85381-4.c: Likewise. --- libgomp/ChangeLog | 8 ++++++++ .../libgomp.oacc-c-c++-common/pr85381-2.c | 20 ++++++++++++++++++- .../libgomp.oacc-c-c++-common/pr85381-4.c | 5 ++++- 3 files changed, 31 insertions(+), 2 deletions(-) diff --git a/libgomp/ChangeLog b/libgomp/ChangeLog index 6c437930b02f..3716f559aa1c 100644 --- a/libgomp/ChangeLog +++ b/libgomp/ChangeLog @@ -1,3 +1,11 @@ +2020-04-03 Thomas Schwinge + + PR tree-optimization/89713 + PR c/94392 + * testsuite/libgomp.oacc-c-c++-common/pr85381-2.c: Again expect + 'bar.sync'. + * testsuite/libgomp.oacc-c-c++-common/pr85381-4.c: Likewise. + 2020-03-31 Tobias Burnus * target.c (GOMP_target_enter_exit_data): Handle PSET/MAP_POINTER. 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 2cb5b95949de..6570c64afff5 100644 --- a/libgomp/testsuite/libgomp.oacc-c-c++-common/pr85381-2.c +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/pr85381-2.c @@ -15,4 +15,22 @@ main (void) return 0; } -/* { dg-final { scan-assembler-times "bar.sync" 0 } } */ +/* Todo: Boths bar.syncs can be removed. + Atm we generate this dead code inbetween forked and joining: + + mov.u32 %r28, %ntid.y; + mov.u32 %r29, %tid.y; + add.u32 %r30, %r29, %r29; + setp.gt.s32 %r31, %r30, 19; + @%r31 bra $L2; + add.u32 %r25, %r28, %r28; + mov.u32 %r24, %r30; + $L3: + add.u32 %r24, %r24, %r25; + setp.le.s32 %r33, %r24, 19; + @%r33 bra $L3; + $L2: + + 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 } } */ 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 e8a433ffc0a5..d955d79718df 100644 --- a/libgomp/testsuite/libgomp.oacc-c-c++-common/pr85381-4.c +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/pr85381-4.c @@ -21,4 +21,7 @@ main (void) return 0; } -/* { dg-final { scan-assembler-times "bar.sync" 0 } } */ +/* 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 } } */ -- 2.25.1