From mboxrd@z Thu Jan 1 00:00:00 1970 Return-Path: Received: (qmail 120179 invoked by alias); 13 Oct 2015 16:00:30 -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 120158 invoked by uid 89); 13 Oct 2015 16:00:27 -0000 Authentication-Results: sourceware.org; auth=none X-Virus-Found: No X-Spam-SWARE-Status: No, score=-2.0 required=5.0 tests=AWL,BAYES_00,SPF_PASS,T_RP_MATCHES_RCVD autolearn=ham version=3.3.2 X-HELO: fencepost.gnu.org Received: from fencepost.gnu.org (HELO fencepost.gnu.org) (208.118.235.10) by sourceware.org (qpsmtpd/0.93/v0.84-503-g423c35a) with (AES128-SHA encrypted) ESMTPS; Tue, 13 Oct 2015 16:00:21 +0000 Received: from eggs.gnu.org ([2001:4830:134:3::10]:36367) by fencepost.gnu.org with esmtps (TLS1.0:RSA_AES_256_CBC_SHA1:256) (Exim 4.82) (envelope-from ) id 1Zm1zv-0001cw-BS for gcc-patches@gnu.org; Tue, 13 Oct 2015 12:00:19 -0400 Received: from Debian-exim by eggs.gnu.org with spam-scanned (Exim 4.71) (envelope-from ) id 1Zm1zq-0000Lg-Jx for gcc-patches@gnu.org; Tue, 13 Oct 2015 12:00:18 -0400 Received: from relay1.mentorg.com ([192.94.38.131]:46891) by eggs.gnu.org with esmtp (Exim 4.71) (envelope-from ) id 1Zm1zq-0000Ks-CD for gcc-patches@gnu.org; Tue, 13 Oct 2015 12:00:14 -0400 Received: from nat-ies.mentorg.com ([192.94.31.2] helo=SVR-IES-FEM-02.mgc.mentorg.com) by relay1.mentorg.com with esmtp id 1Zm1zo-0005ir-Fk from Tom_deVries@mentor.com ; Tue, 13 Oct 2015 09:00:12 -0700 Received: from [127.0.0.1] (137.202.0.76) by SVR-IES-FEM-02.mgc.mentorg.com (137.202.0.106) with Microsoft SMTP Server id 14.3.224.2; Tue, 13 Oct 2015 17:00:11 +0100 To: "gcc-patches@gnu.org" , Jakub Jelinek From: Tom de Vries Subject: [gomp4, committed] Move kernels pass group before pass_fre Message-ID: <561D2A5B.8080002@mentor.com> Date: Tue, 13 Oct 2015 16:00:00 -0000 User-Agent: Mozilla/5.0 (X11; Linux x86_64; rv:38.0) Gecko/20100101 Thunderbird/38.3.0 MIME-Version: 1.0 Content-Type: multipart/mixed; boundary="------------060908080000060804030502" X-detected-operating-system: by eggs.gnu.org: Windows NT kernel [generic] [fuzzy] X-Received-From: 192.94.38.131 X-SW-Source: 2015-10/txt/msg01275.txt.bz2 --------------060908080000060804030502 Content-Type: text/plain; charset="utf-8"; format=flowed Content-Transfer-Encoding: 7bit Content-length: 246 Hi, this patch moves the kernels pass group to before pass_fre. Instead we use pass_dominator_oacc_kernels in the pass group. This fixes an ICE while compiling the test-case included in the patch. Committed to gomp-4_0-branch. Thanks, - Tom --------------060908080000060804030502 Content-Type: text/x-patch; name="0002-Move-kernels-pass-group-before-pass_fre.patch" Content-Transfer-Encoding: 7bit Content-Disposition: inline; filename="0002-Move-kernels-pass-group-before-pass_fre.patch" Content-length: 4276 Move kernels pass group before pass_fre 2015-10-13 Tom de Vries * tree-ssa-dom.c (pass_dominator_oacc_kernels::clone): New function. * passes.def: Move pass group pass_oacc_kernels to before pass_fre. Add pass_dominator_oacc_kernels twice in the pass_oacc_kernels pass group. * c-c++-common/goacc/kernels-acc-on-device-2.c: New test. * c-c++-common/goacc/kernels-counter-var-redundant-load.c: Update. --- gcc/passes.def | 4 ++- .../c-c++-common/goacc/kernels-acc-on-device-2.c | 37 ++++++++++++++++++++++ .../goacc/kernels-counter-var-redundant-load.c | 10 +++--- gcc/tree-ssa-dom.c | 1 + 4 files changed, 47 insertions(+), 5 deletions(-) create mode 100644 gcc/testsuite/c-c++-common/goacc/kernels-acc-on-device-2.c diff --git a/gcc/passes.def b/gcc/passes.def index bc454c0..4ed4ccd 100644 --- a/gcc/passes.def +++ b/gcc/passes.def @@ -86,12 +86,13 @@ along with GCC; see the file COPYING3. If not see /* pass_build_ealias is a dummy pass that ensures that we execute TODO_rebuild_alias at this point. */ NEXT_PASS (pass_build_ealias); - NEXT_PASS (pass_fre); /* Pass group that runs when there are oacc kernels in the function. */ NEXT_PASS (pass_oacc_kernels); PUSH_INSERT_PASSES_WITHIN (pass_oacc_kernels) + NEXT_PASS (pass_dominator_oacc_kernels); NEXT_PASS (pass_ch_oacc_kernels); + NEXT_PASS (pass_dominator_oacc_kernels); NEXT_PASS (pass_tree_loop_init); NEXT_PASS (pass_lim); NEXT_PASS (pass_copy_prop); @@ -105,6 +106,7 @@ along with GCC; see the file COPYING3. If not see NEXT_PASS (pass_expand_omp_ssa); NEXT_PASS (pass_tree_loop_done); POP_INSERT_PASSES () + NEXT_PASS (pass_fre); NEXT_PASS (pass_merge_phi); NEXT_PASS (pass_dse); NEXT_PASS (pass_cd_dce); diff --git a/gcc/testsuite/c-c++-common/goacc/kernels-acc-on-device-2.c b/gcc/testsuite/c-c++-common/goacc/kernels-acc-on-device-2.c new file mode 100644 index 0000000..2c7297b --- /dev/null +++ b/gcc/testsuite/c-c++-common/goacc/kernels-acc-on-device-2.c @@ -0,0 +1,37 @@ +/* { dg-additional-options "-O2" } */ + +#include "openacc.h" + +#define N 32 + +void +foo (float *a, float *b) +{ +#pragma acc kernels copyin(a[0:N]) copyout(b[0:N]) + { + int ii; + int on_host = acc_on_device (acc_device_X); + + for (ii = 0; ii < N; ii++) + { + if (on_host) + b[ii] = a[ii] + 1; + else + b[ii] = a[ii]; + } + } + +#pragma acc kernels copyin(a[0:N]) copyout(b[0:N]) + { + int ii; + int on_host = acc_on_device (acc_device_X); + + for (ii = 0; ii < N; ii++) + { + if (on_host) + b[ii] = a[ii] + 2; + else + b[ii] = a[ii]; + } + } +} diff --git a/gcc/testsuite/c-c++-common/goacc/kernels-counter-var-redundant-load.c b/gcc/testsuite/c-c++-common/goacc/kernels-counter-var-redundant-load.c index 84dee69..c4ffc1d 100644 --- a/gcc/testsuite/c-c++-common/goacc/kernels-counter-var-redundant-load.c +++ b/gcc/testsuite/c-c++-common/goacc/kernels-counter-var-redundant-load.c @@ -1,5 +1,5 @@ /* { dg-additional-options "-O2" } */ -/* { dg-additional-options "-fdump-tree-dom_oacc_kernels" } */ +/* { dg-additional-options "-fdump-tree-dom_oacc_kernels3" } */ #include @@ -28,7 +28,9 @@ foo (unsigned int *c) _15 = .omp_data_i_10->c; c.1_16 = *_15; - Check that there's only one load from anonymous ssa-name (which we assume to - be the one to read c), and that there's no such load for ii. */ + Check that there are two loads from anonymous ssa-names, which we assume to + be: + - the one to read c + - the one to read ii after the kernels region. */ -/* { dg-final { scan-tree-dump-times "(?n)\\*_\[0-9\]\[0-9\]*;$" 1 "dom_oacc_kernels" } } */ +/* { dg-final { scan-tree-dump-times "(?n)\\*_\[0-9\]\[0-9\]*;$" 2 "dom_oacc_kernels3" } } */ diff --git a/gcc/tree-ssa-dom.c b/gcc/tree-ssa-dom.c index c7dc7b0..87f9daa 100644 --- a/gcc/tree-ssa-dom.c +++ b/gcc/tree-ssa-dom.c @@ -788,6 +788,7 @@ public: {} /* opt_pass methods: */ + opt_pass * clone () { return new pass_dominator_oacc_kernels (m_ctxt); } virtual bool gate (function *) { return true; } private: -- 1.9.1 --------------060908080000060804030502--