From mboxrd@z Thu Jan 1 00:00:00 1970 Return-Path: Received: by sourceware.org (Postfix, from userid 1643) id 23E92385381B; Fri, 21 Oct 2022 09:29:10 +0000 (GMT) DKIM-Filter: OpenDKIM Filter v2.11.0 sourceware.org 23E92385381B DKIM-Signature: v=1; a=rsa-sha256; c=relaxed/relaxed; d=gcc.gnu.org; s=default; t=1666344550; bh=UkvRvD/dWrYqN7bnNQtbdtCheIlblHGu1kQqoZdDG+U=; h=From:To:Subject:Date:From; b=Acop1UPctaZZPjV5ssY8MKPTxWRoKt3RKsTzMicR8gzcaFu2lyJ7UTNOYUVl4YH1A zmgs5eP0UlyYUxxC1izG8o+KmzCn8RMne4lrJ7vrUFDIynEvYe4h8IU+jUr60MchQK a+ydzv2s+t9/sho3rgelXl/60wA3ULyZT5uLJRaY= MIME-Version: 1.0 Content-Transfer-Encoding: 7bit Content-Type: text/plain; charset="utf-8" From: Thomas Schwinge To: gcc-cvs@gcc.gnu.org Subject: [gcc r13-3434] Restore 'libgomp.oacc-c-c++-common/nvptx-sese-1.c' SESE regions checking [PR107195, PR107344] X-Act-Checkin: gcc X-Git-Author: Thomas Schwinge X-Git-Refname: refs/heads/master X-Git-Oldrev: 436c4a210e8ff936063579dab3641632c06c8a7d X-Git-Newrev: a9de836c2b22f878cff592b96e11c1b95d4d36ee Message-Id: <20221021092910.23E92385381B@sourceware.org> Date: Fri, 21 Oct 2022 09:29:10 +0000 (GMT) List-Id: https://gcc.gnu.org/g:a9de836c2b22f878cff592b96e11c1b95d4d36ee commit r13-3434-ga9de836c2b22f878cff592b96e11c1b95d4d36ee Author: Thomas Schwinge Date: Sun Oct 16 00:07:20 2022 +0200 Restore 'libgomp.oacc-c-c++-common/nvptx-sese-1.c' SESE regions checking [PR107195, PR107344] That is, adjust for optimization introduced with recent commit r13-3217-gc4d15dddf6b9eacb36f535807ad2ee364af46e04 "[PR107195] Set range to zero when nonzero mask is 0", where GCC now understands that after 'r *= 2;', 'r & 1' will never hold here, and thus transforms/optimizes/"disturbs" the original code such that GCC/nvptx's later "Neuter whole SESE regions" optimization no longer is applicable to it: UNSUPPORTED: libgomp.oacc-c/../libgomp.oacc-c-c++-common/nvptx-sese-1.c -DACC_DEVICE_TYPE_nvidia=1 -DACC_MEM_SHARED=0 -foffload=nvptx-none -O0 PASS: libgomp.oacc-c/../libgomp.oacc-c-c++-common/nvptx-sese-1.c -DACC_DEVICE_TYPE_nvidia=1 -DACC_MEM_SHARED=0 -foffload=nvptx-none -O2 (test for excess errors) PASS: libgomp.oacc-c/../libgomp.oacc-c-c++-common/nvptx-sese-1.c -DACC_DEVICE_TYPE_nvidia=1 -DACC_MEM_SHARED=0 -foffload=nvptx-none -O2 execution test [-PASS:-]{+FAIL:+} libgomp.oacc-c/../libgomp.oacc-c-c++-common/nvptx-sese-1.c -DACC_DEVICE_TYPE_nvidia=1 -DACC_MEM_SHARED=0 -foffload=nvptx-none -O2 scan-nvptx-none-offload-rtl-dump mach "SESE regions:.* [0-9]+{[0-9]+->[0-9]+(\\.[0-9]+)+}" Same for C++. It's unclear to me if this is an actual "problem", which optimization is "more important", so I've filed PR107344 "GCC/nvptx SESE region optimization" to capture this question, and here restore what we intend to be testing (to my understanding) in 'libgomp.oacc-c-c++-common/nvptx-sese-1.c'. PR tree-optimization/107195 PR target/107344 libgomp/ * testsuite/libgomp.oacc-c-c++-common/nvptx-sese-1.c: Restore SESE regions checking. Diff: --- libgomp/testsuite/libgomp.oacc-c-c++-common/nvptx-sese-1.c | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/nvptx-sese-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/nvptx-sese-1.c index 9583265c775..6507e1affb0 100644 --- a/libgomp/testsuite/libgomp.oacc-c-c++-common/nvptx-sese-1.c +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/nvptx-sese-1.c @@ -22,7 +22,7 @@ int main () if (foo (r)) r *= 2; - if (r & 1) /* to here. */ + if (r & 8) /* to here. */ #pragma acc loop vector reduction (+:r) for (int i = 00; i < 40; i++) r += i;