From mboxrd@z Thu Jan 1 00:00:00 1970 Return-Path: Received: (qmail 5787 invoked by alias); 6 Apr 2018 13:49:09 -0000 Mailing-List: contact gcc-patches-help@gcc.gnu.org; run by ezmlm Precedence: bulk List-Id: List-Archive: List-Post: List-Help: Sender: gcc-patches-owner@gcc.gnu.org Received: (qmail 3388 invoked by uid 89); 6 Apr 2018 13:49:02 -0000 Authentication-Results: sourceware.org; auth=none X-Virus-Found: No X-Spam-SWARE-Status: No, score=-24.7 required=5.0 tests=AWL,BAYES_00,GIT_PATCH_0,GIT_PATCH_1,GIT_PATCH_2,GIT_PATCH_3,KAM_NUMSUBJECT,RCVD_IN_DNSWL_NONE,SPF_PASS,URIBL_RED autolearn=ham version=3.3.2 spammy=tiled, splits X-HELO: relay1.mentorg.com Received: from relay1.mentorg.com (HELO relay1.mentorg.com) (192.94.38.131) by sourceware.org (qpsmtpd/0.93/v0.84-503-g423c35a) with ESMTP; Fri, 06 Apr 2018 13:48:57 +0000 Received: from svr-orw-mbx-04.mgc.mentorg.com ([147.34.90.204]) by relay1.mentorg.com with esmtps (TLSv1.2:ECDHE-RSA-AES256-SHA384:256) id 1f4Rjb-0003jP-G4 from Cesar_Philippidis@mentor.com for gcc-patches@gcc.gnu.org; Fri, 06 Apr 2018 06:48:55 -0700 Received: from [127.0.0.1] (147.34.91.1) by SVR-ORW-MBX-04.mgc.mentorg.com (147.34.90.204) with Microsoft SMTP Server (TLS) id 15.0.1320.4; Fri, 6 Apr 2018 06:48:53 -0700 From: Cesar Philippidis Subject: [PATCH] Handle empty infinite loops in OpenACC for PR84955 To: "gcc-patches@gcc.gnu.org" Message-ID: Date: Fri, 06 Apr 2018 13:49:00 -0000 User-Agent: Mozilla/5.0 (X11; Linux x86_64; rv:52.0) Gecko/20100101 Thunderbird/52.7.0 MIME-Version: 1.0 Content-Type: multipart/mixed; boundary="------------A91886DA6924B6DB391BDAA4" X-ClientProxiedBy: svr-orw-mbx-04.mgc.mentorg.com (147.34.90.204) To SVR-ORW-MBX-04.mgc.mentorg.com (147.34.90.204) X-SW-Source: 2018-04/txt/msg00316.txt.bz2 --------------A91886DA6924B6DB391BDAA4 Content-Type: text/plain; charset="utf-8" Content-Transfer-Encoding: 7bit Content-length: 1775 As shown in PR84955, GCC has problems empty infinite loops of the form #pragma acc loop tile (2, 2) for (...) for (...) { for (...) /* do nothing */ ; } I also observed that it generates bad code for the following loop #pragma acc loop for (...) { for (...) /* do nothing */ ; } There are two problems here. First, for the tiled loop, as expand_oacc_for splits the first basic block of the loop body to add a branch to perform the tiling control logic, it does not consider fact that the body may not have an exit when it contains an infinite loop. This situation can occur when region->cont is NULL. The fix here was to add a dummy false edge to exit_bb. The second problem involves fixing up the CFG. Unlike OpenMP which defers a lot of the scheduling control to libgomp, OpenACC makes heavy use of calls to internal functions. The problem here is that the fixup_cfg pass was largely ignoring calls to internal functions when it decides to set PROP_cleanup_cfg. When the CFG isn't cleaned up, the LTO streamer will report that there are fewer loops than there actually are and that causes an ICE because the empty infinite loop is never accounted for. This patch resolves this issue by relaxing the fixup_cfg pass to treat all functions calls, including those to internal functions, the same. An alternative approach to resolve this issue would have been to teach ipa_write_summaries to check if the loop structure needs to be fixed up and call cleanup_cfg as necessary. But I wanted to keep the OMP and OACC code paths similar, so I took the former approach. I regression tested this patch on x86_64-linux using nvptx offloading. Is this patch OK for trunk and GCC 7 (and probably GCC 6). Thanks, Cesar --------------A91886DA6924B6DB391BDAA4 Content-Type: text/x-patch; name="trunk-pr84955.diff" Content-Transfer-Encoding: 7bit Content-Disposition: attachment; filename="trunk-pr84955.diff" Content-length: 3218 Fix PR84955 2018-04-06 Cesar Philippidis PR middle-end/84955 gcc/ * cfgloop.c (flow_loops_find): Add assert. * omp-expand.c (expand_oacc_for): Add dummy false branch for tiled basic blocks without omp continue statements. * tree-cfg.c (execute_fixup_cfg): Handle calls to internal functions like regular functions. libgomp/ * testsuite/libgomp.oacc-c-c++-common/pr84955.c: New test. * testsuite/libgomp.oacc-fortran/pr84955.f90: New test. diff --git a/gcc/cfgloop.c b/gcc/cfgloop.c index 8af793c6015..6e68639452c 100644 --- a/gcc/cfgloop.c +++ b/gcc/cfgloop.c @@ -462,6 +462,9 @@ flow_loops_find (struct loops *loops) { struct loop *loop; + if (!from_scratch) + gcc_assert (header->loop_father != NULL); + /* The current active loop tree has valid loop-fathers for header blocks. */ if (!from_scratch diff --git a/gcc/omp-expand.c b/gcc/omp-expand.c index bb204906ea6..1c7b68fbd8c 100644 --- a/gcc/omp-expand.c +++ b/gcc/omp-expand.c @@ -5439,6 +5439,13 @@ expand_oacc_for (struct omp_region *region, struct omp_for_data *fd) split->flags ^= EDGE_FALLTHRU | EDGE_TRUE_VALUE; + /* Add a dummy exit for the tiled block when cont_bb is missing. */ + if (cont_bb == NULL) + { + edge e = make_edge (body_bb, exit_bb, EDGE_FALSE_VALUE); + e->probability = profile_probability::even (); + } + /* Initialize the user's loop vars. */ gsi = gsi_start_bb (elem_body_bb); expand_oacc_collapse_vars (fd, true, &gsi, counts, e_offset); diff --git a/gcc/tree-cfg.c b/gcc/tree-cfg.c index 9485f73f341..cb676d78128 100644 --- a/gcc/tree-cfg.c +++ b/gcc/tree-cfg.c @@ -9586,10 +9586,7 @@ execute_fixup_cfg (void) for (gsi = gsi_start_bb (bb); !gsi_end_p (gsi);) { gimple *stmt = gsi_stmt (gsi); - tree decl = is_gimple_call (stmt) - ? gimple_call_fndecl (stmt) - : NULL; - if (decl) + if (is_gimple_call (stmt)) { int flags = gimple_call_flags (stmt); if (flags & (ECF_CONST | ECF_PURE | ECF_LOOPING_CONST_OR_PURE)) diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/pr84955.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/pr84955.c new file mode 100644 index 00000000000..5910b57b68d --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/pr84955.c @@ -0,0 +1,20 @@ +/* { dg-do compile } */ + +int +main () +{ + int i, j; + +#pragma acc parallel loop tile(2,3) + for (i = 1; i < 10; i++) + for (j = 1; j < 10; j++) + for (;;) + ; + +#pragma acc parallel loop + for (i = 1; i < 10; i++) + for (;;) + ; + + return i + j; +} diff --git a/libgomp/testsuite/libgomp.oacc-fortran/pr84955.f90 b/libgomp/testsuite/libgomp.oacc-fortran/pr84955.f90 new file mode 100644 index 00000000000..878d8a89f41 --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-fortran/pr84955.f90 @@ -0,0 +1,20 @@ +! { dg-do compile } + +subroutine s + integer :: i, j + !$acc parallel loop tile(2,3) + do i = 1, 10 + do j = 1, 10 + do + end do + end do + end do + !$acc end parallel loop + + !$acc parallel loop + do i = 1, 10 + do + end do + end do + !$acc end parallel loop +end subroutine s --------------A91886DA6924B6DB391BDAA4--