From patchwork Wed Dec 10 09:57:58 2014 Content-Type: text/plain; charset="utf-8" MIME-Version: 1.0 Content-Transfer-Encoding: 8bit X-Patchwork-Submitter: Thomas Schwinge X-Patchwork-Id: 419504 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 31E631400B7 for ; Wed, 10 Dec 2014 20:58:28 +1100 (AEDT) DomainKey-Signature: a=rsa-sha1; c=nofws; d=gcc.gnu.org; h=list-id :list-unsubscribe:list-archive:list-post:list-help:sender:from :to:cc:subject:in-reply-to:references:date:message-id :mime-version:content-type; q=dns; s=default; b=W8Pd8gN86Mb1WUPd XR7mpZJZwgOdPSzqqV+n93HuZ4/wgh+0bDURr3npj3BVRwOr189752FGm4SOrlaA nAzUgAF5jb8InvYmRIvhw+5pkCIKv7IeZ2A6T36YhonPvA23r0URVLMXN3p+AxRn bsI6Dtj0BqUe5fkhD6DHlIwA8Qo= 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:from :to:cc:subject:in-reply-to:references:date:message-id :mime-version:content-type; s=default; bh=10x+HQuIq7lB/+jnf/RVih 25Stc=; b=R5zw9YoUktMmGCkv19Qq9sfQE5CLVK2bwmcOz6Qg2A1zRvaCtL3KjC JOc3qQSKXpD9z5BpxmkTcgOLhrPoux0SbBwLH3PFNzMPeSTOLH9++Stad1ms8u1E kYo1qOQpj+jbTGTJUzC1ZhiII4cQeAqxV/+1Q+JzsRxmrAhsoD0uU= Received: (qmail 9799 invoked by alias); 10 Dec 2014 09:58:20 -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 9787 invoked by uid 89); 10 Dec 2014 09:58:19 -0000 Authentication-Results: sourceware.org; auth=none X-Virus-Found: No X-Spam-SWARE-Status: No, score=-1.9 required=5.0 tests=AWL, BAYES_00, RCVD_IN_DNSWL_NONE autolearn=ham version=3.3.2 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; Wed, 10 Dec 2014 09:58:12 +0000 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 1Xye23-0002wm-Di from Thomas_Schwinge@mentor.com ; Wed, 10 Dec 2014 01:58:08 -0800 Received: from feldtkeller.schwinge.homeip.net (137.202.0.76) by SVR-IES-FEM-02.mgc.mentorg.com (137.202.0.106) with Microsoft SMTP Server id 14.3.181.6; Wed, 10 Dec 2014 09:58:04 +0000 From: Thomas Schwinge To: , Jakub Jelinek CC: Cesar Philippidis , Tom de Vries , Bernd Schmidt Subject: Re: OpenACC GIMPLE_OACC_* -- or not? In-Reply-To: <20141112134502.GC5026@tucnak.redhat.com> References: <1383766943-8863-2-git-send-email-thomas@codesourcery.com> <1383766943-8863-3-git-send-email-thomas@codesourcery.com> <1383766943-8863-4-git-send-email-thomas@codesourcery.com> <1383766943-8863-5-git-send-email-thomas@codesourcery.com> <1383766943-8863-6-git-send-email-thomas@codesourcery.com> <1383766943-8863-7-git-send-email-thomas@codesourcery.com> <1383766943-8863-8-git-send-email-thomas@codesourcery.com> <1383766943-8863-9-git-send-email-thomas@codesourcery.com> <8761s5joir.fsf@schwinge.name> <87sihoczm0.fsf@kepler.schwinge.homeip.net> <20141112134502.GC5026@tucnak.redhat.com> User-Agent: Notmuch/0.9-101-g81dad07 (http://notmuchmail.org) Emacs/24.4.1 (i586-pc-linux-gnu) Date: Wed, 10 Dec 2014 10:57:58 +0100 Message-ID: <87wq5z7ro9.fsf@kepler.schwinge.homeip.net> MIME-Version: 1.0 Hi! On Wed, 12 Nov 2014 14:45:02 +0100, Jakub Jelinek wrote: > On Wed, Nov 12, 2014 at 02:33:43PM +0100, Thomas Schwinge wrote: > > Months later, with months' worth of GCC internals experience, I now came > > to realize that maybe this has not actually been a useful thing to do > > (and likewise for the GIMPLE_OACC_KERNELS also added later on, > > ). > > All handling of GIMPLE_OACC_PARALLEL and GIMPLE_OACC_KERNELS closely > > follows that of GIMPLE_OMP_TARGET's GF_OMP_TARGET_KIND_REGION, with only > > minor divergence. What I did not understand back then, has not been > > obvious to me, was that the underlying structure of all those codes will > > in fact be the same (as already made apparent by using the one > > GIMPLE_OMP_TARGET for all of: OpenMP target offloading regions, OpenMP > > target data regions, OpenMP target data maintenenace "executable" > > statements), and any "customization" then happens via the clauses > > attached to GIMPLE_OMP_TARGET. > > I'm fine with merging them into kinds, [...] > > > So, sanity check: should we now merge GIMPLE_OACC_PARALLEL and > > GIMPLE_OACC_KERNELS into being "subtypes" of GIMPLE_OMP_TARGET (like > > GF_OMP_TARGET_KIND_REGION), as already done for > > GF_OMP_TARGET_KIND_OACC_DATA (like GF_OMP_TARGET_KIND_DATA), and > > GF_OMP_TARGET_KIND_OACC_UPDATE and > > GF_OMP_TARGET_KIND_OACC_ENTER_EXIT_DATA (like GF_OMP_TARGET_KIND_UPDATE). > > Yep. In r218568, I applied the following to gomp-4_0-branch: commit 28629d718a63a782170cfb06a4d0278de0779039 Author: tschwinge Date: Wed Dec 10 09:52:28 2014 +0000 Merge GIMPLE_OACC_KERNELS and GIMPLE_OACC_PARALLEL into GIMPLE_OMP_TARGET. gcc/ * gimple.def (GIMPLE_OACC_KERNELS, GIMPLE_OACC_PARALLEL): Merge into GIMPLE_OMP_TARGET. Update all users. gcc/ * cgraphbuild.c (pass_build_cgraph_edges::execute): Remove handling of GIMPLE_OACC_PARALLEL. * gimple-pretty-print.c (dump_gimple_omp_target): Dump a bit more data, pretty-printing. git-svn-id: svn+ssh://gcc.gnu.org/svn/gcc/branches/gomp-4_0-branch@218568 138bc75d-0d04-0410-961f-82ee72b054a4 --- gcc/ChangeLog.gomp | 11 ++ gcc/cgraphbuild.c | 13 +- gcc/doc/gimple.texi | 15 -- gcc/gimple-low.c | 2 - gcc/gimple-pretty-print.c | 118 +++-------- gcc/gimple-walk.c | 32 --- gcc/gimple.c | 40 +--- gcc/gimple.def | 35 +--- gcc/gimple.h | 273 ++------------------------ gcc/gimplify.c | 6 +- gcc/omp-low.c | 225 +++++++-------------- gcc/testsuite/gfortran.dg/goacc/private-1.f95 | 2 +- gcc/tree-inline.c | 6 - gcc/tree-nested.c | 16 -- 14 files changed, 136 insertions(+), 658 deletions(-) Grüße, Thomas diff --git gcc/ChangeLog.gomp gcc/ChangeLog.gomp index bece7c1..06e8583 100644 --- gcc/ChangeLog.gomp +++ gcc/ChangeLog.gomp @@ -1,4 +1,15 @@ 2014-12-10 Thomas Schwinge + Bernd Schmidt + + * gimple.def (GIMPLE_OACC_KERNELS, GIMPLE_OACC_PARALLEL): Merge + into GIMPLE_OMP_TARGET. Update all users. + +2014-12-10 Thomas Schwinge + + * cgraphbuild.c (pass_build_cgraph_edges::execute): Remove + handling of GIMPLE_OACC_PARALLEL. + * gimple-pretty-print.c (dump_gimple_omp_target): Dump a bit more + data, pretty-printing. * omp-low.c (build_omp_regions_1, make_gimple_omp_edges) : Handle diff --git gcc/cgraphbuild.c gcc/cgraphbuild.c index 9b078bc..c72ceab 100644 --- gcc/cgraphbuild.c +++ gcc/cgraphbuild.c @@ -368,21 +368,14 @@ pass_build_cgraph_edges::execute (function *fun) bb->count, freq); } node->record_stmt_references (stmt); - if (gimple_code (stmt) == GIMPLE_OACC_PARALLEL - && gimple_oacc_parallel_child_fn (stmt)) - { - tree fn = gimple_oacc_parallel_child_fn (stmt); - node->create_reference (cgraph_node::get_create (fn), - IPA_REF_ADDR, stmt); - } - else if (gimple_code (stmt) == GIMPLE_OMP_PARALLEL - && gimple_omp_parallel_child_fn (stmt)) + if (gimple_code (stmt) == GIMPLE_OMP_PARALLEL + && gimple_omp_parallel_child_fn (stmt)) { tree fn = gimple_omp_parallel_child_fn (stmt); node->create_reference (cgraph_node::get_create (fn), IPA_REF_ADDR, stmt); } - else if (gimple_code (stmt) == GIMPLE_OMP_TASK) + if (gimple_code (stmt) == GIMPLE_OMP_TASK) { tree fn = gimple_omp_task_child_fn (stmt); if (fn) diff --git gcc/doc/gimple.texi gcc/doc/gimple.texi index 4c59748..696c10e 100644 --- gcc/doc/gimple.texi +++ gcc/doc/gimple.texi @@ -439,8 +439,6 @@ The following table briefly describes the GIMPLE instruction set. @item @code{GIMPLE_GOTO} @tab x @tab x @item @code{GIMPLE_LABEL} @tab x @tab x @item @code{GIMPLE_NOP} @tab x @tab x -@item @code{GIMPLE_OACC_KERNELS} @tab x @tab x -@item @code{GIMPLE_OACC_PARALLEL} @tab x @tab x @item @code{GIMPLE_OMP_ATOMIC_LOAD} @tab x @tab x @item @code{GIMPLE_OMP_ATOMIC_STORE} @tab x @tab x @item @code{GIMPLE_OMP_CONTINUE} @tab x @tab x @@ -1008,8 +1006,6 @@ Return a deep copy of statement @code{STMT}. * @code{GIMPLE_EH_FILTER}:: * @code{GIMPLE_LABEL}:: * @code{GIMPLE_NOP}:: -* @code{GIMPLE_OACC_KERNELS}:: -* @code{GIMPLE_OACC_PARALLEL}:: * @code{GIMPLE_OMP_ATOMIC_LOAD}:: * @code{GIMPLE_OMP_ATOMIC_STORE}:: * @code{GIMPLE_OMP_CONTINUE}:: @@ -1655,17 +1651,6 @@ Build a @code{GIMPLE_NOP} statement. Returns @code{TRUE} if statement @code{G} is a @code{GIMPLE_NOP}. @end deftypefn - -@node @code{GIMPLE_OACC_KERNELS} -@subsection @code{GIMPLE_OACC_KERNELS} -@cindex @code{GIMPLE_OACC_KERNELS} - - -@node @code{GIMPLE_OACC_PARALLEL} -@subsection @code{GIMPLE_OACC_PARALLEL} -@cindex @code{GIMPLE_OACC_PARALLEL} - - @node @code{GIMPLE_OMP_ATOMIC_LOAD} @subsection @code{GIMPLE_OMP_ATOMIC_LOAD} @cindex @code{GIMPLE_OMP_ATOMIC_LOAD} diff --git gcc/gimple-low.c gcc/gimple-low.c index 60a7792e..3507d3c 100644 --- gcc/gimple-low.c +++ gcc/gimple-low.c @@ -368,8 +368,6 @@ lower_stmt (gimple_stmt_iterator *gsi, struct lower_data *data) } break; - case GIMPLE_OACC_KERNELS: - case GIMPLE_OACC_PARALLEL: case GIMPLE_OMP_PARALLEL: case GIMPLE_OMP_TASK: case GIMPLE_OMP_TARGET: diff --git gcc/gimple-pretty-print.c gcc/gimple-pretty-print.c index 72dfac6..38d39f7 100644 --- gcc/gimple-pretty-print.c +++ gcc/gimple-pretty-print.c @@ -1338,15 +1338,21 @@ dump_gimple_omp_target (pretty_printer *buffer, gimple gs, int spc, int flags) case GF_OMP_TARGET_KIND_UPDATE: kind = " update"; break; + case GF_OMP_TARGET_KIND_OACC_KERNELS: + kind = " oacc_kernels"; + break; + case GF_OMP_TARGET_KIND_OACC_PARALLEL: + kind = " oacc_parallel"; + break; case GF_OMP_TARGET_KIND_OACC_DATA: kind = " oacc_data"; break; - case GF_OMP_TARGET_KIND_OACC_ENTER_EXIT_DATA: - kind = " oacc_enter_exit_data"; - break; case GF_OMP_TARGET_KIND_OACC_UPDATE: kind = " oacc_update"; break; + case GF_OMP_TARGET_KIND_OACC_ENTER_EXIT_DATA: + kind = " oacc_enter_exit_data"; + break; default: gcc_unreachable (); } @@ -1355,7 +1361,9 @@ dump_gimple_omp_target (pretty_printer *buffer, gimple gs, int spc, int flags) dump_gimple_fmt (buffer, spc, flags, "%G%s <%+BODY <%S>%nCLAUSES <", gs, kind, gimple_omp_body (gs)); dump_omp_clauses (buffer, gimple_omp_target_clauses (gs), spc, flags); - dump_gimple_fmt (buffer, spc, flags, " >"); + dump_gimple_fmt (buffer, spc, flags, " >, %T, %T%n>", + gimple_omp_target_child_fn (gs), + gimple_omp_target_data_arg (gs)); } else { @@ -1367,16 +1375,28 @@ dump_gimple_omp_target (pretty_printer *buffer, gimple gs, int spc, int flags) pp_string (buffer, " [child fn: "); dump_generic_node (buffer, gimple_omp_target_child_fn (gs), spc, flags, false); - pp_right_bracket (buffer); + pp_string (buffer, " ("); + if (gimple_omp_target_data_arg (gs)) + dump_generic_node (buffer, gimple_omp_target_data_arg (gs), + spc, flags, false); + else + pp_string (buffer, "???"); + pp_string (buffer, ")]"); } - if (!gimple_seq_empty_p (gimple_omp_body (gs))) + gimple_seq body = gimple_omp_body (gs); + if (body && gimple_code (gimple_seq_first_stmt (body)) != GIMPLE_BIND) { newline_and_indent (buffer, spc + 2); - pp_character (buffer, '{'); + pp_left_brace (buffer); pp_newline (buffer); - dump_gimple_seq (buffer, gimple_omp_body (gs), spc + 4, flags); + dump_gimple_seq (buffer, body, spc + 4, flags); newline_and_indent (buffer, spc + 2); - pp_character (buffer, '}'); + pp_right_brace (buffer); + } + else if (body) + { + pp_newline (buffer); + dump_gimple_seq (buffer, body, spc + 2, flags); } } } @@ -1878,81 +1898,6 @@ dump_gimple_phi (pretty_printer *buffer, gimple phi, int spc, bool comment, } -/* Dump an OpenACC offload tuple on the pretty_printer BUFFER, SPC spaces - of indent. FLAGS specifies details to show in the dump (see TDF_* in - dumpfile.h). */ - -static void -dump_gimple_oacc_offload (pretty_printer *buffer, gimple gs, int spc, - int flags) -{ - tree (*gimple_omp_clauses) (const_gimple); - tree (*gimple_omp_child_fn) (const_gimple); - tree (*gimple_omp_data_arg) (const_gimple); - const char *kind; - switch (gimple_code (gs)) - { - case GIMPLE_OACC_KERNELS: - gimple_omp_clauses = gimple_oacc_kernels_clauses; - gimple_omp_child_fn = gimple_oacc_kernels_child_fn; - gimple_omp_data_arg = gimple_oacc_kernels_data_arg; - kind = "kernels"; - break; - case GIMPLE_OACC_PARALLEL: - gimple_omp_clauses = gimple_oacc_parallel_clauses; - gimple_omp_child_fn = gimple_oacc_parallel_child_fn; - gimple_omp_data_arg = gimple_oacc_parallel_data_arg; - kind = "parallel"; - break; - default: - gcc_unreachable (); - } - if (flags & TDF_RAW) - { - dump_gimple_fmt (buffer, spc, flags, "%G <%+BODY <%S>%nCLAUSES <", gs, - gimple_omp_body (gs)); - dump_omp_clauses (buffer, gimple_omp_clauses (gs), spc, flags); - dump_gimple_fmt (buffer, spc, flags, " >, %T, %T%n>", - gimple_omp_child_fn (gs), gimple_omp_data_arg (gs)); - } - else - { - gimple_seq body; - pp_string (buffer, "#pragma acc "); - pp_string (buffer, kind); - dump_omp_clauses (buffer, gimple_omp_clauses (gs), spc, flags); - if (gimple_omp_child_fn (gs)) - { - pp_string (buffer, " [child fn: "); - dump_generic_node (buffer, gimple_omp_child_fn (gs), - spc, flags, false); - pp_string (buffer, " ("); - if (gimple_omp_data_arg (gs)) - dump_generic_node (buffer, gimple_omp_data_arg (gs), - spc, flags, false); - else - pp_string (buffer, "???"); - pp_string (buffer, ")]"); - } - body = gimple_omp_body (gs); - if (body && gimple_code (gimple_seq_first_stmt (body)) != GIMPLE_BIND) - { - newline_and_indent (buffer, spc + 2); - pp_left_brace (buffer); - pp_newline (buffer); - dump_gimple_seq (buffer, body, spc + 4, flags); - newline_and_indent (buffer, spc + 2); - pp_right_brace (buffer); - } - else if (body) - { - pp_newline (buffer); - dump_gimple_seq (buffer, body, spc + 2, flags); - } - } -} - - /* Dump a GIMPLE_OMP_PARALLEL tuple on the pretty_printer BUFFER, SPC spaces of indent. FLAGS specifies details to show in the dump (see TDF_* in dumpfile.h). */ @@ -2237,11 +2182,6 @@ pp_gimple_stmt_1 (pretty_printer *buffer, gimple gs, int spc, int flags) dump_gimple_phi (buffer, gs, spc, false, flags); break; - case GIMPLE_OACC_KERNELS: - case GIMPLE_OACC_PARALLEL: - dump_gimple_oacc_offload (buffer, gs, spc, flags); - break; - case GIMPLE_OMP_PARALLEL: dump_gimple_omp_parallel (buffer, gs, spc, flags); break; diff --git gcc/gimple-walk.c gcc/gimple-walk.c index cc74d34..bfa3532 100644 --- gcc/gimple-walk.c +++ gcc/gimple-walk.c @@ -304,36 +304,6 @@ walk_gimple_op (gimple stmt, walk_tree_fn callback_op, return ret; break; - case GIMPLE_OACC_KERNELS: - ret = walk_tree (gimple_oacc_kernels_clauses_ptr (stmt), callback_op, - wi, pset); - if (ret) - return ret; - ret = walk_tree (gimple_oacc_kernels_child_fn_ptr (stmt), callback_op, - wi, pset); - if (ret) - return ret; - ret = walk_tree (gimple_oacc_kernels_data_arg_ptr (stmt), callback_op, - wi, pset); - if (ret) - return ret; - break; - - case GIMPLE_OACC_PARALLEL: - ret = walk_tree (gimple_oacc_parallel_clauses_ptr (stmt), callback_op, - wi, pset); - if (ret) - return ret; - ret = walk_tree (gimple_oacc_parallel_child_fn_ptr (stmt), callback_op, - wi, pset); - if (ret) - return ret; - ret = walk_tree (gimple_oacc_parallel_data_arg_ptr (stmt), callback_op, - wi, pset); - if (ret) - return ret; - break; - case GIMPLE_OMP_CONTINUE: ret = walk_tree (gimple_omp_continue_control_def_ptr (stmt), callback_op, wi, pset); @@ -629,8 +599,6 @@ walk_gimple_stmt (gimple_stmt_iterator *gsi, walk_stmt_fn callback_stmt, return wi->callback_result; /* FALL THROUGH. */ - case GIMPLE_OACC_KERNELS: - case GIMPLE_OACC_PARALLEL: case GIMPLE_OMP_CRITICAL: case GIMPLE_OMP_MASTER: case GIMPLE_OMP_TASKGROUP: diff --git gcc/gimple.c gcc/gimple.c index e6de836..32615e8 100644 --- gcc/gimple.c +++ gcc/gimple.c @@ -811,40 +811,6 @@ gimple_build_debug_source_bind_stat (tree var, tree value, } -/* Build a GIMPLE_OACC_KERNELS statement. - - BODY is sequence of statements which are executed as kernels. - CLAUSES are the OpenACC kernels construct's clauses. */ - -gimple -gimple_build_oacc_kernels (gimple_seq body, tree clauses) -{ - gimple p = gimple_alloc (GIMPLE_OACC_KERNELS, 0); - if (body) - gimple_omp_set_body (p, body); - gimple_oacc_kernels_set_clauses (p, clauses); - - return p; -} - - -/* Build a GIMPLE_OACC_PARALLEL statement. - - BODY is sequence of statements which are executed in parallel. - CLAUSES are the OpenACC parallel construct's clauses. */ - -gimple -gimple_build_oacc_parallel (gimple_seq body, tree clauses) -{ - gimple p = gimple_alloc (GIMPLE_OACC_PARALLEL, 0); - if (body) - gimple_omp_set_body (p, body); - gimple_oacc_parallel_set_clauses (p, clauses); - - return p; -} - - /* Build a GIMPLE_OMP_CRITICAL statement. BODY is the sequence of statements for which only one thread can execute. @@ -1077,7 +1043,7 @@ gimple_build_omp_single (gimple_seq body, tree clauses) /* Build a GIMPLE_OMP_TARGET statement. BODY is the sequence of statements that will be executed. - KIND is the kind of target region. + KIND is the kind of the region. CLAUSES are any of the construct's clauses. */ gimple @@ -1719,10 +1685,6 @@ gimple_copy (gimple stmt) gimple_try_set_cleanup (copy, new_seq); break; - case GIMPLE_OACC_KERNELS: - case GIMPLE_OACC_PARALLEL: - gcc_unreachable (); - case GIMPLE_OMP_FOR: gcc_assert (!is_gimple_omp_oacc_specifically (stmt)); new_seq = gimple_seq_copy (gimple_omp_for_pre_body (stmt)); diff --git gcc/gimple.def gcc/gimple.def index e2e912c..269c2d7 100644 --- gcc/gimple.def +++ gcc/gimple.def @@ -205,33 +205,8 @@ DEFGSCODE(GIMPLE_NOP, "gimple_nop", GSS_BASE) /* IMPORTANT. - Do not rearrange any of the GIMPLE_OACC_* and GIMPLE_OMP_* codes. This - ordering is exposed by the range check in gimple_omp_subcode. */ - - -/* GIMPLE_OACC_KERNELS represents - #pragma acc kernels [CLAUSES] - BODY is the sequence of statements inside the kernels construct. - CLAUSES is an OMP_CLAUSE chain holding the associated clauses. - CHILD_FN is set when outlining the body of the kernels region. - All the statements in BODY are moved into this newly created - function when converting OMP constructs into low-GIMPLE. - DATA_ARG is a vec of 3 local variables in the parent function - containing data to be mapped to CHILD_FN. This is used to - implement the MAP clauses. */ -DEFGSCODE(GIMPLE_OACC_KERNELS, "gimple_oacc_kernels", GSS_OMP_PARALLEL_LAYOUT) - -/* GIMPLE_OACC_PARALLEL represents - #pragma acc parallel [CLAUSES] - BODY is the sequence of statements inside the parallel construct. - CLAUSES is an OMP_CLAUSE chain holding the associated clauses. - CHILD_FN is set when outlining the body of the parallel region. - All the statements in BODY are moved into this newly created - function when converting OMP constructs into low-GIMPLE. - DATA_ARG is a vec of 3 local variables in the parent function - containing data to be mapped to CHILD_FN. This is used to - implement the MAP clauses. */ -DEFGSCODE(GIMPLE_OACC_PARALLEL, "gimple_oacc_parallel", GSS_OMP_PARALLEL_LAYOUT) + Do not rearrange any of the GIMPLE_OMP_* codes. This ordering is + exposed by the range check in gimple_omp_subcode. */ /* Tuples used for lowering of OMP_ATOMIC. Although the form of the OMP_ATOMIC expression is very simple (just in form mem op= expr), various implicit @@ -381,12 +356,12 @@ DEFGSCODE(GIMPLE_OMP_SECTIONS_SWITCH, "gimple_omp_sections_switch", GSS_BASE) DEFGSCODE(GIMPLE_OMP_SINGLE, "gimple_omp_single", GSS_OMP_SINGLE_LAYOUT) /* GIMPLE_OMP_TARGET represents - #pragma acc data + #pragma acc {kernels,parallel,data} #pragma omp target {,data,update} - BODY is the sequence of statements inside the target construct + BODY is the sequence of statements inside the construct (NULL for target update). CLAUSES is an OMP_CLAUSE chain holding the associated clauses. - CHILD_FN is set when outlining the body of the target region. + CHILD_FN is set when outlining the body of the offloaded region. All the statements in BODY are moved into this newly created function when converting OMP constructs into low-GIMPLE. DATA_ARG is a vec of 3 local variables in the parent function diff --git gcc/gimple.h gcc/gimple.h index a91bd4e..60c2469 100644 --- gcc/gimple.h +++ gcc/gimple.h @@ -108,9 +108,11 @@ enum gf_mask { GF_OMP_TARGET_KIND_REGION = 0, GF_OMP_TARGET_KIND_DATA = 1, GF_OMP_TARGET_KIND_UPDATE = 2, - GF_OMP_TARGET_KIND_OACC_DATA = 3, - GF_OMP_TARGET_KIND_OACC_UPDATE = 4, - GF_OMP_TARGET_KIND_OACC_ENTER_EXIT_DATA = 5, + GF_OMP_TARGET_KIND_OACC_PARALLEL = 3, + GF_OMP_TARGET_KIND_OACC_KERNELS = 4, + GF_OMP_TARGET_KIND_OACC_DATA = 5, + GF_OMP_TARGET_KIND_OACC_UPDATE = 6, + GF_OMP_TARGET_KIND_OACC_ENTER_EXIT_DATA = 7, /* True on an GIMPLE_OMP_RETURN statement if the return does not require a thread synchronization via some sort of barrier. The exact barrier @@ -560,8 +562,8 @@ struct GTY((tag("GSS_OMP_FOR"))) }; -/* GIMPLE_OACC_KERNELS, GIMPLE_OACC_PARALLEL, GIMPLE_OMP_PARALLEL, - GIMPLE_OMP_TARGET, GIMPLE_OMP_TASK */ +/* GIMPLE_OMP_PARALLEL, GIMPLE_OMP_TARGET, GIMPLE_OMP_TASK */ + struct GTY((tag("GSS_OMP_PARALLEL_LAYOUT"))) gimple_statement_omp_parallel_layout : public gimple_statement_omp { @@ -580,22 +582,6 @@ struct GTY((tag("GSS_OMP_PARALLEL_LAYOUT"))) tree data_arg; }; -/* GIMPLE_OACC_KERNELS */ -struct GTY((tag("GSS_OMP_PARALLEL_LAYOUT"))) - gimple_statement_oacc_kernels : public gimple_statement_omp_parallel_layout -{ - /* No extra fields; adds invariant: - stmt->code == GIMPLE_OACC_KERNELS. */ -}; - -/* GIMPLE_OACC_PARALLEL */ -struct GTY((tag("GSS_OMP_PARALLEL_LAYOUT"))) - gimple_statement_oacc_parallel : public gimple_statement_omp_parallel_layout -{ - /* No extra fields; adds invariant: - stmt->code == GIMPLE_OACC_PARALLEL. */ -}; - /* GIMPLE_OMP_PARALLEL or GIMPLE_TASK */ struct GTY((tag("GSS_OMP_PARALLEL_LAYOUT"))) gimple_statement_omp_taskreg : public gimple_statement_omp_parallel_layout @@ -913,22 +899,6 @@ is_a_helper ::test (gimple gs) template <> template <> inline bool -is_a_helper ::test (gimple gs) -{ - return gs->code == GIMPLE_OACC_KERNELS; -} - -template <> -template <> -inline bool -is_a_helper ::test (gimple gs) -{ - return gs->code == GIMPLE_OACC_PARALLEL; -} - -template <> -template <> -inline bool is_a_helper ::test (gimple gs) { return gs->code == GIMPLE_OMP_PARALLEL || gs->code == GIMPLE_OMP_TASK; @@ -1121,22 +1091,6 @@ is_a_helper ::test (const_gimple gs) template <> template <> inline bool -is_a_helper ::test (const_gimple gs) -{ - return gs->code == GIMPLE_OACC_KERNELS; -} - -template <> -template <> -inline bool -is_a_helper ::test (const_gimple gs) -{ - return gs->code == GIMPLE_OACC_PARALLEL; -} - -template <> -template <> -inline bool is_a_helper ::test (const_gimple gs) { return gs->code == GIMPLE_OMP_PARALLEL || gs->code == GIMPLE_OMP_TASK; @@ -1260,8 +1214,6 @@ gimple gimple_build_debug_bind_stat (tree, tree, gimple MEM_STAT_DECL); gimple gimple_build_debug_source_bind_stat (tree, tree, gimple MEM_STAT_DECL); #define gimple_build_debug_source_bind(var,val,stmt) \ gimple_build_debug_source_bind_stat ((var), (val), (stmt) MEM_STAT_INFO) -gimple gimple_build_oacc_kernels (gimple_seq, tree); -gimple gimple_build_oacc_parallel (gimple_seq, tree); gimple gimple_build_omp_critical (gimple_seq, tree); gimple gimple_build_omp_for (gimple_seq, int, tree, size_t, gimple_seq); gimple gimple_build_omp_parallel (gimple_seq, tree, tree, tree); @@ -1500,8 +1452,6 @@ gimple_has_substatements (gimple g) case GIMPLE_EH_FILTER: case GIMPLE_EH_ELSE: case GIMPLE_TRY: - case GIMPLE_OACC_KERNELS: - case GIMPLE_OACC_PARALLEL: case GIMPLE_OMP_FOR: case GIMPLE_OMP_MASTER: case GIMPLE_OMP_TASKGROUP: @@ -4362,197 +4312,6 @@ gimple_omp_set_body (gimple gs, gimple_seq body) } -/* Return the clauses associated with OACC_KERNELS statement GS. */ - -static inline tree -gimple_oacc_kernels_clauses (const_gimple gs) -{ - const gimple_statement_oacc_kernels *oacc_kernels_stmt = - as_a (gs); - return oacc_kernels_stmt->clauses; -} - -/* Return a pointer to the clauses associated with OACC_KERNELS statement GS. */ - -static inline tree * -gimple_oacc_kernels_clauses_ptr (gimple gs) -{ - gimple_statement_oacc_kernels *oacc_kernels_stmt = - as_a (gs); - return &oacc_kernels_stmt->clauses; -} - -/* Set CLAUSES to be the list of clauses associated with OACC_KERNELS statement - GS. */ - -static inline void -gimple_oacc_kernels_set_clauses (gimple gs, tree clauses) -{ - gimple_statement_oacc_kernels *oacc_kernels_stmt = - as_a (gs); - oacc_kernels_stmt->clauses = clauses; -} - -/* Return the child function used to hold the body of OACC_KERNELS statement - GS. */ - -static inline tree -gimple_oacc_kernels_child_fn (const_gimple gs) -{ - const gimple_statement_oacc_kernels *oacc_kernels_stmt = - as_a (gs); - return oacc_kernels_stmt->child_fn; -} - -/* Return a pointer to the child function used to hold the body of OACC_KERNELS - statement GS. */ - -static inline tree * -gimple_oacc_kernels_child_fn_ptr (gimple gs) -{ - gimple_statement_oacc_kernels *oacc_kernels_stmt = - as_a (gs); - return &oacc_kernels_stmt->child_fn; -} - -/* Set CHILD_FN to be the child function for OACC_KERNELS statement GS. */ - -static inline void -gimple_oacc_kernels_set_child_fn (gimple gs, tree child_fn) -{ - gimple_statement_oacc_kernels *oacc_kernels_stmt = - as_a (gs); - oacc_kernels_stmt->child_fn = child_fn; -} - -/* Return the artificial argument used to send variables and values - from the parent to the children threads in OACC_KERNELS statement GS. */ - -static inline tree -gimple_oacc_kernels_data_arg (const_gimple gs) -{ - const gimple_statement_oacc_kernels *oacc_kernels_stmt = - as_a (gs); - return oacc_kernels_stmt->data_arg; -} - -/* Return a pointer to the data argument for OACC_KERNELS statement GS. */ - -static inline tree * -gimple_oacc_kernels_data_arg_ptr (gimple gs) -{ - gimple_statement_oacc_kernels *oacc_kernels_stmt = - as_a (gs); - return &oacc_kernels_stmt->data_arg; -} - -/* Set DATA_ARG to be the data argument for OACC_KERNELS statement GS. */ - -static inline void -gimple_oacc_kernels_set_data_arg (gimple gs, tree data_arg) -{ - gimple_statement_oacc_kernels *oacc_kernels_stmt = - as_a (gs); - oacc_kernels_stmt->data_arg = data_arg; -} - - -/* Return the clauses associated with OACC_PARALLEL statement GS. */ - -static inline tree -gimple_oacc_parallel_clauses (const_gimple gs) -{ - const gimple_statement_oacc_parallel *oacc_parallel_stmt = - as_a (gs); - return oacc_parallel_stmt->clauses; -} - -/* Return a pointer to the clauses associated with OACC_PARALLEL statement - GS. */ - -static inline tree * -gimple_oacc_parallel_clauses_ptr (gimple gs) -{ - gimple_statement_oacc_parallel *oacc_parallel_stmt = - as_a (gs); - return &oacc_parallel_stmt->clauses; -} - -/* Set CLAUSES to be the list of clauses associated with OACC_PARALLEL - statement GS. */ - -static inline void -gimple_oacc_parallel_set_clauses (gimple gs, tree clauses) -{ - gimple_statement_oacc_parallel *oacc_parallel_stmt = - as_a (gs); - oacc_parallel_stmt->clauses = clauses; -} - -/* Return the child function used to hold the body of OACC_PARALLEL statement - GS. */ - -static inline tree -gimple_oacc_parallel_child_fn (const_gimple gs) -{ - const gimple_statement_oacc_parallel *oacc_parallel_stmt = - as_a (gs); - return oacc_parallel_stmt->child_fn; -} - -/* Return a pointer to the child function used to hold the body of - OACC_PARALLEL statement GS. */ - -static inline tree * -gimple_oacc_parallel_child_fn_ptr (gimple gs) -{ - gimple_statement_oacc_parallel *oacc_parallel_stmt = - as_a (gs); - return &oacc_parallel_stmt->child_fn; -} - -/* Set CHILD_FN to be the child function for OACC_PARALLEL statement GS. */ - -static inline void -gimple_oacc_parallel_set_child_fn (gimple gs, tree child_fn) -{ - gimple_statement_oacc_parallel *oacc_parallel_stmt = - as_a (gs); - oacc_parallel_stmt->child_fn = child_fn; -} - -/* Return the artificial argument used to send variables and values - from the parent to the children threads in OACC_PARALLEL statement GS. */ - -static inline tree -gimple_oacc_parallel_data_arg (const_gimple gs) -{ - const gimple_statement_oacc_parallel *oacc_parallel_stmt = - as_a (gs); - return oacc_parallel_stmt->data_arg; -} - -/* Return a pointer to the data argument for OACC_PARALLEL statement GS. */ - -static inline tree * -gimple_oacc_parallel_data_arg_ptr (gimple gs) -{ - gimple_statement_oacc_parallel *oacc_parallel_stmt = - as_a (gs); - return &oacc_parallel_stmt->data_arg; -} - -/* Set DATA_ARG to be the data argument for OACC_PARALLEL statement GS. */ - -static inline void -gimple_oacc_parallel_set_data_arg (gimple gs, tree data_arg) -{ - gimple_statement_oacc_parallel *oacc_parallel_stmt = - as_a (gs); - oacc_parallel_stmt->data_arg = data_arg; -} - - /* Return the name associated with OMP_CRITICAL statement GS. */ static inline tree @@ -5374,7 +5133,7 @@ gimple_omp_target_set_clauses (gimple gs, tree clauses) } -/* Return the kind of OMP target statemement. */ +/* Return the kind of the OMP_TARGET G. */ static inline int gimple_omp_target_kind (const_gimple g) @@ -5384,7 +5143,7 @@ gimple_omp_target_kind (const_gimple g) } -/* Set the OMP target kind. */ +/* Set the kind of the OMP_TARGET G. */ static inline void gimple_omp_target_set_kind (gimple g, int kind) @@ -5854,8 +5613,6 @@ gimple_return_set_retbnd (gimple gs, tree retval) /* Returns true when the gimple statement STMT is any of the OpenMP types. */ #define CASE_GIMPLE_OMP \ - case GIMPLE_OACC_KERNELS: \ - case GIMPLE_OACC_PARALLEL: \ case GIMPLE_OMP_PARALLEL: \ case GIMPLE_OMP_TASK: \ case GIMPLE_OMP_FOR: \ @@ -5898,9 +5655,6 @@ is_gimple_omp_oacc_specifically (const_gimple stmt) gcc_assert (is_gimple_omp (stmt)); switch (gimple_code (stmt)) { - case GIMPLE_OACC_KERNELS: - case GIMPLE_OACC_PARALLEL: - return true; case GIMPLE_OMP_FOR: switch (gimple_omp_for_kind (stmt)) { @@ -5908,10 +5662,12 @@ is_gimple_omp_oacc_specifically (const_gimple stmt) return true; default: return false; - } + } case GIMPLE_OMP_TARGET: switch (gimple_omp_target_kind (stmt)) { + case GF_OMP_TARGET_KIND_OACC_PARALLEL: + case GF_OMP_TARGET_KIND_OACC_KERNELS: case GF_OMP_TARGET_KIND_OACC_DATA: case GF_OMP_TARGET_KIND_OACC_UPDATE: case GF_OMP_TARGET_KIND_OACC_ENTER_EXIT_DATA: @@ -5933,13 +5689,12 @@ is_gimple_omp_offloaded (const_gimple stmt) gcc_assert (is_gimple_omp (stmt)); switch (gimple_code (stmt)) { - case GIMPLE_OACC_KERNELS: - case GIMPLE_OACC_PARALLEL: - return true; case GIMPLE_OMP_TARGET: switch (gimple_omp_target_kind (stmt)) { case GF_OMP_TARGET_KIND_REGION: + case GF_OMP_TARGET_KIND_OACC_PARALLEL: + case GF_OMP_TARGET_KIND_OACC_KERNELS: return true; default: return false; diff --git gcc/gimplify.c gcc/gimplify.c index ad48d51..eb9d930 100644 --- gcc/gimplify.c +++ gcc/gimplify.c @@ -7315,10 +7315,12 @@ gimplify_omp_workshare (tree *expr_p, gimple_seq *pre_p) OACC_DATA_CLAUSES (expr)); break; case OACC_KERNELS: - stmt = gimple_build_oacc_kernels (body, OACC_KERNELS_CLAUSES (expr)); + stmt = gimple_build_omp_target (body, GF_OMP_TARGET_KIND_OACC_KERNELS, + OMP_CLAUSES (expr)); break; case OACC_PARALLEL: - stmt = gimple_build_oacc_parallel (body, OACC_PARALLEL_CLAUSES (expr)); + stmt = gimple_build_omp_target (body, GF_OMP_TARGET_KIND_OACC_PARALLEL, + OMP_CLAUSES (expr)); break; case OMP_SECTIONS: stmt = gimple_build_omp_sections (body, OMP_CLAUSES (expr)); diff --git gcc/omp-low.c gcc/omp-low.c index 6fed38f..39e2f22 100644 --- gcc/omp-low.c +++ gcc/omp-low.c @@ -272,10 +272,12 @@ oacc_max_threads (omp_context *ctx) Scan for the innermost vector_length clause. */ for (omp_context *oc = ctx; oc; oc = oc->outer) { - if (gimple_code (oc->stmt) != GIMPLE_OACC_PARALLEL) + if (gimple_code (oc->stmt) != GIMPLE_OMP_TARGET + || (gimple_omp_target_kind (oc->stmt) + != GF_OMP_TARGET_KIND_OACC_PARALLEL)) continue; - clauses = gimple_oacc_parallel_clauses (oc->stmt); + clauses = gimple_omp_target_clauses (oc->stmt); vector_length = find_omp_clause (clauses, OMP_CLAUSE_VECTOR_LENGTH); if (vector_length) @@ -2643,28 +2645,7 @@ scan_omp_target (gimple stmt, omp_context *outer_ctx) { omp_context *ctx; tree name; - bool offloaded; - void (*gimple_omp_set_child_fn) (gimple, tree); - tree (*gimple_omp_clauses) (const_gimple); - - offloaded = is_gimple_omp_offloaded (stmt); - switch (gimple_code (stmt)) - { - case GIMPLE_OACC_KERNELS: - gimple_omp_set_child_fn = gimple_oacc_kernels_set_child_fn; - gimple_omp_clauses = gimple_oacc_kernels_clauses; - break; - case GIMPLE_OACC_PARALLEL: - gimple_omp_set_child_fn = gimple_oacc_parallel_set_child_fn; - gimple_omp_clauses = gimple_oacc_parallel_clauses; - break; - case GIMPLE_OMP_TARGET: - gimple_omp_set_child_fn = gimple_omp_target_set_child_fn; - gimple_omp_clauses = gimple_omp_target_clauses; - break; - default: - gcc_unreachable (); - } + bool offloaded = is_gimple_omp_offloaded (stmt); if (is_gimple_omp_oacc_specifically (stmt)) { @@ -2689,10 +2670,10 @@ scan_omp_target (gimple stmt, omp_context *outer_ctx) 0, 0); create_omp_child_function (ctx, false); - gimple_omp_set_child_fn (stmt, ctx->cb.dst_fn); + gimple_omp_target_set_child_fn (stmt, ctx->cb.dst_fn); } - scan_sharing_clauses (gimple_omp_clauses (stmt), ctx); + scan_sharing_clauses (gimple_omp_target_clauses (stmt), ctx); scan_omp (gimple_omp_body_ptr (stmt), ctx); if (TYPE_FIELDS (ctx->record_type) == NULL) @@ -2944,8 +2925,6 @@ check_omp_nesting_restrictions (gimple stmt, omp_context *ctx) "of work-sharing, critical, ordered, master or explicit " "task region"); return false; - case GIMPLE_OACC_KERNELS: - case GIMPLE_OACC_PARALLEL: case GIMPLE_OMP_PARALLEL: return true; default: @@ -3203,8 +3182,6 @@ scan_omp_1_stmt (gimple_stmt_iterator *gsi, bool *handled_ops_p, scan_omp (gimple_omp_body_ptr (stmt), ctx); break; - case GIMPLE_OACC_KERNELS: - case GIMPLE_OACC_PARALLEL: case GIMPLE_OMP_TARGET: scan_omp_target (stmt, ctx); break; @@ -8780,42 +8757,24 @@ expand_omp_target (struct omp_region *region) gimple entry_stmt, stmt; edge e; bool offloaded, data_region; - tree (*gimple_omp_child_fn) (const_gimple); - tree (*gimple_omp_data_arg) (const_gimple); entry_stmt = last_stmt (region->entry); new_bb = region->entry; offloaded = is_gimple_omp_offloaded (entry_stmt); - data_region = false; - switch (region->type) + switch (gimple_omp_target_kind (entry_stmt)) { - case GIMPLE_OACC_KERNELS: - gimple_omp_child_fn = gimple_oacc_kernels_child_fn; - gimple_omp_data_arg = gimple_oacc_kernels_data_arg; + case GF_OMP_TARGET_KIND_REGION: + case GF_OMP_TARGET_KIND_UPDATE: + case GF_OMP_TARGET_KIND_OACC_PARALLEL: + case GF_OMP_TARGET_KIND_OACC_KERNELS: + case GF_OMP_TARGET_KIND_OACC_UPDATE: + case GF_OMP_TARGET_KIND_OACC_ENTER_EXIT_DATA: + data_region = false; break; - case GIMPLE_OACC_PARALLEL: - gimple_omp_child_fn = gimple_oacc_parallel_child_fn; - gimple_omp_data_arg = gimple_oacc_parallel_data_arg; - break; - case GIMPLE_OMP_TARGET: - switch (gimple_omp_target_kind (entry_stmt)) - { - case GF_OMP_TARGET_KIND_DATA: - case GF_OMP_TARGET_KIND_OACC_DATA: - data_region = true; - break; - case GF_OMP_TARGET_KIND_REGION: - case GF_OMP_TARGET_KIND_UPDATE: - case GF_OMP_TARGET_KIND_OACC_UPDATE: - case GF_OMP_TARGET_KIND_OACC_ENTER_EXIT_DATA: - break; - default: - gcc_unreachable (); - } - - gimple_omp_child_fn = gimple_omp_target_child_fn; - gimple_omp_data_arg = gimple_omp_target_data_arg; + case GF_OMP_TARGET_KIND_DATA: + case GF_OMP_TARGET_KIND_OACC_DATA: + data_region = true; break; default: gcc_unreachable (); @@ -8825,7 +8784,7 @@ expand_omp_target (struct omp_region *region) child_cfun = NULL; if (offloaded) { - child_fn = gimple_omp_child_fn (entry_stmt); + child_fn = gimple_omp_target_child_fn (entry_stmt); child_cfun = DECL_STRUCT_FUNCTION (child_fn); } @@ -8854,13 +8813,14 @@ expand_omp_target (struct omp_region *region) a function call that has been inlined, the original PARM_DECL .OMP_DATA_I may have been converted into a different local variable. In which case, we need to keep the assignment. */ - if (gimple_omp_data_arg (entry_stmt)) + tree data_arg = gimple_omp_target_data_arg (entry_stmt); + if (data_arg) { basic_block entry_succ_bb = single_succ (entry_bb); gimple_stmt_iterator gsi; tree arg; gimple tgtcopy_stmt = NULL; - tree sender = TREE_VEC_ELT (gimple_omp_data_arg (entry_stmt), 0); + tree sender = TREE_VEC_ELT (data_arg, 0); for (gsi = gsi_start_bb (entry_succ_bb); ; gsi_next (&gsi)) { @@ -8994,49 +8954,38 @@ expand_omp_target (struct omp_region *region) tree t1, t2, t3, t4, device, cond, c, clauses; enum built_in_function start_ix; location_t clause_loc; - tree (*gimple_omp_clauses) (const_gimple); - switch (region->type) + switch (gimple_omp_target_kind (entry_stmt)) { - case GIMPLE_OACC_KERNELS: - gimple_omp_clauses = gimple_oacc_kernels_clauses; - start_ix = BUILT_IN_GOACC_KERNELS; + case GF_OMP_TARGET_KIND_REGION: + start_ix = BUILT_IN_GOMP_TARGET; break; - case GIMPLE_OACC_PARALLEL: - gimple_omp_clauses = gimple_oacc_parallel_clauses; + case GF_OMP_TARGET_KIND_DATA: + start_ix = BUILT_IN_GOMP_TARGET_DATA; + break; + case GF_OMP_TARGET_KIND_UPDATE: + start_ix = BUILT_IN_GOMP_TARGET_UPDATE; + break; + case GF_OMP_TARGET_KIND_OACC_PARALLEL: start_ix = BUILT_IN_GOACC_PARALLEL; break; - case GIMPLE_OMP_TARGET: - gimple_omp_clauses = gimple_omp_target_clauses; - switch (gimple_omp_target_kind (entry_stmt)) - { - case GF_OMP_TARGET_KIND_REGION: - start_ix = BUILT_IN_GOMP_TARGET; - break; - case GF_OMP_TARGET_KIND_DATA: - start_ix = BUILT_IN_GOMP_TARGET_DATA; - break; - case GF_OMP_TARGET_KIND_UPDATE: - start_ix = BUILT_IN_GOMP_TARGET_UPDATE; - break; - case GF_OMP_TARGET_KIND_OACC_DATA: - start_ix = BUILT_IN_GOACC_DATA_START; - break; - case GF_OMP_TARGET_KIND_OACC_ENTER_EXIT_DATA: - start_ix = BUILT_IN_GOACC_ENTER_EXIT_DATA; - break; - case GF_OMP_TARGET_KIND_OACC_UPDATE: - start_ix = BUILT_IN_GOACC_UPDATE; - break; - default: - gcc_unreachable (); - } + case GF_OMP_TARGET_KIND_OACC_KERNELS: + start_ix = BUILT_IN_GOACC_KERNELS; + break; + case GF_OMP_TARGET_KIND_OACC_DATA: + start_ix = BUILT_IN_GOACC_DATA_START; + break; + case GF_OMP_TARGET_KIND_OACC_UPDATE: + start_ix = BUILT_IN_GOACC_UPDATE; + break; + case GF_OMP_TARGET_KIND_OACC_ENTER_EXIT_DATA: + start_ix = BUILT_IN_GOACC_ENTER_EXIT_DATA; break; default: gcc_unreachable (); } - clauses = gimple_omp_clauses (entry_stmt); + clauses = gimple_omp_target_clauses (entry_stmt); /* By default, the value of DEVICE is -1 (let runtime library choose) and there is no conditional. */ @@ -9119,7 +9068,7 @@ expand_omp_target (struct omp_region *region) } gsi = gsi_last_bb (new_bb); - t = gimple_omp_data_arg (entry_stmt); + t = gimple_omp_target_data_arg (entry_stmt); if (t == NULL) { t1 = size_zero_node; @@ -9331,8 +9280,6 @@ expand_omp (struct omp_region *region) expand_omp_atomic (region); break; - case GIMPLE_OACC_KERNELS: - case GIMPLE_OACC_PARALLEL: case GIMPLE_OMP_TARGET: expand_omp_target (region); break; @@ -9679,22 +9626,18 @@ oacc_initialize_reduction_data (tree clauses, tree nthreads, tree c, t, oc; gimple stmt; omp_context *octx; - tree (*gimple_omp_clauses) (const_gimple); - void (*gimple_omp_set_clauses) (gimple, tree); /* Find the innermost PARALLEL openmp context. FIXME: OpenACC kernels may require extra care unless they are converted to openmp for loops. */ - - if (gimple_code (ctx->stmt) == GIMPLE_OACC_PARALLEL) + if (gimple_code (ctx->stmt) == GIMPLE_OMP_TARGET + && (gimple_omp_target_kind (ctx->stmt) + == GF_OMP_TARGET_KIND_OACC_PARALLEL)) octx = ctx; else octx = ctx->outer; - gimple_omp_clauses = gimple_oacc_parallel_clauses; - gimple_omp_set_clauses = gimple_oacc_parallel_set_clauses; - /* Extract the clauses. */ - oc = gimple_omp_clauses (octx->stmt); + oc = gimple_omp_target_clauses (octx->stmt); /* Find the last outer clause. */ for (; oc && OMP_CLAUSE_CHAIN (oc); oc = OMP_CLAUSE_CHAIN (oc)) @@ -9744,7 +9687,7 @@ oacc_initialize_reduction_data (tree clauses, tree nthreads, if (oc) OMP_CLAUSE_CHAIN (oc) = t; else - gimple_omp_set_clauses (octx->stmt, t); + gimple_omp_target_set_clauses (octx->stmt, t); OMP_CLAUSE_SIZE (t) = size; oc = t; } @@ -11195,45 +11138,27 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx) location_t loc = gimple_location (stmt); bool offloaded, data_region; unsigned int map_cnt = 0; - tree (*gimple_omp_clauses) (const_gimple); - void (*gimple_omp_set_data_arg) (gimple, tree); offloaded = is_gimple_omp_offloaded (stmt); - data_region = false; - switch (gimple_code (stmt)) + switch (gimple_omp_target_kind (stmt)) { - case GIMPLE_OACC_KERNELS: - gimple_omp_clauses = gimple_oacc_kernels_clauses; - gimple_omp_set_data_arg = gimple_oacc_kernels_set_data_arg; + case GF_OMP_TARGET_KIND_REGION: + case GF_OMP_TARGET_KIND_UPDATE: + case GF_OMP_TARGET_KIND_OACC_PARALLEL: + case GF_OMP_TARGET_KIND_OACC_KERNELS: + case GF_OMP_TARGET_KIND_OACC_UPDATE: + case GF_OMP_TARGET_KIND_OACC_ENTER_EXIT_DATA: + data_region = false; break; - case GIMPLE_OACC_PARALLEL: - gimple_omp_clauses = gimple_oacc_parallel_clauses; - gimple_omp_set_data_arg = gimple_oacc_parallel_set_data_arg; - break; - case GIMPLE_OMP_TARGET: - switch (gimple_omp_target_kind (stmt)) - { - case GF_OMP_TARGET_KIND_DATA: - case GF_OMP_TARGET_KIND_OACC_DATA: - data_region = true; - break; - case GF_OMP_TARGET_KIND_REGION: - case GF_OMP_TARGET_KIND_UPDATE: - case GF_OMP_TARGET_KIND_OACC_UPDATE: - case GF_OMP_TARGET_KIND_OACC_ENTER_EXIT_DATA: - break; - default: - gcc_unreachable (); - } - - gimple_omp_clauses = gimple_omp_target_clauses; - gimple_omp_set_data_arg = gimple_omp_target_set_data_arg; + case GF_OMP_TARGET_KIND_DATA: + case GF_OMP_TARGET_KIND_OACC_DATA: + data_region = true; break; default: gcc_unreachable (); } - clauses = gimple_omp_clauses (stmt); + clauses = gimple_omp_target_clauses (stmt); tgt_bind = NULL; tgt_body = NULL; @@ -11250,15 +11175,9 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx) irlist = NULL; orlist = NULL; - switch (gimple_code (stmt)) - { - case GIMPLE_OACC_KERNELS: - case GIMPLE_OACC_PARALLEL: - oacc_process_reduction_data (&tgt_body, &irlist, &orlist, ctx); - break; - default: - break; - } + if (offloaded + && is_gimple_omp_oacc_specifically (stmt)) + oacc_process_reduction_data (&tgt_body, &irlist, &orlist, ctx); for (c = clauses; c ; c = OMP_CLAUSE_CHAIN (c)) switch (OMP_CLAUSE_CODE (c)) @@ -11390,7 +11309,7 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx) DECL_NAMELESS (TREE_VEC_ELT (t, 2)) = 1; TREE_ADDRESSABLE (TREE_VEC_ELT (t, 2)) = 1; TREE_STATIC (TREE_VEC_ELT (t, 2)) = 1; - gimple_omp_set_data_arg (stmt, t); + gimple_omp_target_set_data_arg (stmt, t); vec *vsize; vec *vkind; @@ -11457,8 +11376,8 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx) || TREE_CODE (TREE_TYPE (ovar)) != ARRAY_TYPE); if (maybe_lookup_oacc_reduction (var, ctx)) { - gcc_assert (gimple_code (stmt) == GIMPLE_OACC_KERNELS - || gimple_code (stmt) == GIMPLE_OACC_PARALLEL); + gcc_assert (offloaded + && is_gimple_omp_oacc_specifically (stmt)); gimplify_assign (x, var, &ilist); } else if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP @@ -11802,8 +11721,6 @@ lower_omp_1 (gimple_stmt_iterator *gsi_p, omp_context *ctx) lower_omp_regimplify_p, ctx ? NULL : &wi, NULL)) gimple_regimplify_operands (stmt, gsi_p); break; - case GIMPLE_OACC_KERNELS: - case GIMPLE_OACC_PARALLEL: case GIMPLE_OMP_TARGET: ctx = maybe_lookup_ctx (stmt); gcc_assert (ctx); @@ -12095,8 +12012,6 @@ diagnose_sb_1 (gimple_stmt_iterator *gsi_p, bool *handled_ops_p, { WALK_SUBSTMTS; - case GIMPLE_OACC_KERNELS: - case GIMPLE_OACC_PARALLEL: case GIMPLE_OMP_PARALLEL: case GIMPLE_OMP_TASK: case GIMPLE_OMP_SECTIONS: @@ -12155,8 +12070,6 @@ diagnose_sb_2 (gimple_stmt_iterator *gsi_p, bool *handled_ops_p, { WALK_SUBSTMTS; - case GIMPLE_OACC_KERNELS: - case GIMPLE_OACC_PARALLEL: case GIMPLE_OMP_PARALLEL: case GIMPLE_OMP_TASK: case GIMPLE_OMP_SECTIONS: @@ -12252,8 +12165,6 @@ make_gimple_omp_edges (basic_block bb, struct omp_region **region, switch (code) { - case GIMPLE_OACC_KERNELS: - case GIMPLE_OACC_PARALLEL: case GIMPLE_OMP_PARALLEL: case GIMPLE_OMP_TASK: case GIMPLE_OMP_FOR: diff --git gcc/testsuite/gfortran.dg/goacc/private-1.f95 gcc/testsuite/gfortran.dg/goacc/private-1.f95 index 54c027d..23ce95a 100644 --- gcc/testsuite/gfortran.dg/goacc/private-1.f95 +++ gcc/testsuite/gfortran.dg/goacc/private-1.f95 @@ -31,7 +31,7 @@ program test end do !$acc end parallel end program test -! { dg-final { scan-tree-dump-times "pragma acc parallel" 3 "omplower" } } +! { dg-final { scan-tree-dump-times "pragma omp target oacc_parallel" 3 "omplower" } } ! { dg-final { scan-tree-dump-times "private\\(i\\)" 3 "omplower" } } ! { dg-final { scan-tree-dump-times "private\\(j\\)" 2 "omplower" } } ! { dg-final { scan-tree-dump-times "private\\(k\\)" 1 "omplower" } } diff --git gcc/tree-inline.c gcc/tree-inline.c index 54b3514..89cb1eb 100644 --- gcc/tree-inline.c +++ gcc/tree-inline.c @@ -1395,10 +1395,6 @@ remap_gimple_stmt (gimple stmt, copy_body_data *id) copy = gimple_build_wce (s1); break; - case GIMPLE_OACC_KERNELS: - case GIMPLE_OACC_PARALLEL: - gcc_unreachable (); - case GIMPLE_OMP_PARALLEL: s1 = remap_gimple_seq (gimple_omp_body (stmt), id); copy = gimple_build_omp_parallel @@ -4112,8 +4108,6 @@ estimate_num_insns (gimple stmt, eni_weights *weights) + estimate_num_insns_seq (gimple_omp_body (stmt), weights) + estimate_num_insns_seq (gimple_omp_for_pre_body (stmt), weights)); - case GIMPLE_OACC_KERNELS: - case GIMPLE_OACC_PARALLEL: case GIMPLE_OMP_PARALLEL: case GIMPLE_OMP_TASK: case GIMPLE_OMP_CRITICAL: diff --git gcc/tree-nested.c gcc/tree-nested.c index b5d6543..4da6297 100644 --- gcc/tree-nested.c +++ gcc/tree-nested.c @@ -1325,10 +1325,6 @@ convert_nonlocal_reference_stmt (gimple_stmt_iterator *gsi, bool *handled_ops_p, } break; - case GIMPLE_OACC_KERNELS: - case GIMPLE_OACC_PARALLEL: - gcc_unreachable (); - case GIMPLE_OMP_PARALLEL: case GIMPLE_OMP_TASK: save_suppress = info->suppress_expansion; @@ -1898,10 +1894,6 @@ convert_local_reference_stmt (gimple_stmt_iterator *gsi, bool *handled_ops_p, switch (gimple_code (stmt)) { - case GIMPLE_OACC_KERNELS: - case GIMPLE_OACC_PARALLEL: - gcc_unreachable (); - case GIMPLE_OMP_PARALLEL: case GIMPLE_OMP_TASK: save_suppress = info->suppress_expansion; @@ -2289,10 +2281,6 @@ convert_tramp_reference_stmt (gimple_stmt_iterator *gsi, bool *handled_ops_p, break; } - case GIMPLE_OACC_KERNELS: - case GIMPLE_OACC_PARALLEL: - gcc_unreachable (); - case GIMPLE_OMP_TARGET: gcc_assert (!is_gimple_omp_oacc_specifically (stmt)); if (gimple_omp_target_kind (stmt) != GF_OMP_TARGET_KIND_REGION) @@ -2360,10 +2348,6 @@ convert_gimple_call (gimple_stmt_iterator *gsi, bool *handled_ops_p, } break; - case GIMPLE_OACC_KERNELS: - case GIMPLE_OACC_PARALLEL: - gcc_unreachable (); - case GIMPLE_OMP_PARALLEL: case GIMPLE_OMP_TASK: save_static_chain_added = info->static_chain_added;