From patchwork Mon Oct 12 15:16:55 2015 Content-Type: text/plain; charset="utf-8" MIME-Version: 1.0 Content-Transfer-Encoding: 7bit X-Patchwork-Submitter: Tom de Vries X-Patchwork-Id: 529172 Return-Path: X-Original-To: incoming@patchwork.ozlabs.org Delivered-To: patchwork-incoming@bilbo.ozlabs.org Received: from sourceware.org (server1.sourceware.org [209.132.180.131]) (using TLSv1.2 with cipher ECDHE-RSA-AES256-GCM-SHA384 (256/256 bits)) (No client certificate requested) by ozlabs.org (Postfix) with ESMTPS id 3FB65140082 for ; Tue, 13 Oct 2015 02:18:01 +1100 (AEDT) Authentication-Results: ozlabs.org; dkim=pass (1024-bit key; unprotected) header.d=gcc.gnu.org header.i=@gcc.gnu.org header.b=oezCp4Cs; dkim-atps=neutral DomainKey-Signature: a=rsa-sha1; c=nofws; d=gcc.gnu.org; h=list-id :list-unsubscribe:list-archive:list-post:list-help:sender :subject:to:references:cc:from:message-id:date:mime-version :in-reply-to:content-type; q=dns; s=default; b=sp5//HLxlsj9xDDhQ n+nxD3gkFM/eXoz70wvP2ZjbdHEsAFw+gCZ3NfCFtvcXHCVY8gFMh78g3s+o+aY8 xEp0lXEj5gZgdfHHyZCs852yRVdGomGkGqmqqytxsHSSPhAFbzFIH4ncCVl9uWLt gK1oXiVzYGedrDK+IG1J0gGAxU= DKIM-Signature: v=1; a=rsa-sha1; c=relaxed; d=gcc.gnu.org; h=list-id :list-unsubscribe:list-archive:list-post:list-help:sender :subject:to:references:cc:from:message-id:date:mime-version :in-reply-to:content-type; s=default; bh=6+L3lUxeaaGy83J+8bmN5Wd rRBo=; b=oezCp4CsMIm357M5zQQqDrPFKGDLNCDgYaiot7zPLYuth4SJT5nAmjh JB++4kBg8QF9+6dRgEFUqNDze8IM5HR23iMSe9TSayj4BX3MvKZ/wOHa+TJh9jta C3wuuP0Op5qL0SFUPFn15z9IGFPOc4ANtYmJCXDoYk7FmFa94ZpU= Received: (qmail 41208 invoked by alias); 12 Oct 2015 15:17:53 -0000 Mailing-List: contact gcc-patches-help@gcc.gnu.org; run by ezmlm Precedence: bulk List-Id: List-Unsubscribe: List-Archive: List-Post: List-Help: Sender: gcc-patches-owner@gcc.gnu.org Delivered-To: mailing list gcc-patches@gcc.gnu.org Received: (qmail 41197 invoked by uid 89); 12 Oct 2015 15:17:53 -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; Mon, 12 Oct 2015 15:17:51 +0000 Received: from eggs.gnu.org ([2001:4830:134:3::10]:58271) by fencepost.gnu.org with esmtps (TLS1.0:RSA_AES_256_CBC_SHA1:256) (Exim 4.82) (envelope-from ) id 1ZlerE-00087T-Jz for gcc-patches@gnu.org; Mon, 12 Oct 2015 11:17:49 -0400 Received: from Debian-exim by eggs.gnu.org with spam-scanned (Exim 4.71) (envelope-from ) id 1ZlerA-00053z-TG for gcc-patches@gnu.org; Mon, 12 Oct 2015 11:17:48 -0400 Received: from relay1.mentorg.com ([192.94.38.131]:33067) by eggs.gnu.org with esmtp (Exim 4.71) (envelope-from ) id 1ZlerA-00053s-JR for gcc-patches@gnu.org; Mon, 12 Oct 2015 11:17:44 -0400 Received: from nat-ies.mentorg.com ([192.94.31.2] helo=SVR-IES-FEM-01.mgc.mentorg.com) by relay1.mentorg.com with esmtp id 1Zler8-0001zt-SQ from Tom_deVries@mentor.com ; Mon, 12 Oct 2015 08:17:43 -0700 Received: from [127.0.0.1] (137.202.0.76) by SVR-IES-FEM-01.mgc.mentorg.com (137.202.0.104) with Microsoft SMTP Server id 14.3.224.2; Mon, 12 Oct 2015 16:17:41 +0100 Subject: [committed. gomp4, 6/6] Add pass_dominator_oacc_kernels To: "gcc-patches@gnu.org" References: <561BC874.2020908@mentor.com> CC: Jakub Jelinek From: Tom de Vries Message-ID: <561BCEE7.1090307@mentor.com> Date: Mon, 12 Oct 2015 17:16:55 +0200 User-Agent: Mozilla/5.0 (X11; Linux x86_64; rv:38.0) Gecko/20100101 Thunderbird/38.2.0 MIME-Version: 1.0 In-Reply-To: <561BC874.2020908@mentor.com> X-detected-operating-system: by eggs.gnu.org: Windows NT kernel [generic] [fuzzy] X-Received-From: 192.94.38.131 On 12/10/15 16:49, Tom de Vries wrote: > Hi, > > I've committed the following patch series to the gomp-4_0-branch. > > 1 Add pass_dominator::jump_threading_p () > 2 Add dom_walker::walk_until > 3 Add pass_dominator::sese_mode_p () > 4 Add skip_stmt parm to pass_dominator::get_sese () > 5 Add oacc kernels related infra functions > 6 Add pass_dominator_oacc_kernels > > The patch series adds a pass pass_dominator_oacc_kernels, which does the > pass_dominator optimizations (with the exception of jump threading) on > each oacc kernels region rather than on the whole function. > > Bootstrapped and reg-tested on x86_64. > > I'll post the patches individually, in reply to this email. This patch : - factors a class dominator_base out of class pass_dominators, - declares a new class pass_dominators_oacc_kernels, that operates on oacc kernels regions, and - adds the new pass before pass_parallelize_loops_oacc_kernels in the oacc kernels pass group. Thanks, - Tom Add pass_dominator_oacc_kernels 2015-10-12 Tom de Vries * passes.def: Add pass_dominator_oacc_kernels to pass group pass_oacc_kernels. Add pass_tree_loop_done before, and pass_tree_loop_init after. * tree-pass.h (make_pass_dominator_oacc_kernels): Declare. * tree-ssa-dom.c (class dominator_base): New class. Factor out of ... (class pass_dominator): ... here. (pass_dominator_oacc_kernels): New pass. (make_pass_dominator_oacc_kernels): New function. * c-c++-common/goacc/kernels-counter-var-redundant-load.c: New test. --- gcc/passes.def | 3 + .../goacc/kernels-counter-var-redundant-load.c | 34 ++++++ gcc/tree-pass.h | 1 + gcc/tree-ssa-dom.c | 117 +++++++++++++++++---- 4 files changed, 134 insertions(+), 21 deletions(-) create mode 100644 gcc/testsuite/c-c++-common/goacc/kernels-counter-var-redundant-load.c diff --git a/gcc/passes.def b/gcc/passes.def index 0498a8b..bc454c0 100644 --- a/gcc/passes.def +++ b/gcc/passes.def @@ -98,6 +98,9 @@ along with GCC; see the file COPYING3. If not see NEXT_PASS (pass_lim); NEXT_PASS (pass_copy_prop); NEXT_PASS (pass_scev_cprop); + NEXT_PASS (pass_tree_loop_done); + NEXT_PASS (pass_dominator_oacc_kernels); + NEXT_PASS (pass_tree_loop_init); NEXT_PASS (pass_parallelize_loops_oacc_kernels); NEXT_PASS (pass_expand_omp_ssa); NEXT_PASS (pass_tree_loop_done); 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 new file mode 100644 index 0000000..84dee69 --- /dev/null +++ b/gcc/testsuite/c-c++-common/goacc/kernels-counter-var-redundant-load.c @@ -0,0 +1,34 @@ +/* { dg-additional-options "-O2" } */ +/* { dg-additional-options "-fdump-tree-dom_oacc_kernels" } */ + +#include + +#define N (1024 * 512) +#define COUNTERTYPE unsigned int + +COUNTERTYPE +foo (unsigned int *c) +{ + COUNTERTYPE ii; + +#pragma acc kernels copyout (c[0:N]) + { + for (ii = 0; ii < N; ii++) + c[ii] = 1; + } + + return ii; +} + +/* We're expecting: + + .omp_data_i_10 = &.omp_data_arr.3; + _11 = .omp_data_i_10->ii; + *_11 = 0; + _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. */ + +/* { dg-final { scan-tree-dump-times "(?n)\\*_\[0-9\]\[0-9\]*;$" 1 "dom_oacc_kernels" } } */ diff --git a/gcc/tree-pass.h b/gcc/tree-pass.h index 52ba3e5..15c8bf6 100644 --- a/gcc/tree-pass.h +++ b/gcc/tree-pass.h @@ -392,6 +392,7 @@ extern gimple_opt_pass *make_pass_build_ssa (gcc::context *ctxt); extern gimple_opt_pass *make_pass_build_alias (gcc::context *ctxt); extern gimple_opt_pass *make_pass_build_ealias (gcc::context *ctxt); extern gimple_opt_pass *make_pass_dominator (gcc::context *ctxt); +extern gimple_opt_pass *make_pass_dominator_oacc_kernels (gcc::context *ctxt); extern gimple_opt_pass *make_pass_dce (gcc::context *ctxt); extern gimple_opt_pass *make_pass_cd_dce (gcc::context *ctxt); extern gimple_opt_pass *make_pass_call_cdce (gcc::context *ctxt); diff --git a/gcc/tree-ssa-dom.c b/gcc/tree-ssa-dom.c index 573e6fc..c7dc7b0 100644 --- a/gcc/tree-ssa-dom.c +++ b/gcc/tree-ssa-dom.c @@ -45,6 +45,7 @@ along with GCC; see the file COPYING3. If not see #include "gimplify.h" #include "tree-cfgcleanup.h" #include "cfgcleanup.h" +#include "omp-low.h" /* This file implements optimizations on the dominator tree. */ @@ -526,6 +527,31 @@ private: namespace { +class dominator_base : public gimple_opt_pass +{ + protected: + dominator_base (pass_data data, gcc::context *ctxt) + : gimple_opt_pass (data, ctxt) + {} + + unsigned int execute (function *); + + /* Return true if pass should perform jump threading. */ + virtual bool jump_threading_p (void) { return !sese_mode_p (); } + + /* Return true if pass should visit a series of seses rather than the whole + dominator tree. */ + virtual bool sese_mode_p (void) { return false; } + + /* In sese mode, return true if there's another sese to visit. Return the + sese to visit in SESE_ENTRY and SESE_EXIT. If a stmt in the sese should + not be optimized, return it in SKIP_STMT. */ + virtual bool get_sese (basic_block *sese_entry ATTRIBUTE_UNUSED, + basic_block *sese_exit ATTRIBUTE_UNUSED, + gimple **skip_stmt ATTRIBUTE_UNUSED) + { gcc_unreachable (); } +}; // class dominator_base + const pass_data pass_data_dominator = { GIMPLE_PASS, /* type */ @@ -539,38 +565,20 @@ const pass_data pass_data_dominator = ( TODO_cleanup_cfg | TODO_update_ssa ), /* todo_flags_finish */ }; -class pass_dominator : public gimple_opt_pass +class pass_dominator : public dominator_base { public: pass_dominator (gcc::context *ctxt) - : gimple_opt_pass (pass_data_dominator, ctxt) + : dominator_base (pass_data_dominator, ctxt) {} /* opt_pass methods: */ opt_pass * clone () { return new pass_dominator (m_ctxt); } virtual bool gate (function *) { return flag_tree_dom != 0; } - virtual unsigned int execute (function *); - - protected: - /* Return true if pass should perform jump threading. */ - virtual bool jump_threading_p (void) { return !sese_mode_p (); } - - /* Return true if pass should visit a series of seses rather than the whole - dominator tree. */ - virtual bool sese_mode_p (void) { return false; } - - /* In sese mode, return true if there's another sese to visit. Return the - sese to visit in SESE_ENTRY and SESE_EXIT. If a stmt in the sese should - not be optimized, return it in SKIP_STMT. */ - virtual bool get_sese (basic_block *sese_entry ATTRIBUTE_UNUSED, - basic_block *sese_exit ATTRIBUTE_UNUSED, - gimple **skip_stmt ATTRIBUTE_UNUSED) - { gcc_unreachable (); } - }; // class pass_dominator unsigned int -pass_dominator::execute (function *fun) +dominator_base::execute (function *fun) { memset (&opt_stats, 0, sizeof (opt_stats)); @@ -759,6 +767,68 @@ pass_dominator::execute (function *fun) return 0; } +const pass_data pass_data_dominator_oacc_kernels = +{ + GIMPLE_PASS, /* type */ + "dom_oacc_kernels", /* name */ + OPTGROUP_NONE, /* optinfo_flags */ + TV_TREE_SSA_DOMINATOR_OPTS, /* tv_id */ + ( PROP_cfg | PROP_ssa ), /* properties_required */ + 0, /* properties_provided */ + 0, /* properties_destroyed */ + 0, /* todo_flags_start */ + ( TODO_cleanup_cfg | TODO_update_ssa ), /* todo_flags_finish */ +}; + +class pass_dominator_oacc_kernels : public dominator_base +{ +public: + pass_dominator_oacc_kernels (gcc::context *ctxt) + : dominator_base (pass_data_dominator_oacc_kernels, ctxt), m_regions (NULL) + {} + + /* opt_pass methods: */ + virtual bool gate (function *) { return true; } + + private: + bitmap m_regions; + +protected: + /* dominator_base methods: */ + virtual bool sese_mode_p (void) { return true; } + virtual bool get_sese (basic_block *sese_entry, basic_block *sese_exit, + gimple **skip_stmt) + { + if (m_regions == NULL) + { + m_regions = BITMAP_ALLOC (NULL); + basic_block bb; + FOR_EACH_BB_FN (bb, cfun) + if (oacc_kernels_region_entry_p (bb, NULL)) + bitmap_set_bit (m_regions, bb->index); + } + + if (bitmap_empty_p (m_regions)) + { + BITMAP_FREE (m_regions); + return false; + } + + unsigned int index = bitmap_first_set_bit (m_regions); + bitmap_clear_bit (m_regions, index); + + *sese_entry = BASIC_BLOCK_FOR_FN (cfun, index); + *sese_exit = get_oacc_kernels_region_exit (*sese_entry); + + tree omp_data_i = get_omp_data_i (single_pred (*sese_entry)); + if (omp_data_i != NULL_TREE) + *skip_stmt = SSA_NAME_DEF_STMT (omp_data_i); + + return true; + } + +}; // class pass_dominator_oacc_kernels + } // anon namespace gimple_opt_pass * @@ -767,6 +837,11 @@ make_pass_dominator (gcc::context *ctxt) return new pass_dominator (ctxt); } +gimple_opt_pass * +make_pass_dominator_oacc_kernels (gcc::context *ctxt) +{ + return new pass_dominator_oacc_kernels (ctxt); +} /* Given a conditional statement CONDSTMT, convert the condition to a canonical form. */ -- 1.9.1