From patchwork Wed Nov 6 19:53:00 2013 Content-Type: text/plain; charset="utf-8" MIME-Version: 1.0 Content-Transfer-Encoding: 8bit X-Patchwork-Submitter: Thomas Schwinge X-Patchwork-Id: 288998 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)) (Client did not present a certificate) by ozlabs.org (Postfix) with ESMTPS id D8BB92C00A8 for ; Thu, 7 Nov 2013 06:53:47 +1100 (EST) 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:subject:in-reply-to:references:date:message-id:mime-version :content-type; q=dns; s=default; b=lLoeMjsz1q4LJoNFL/udXV0uFxIft yX9eSAJGG5/7Eznohfct874sWabpvTmfFLOp0PqHMRUrGrOXCROo4kIC8wRM4Ng/ pgzxXz759kuHd/tTS/Vr5usFFbuiq7bZiewLZ+VGd1jmsA/IruaVgYgotvwB3N7X mlz+kYID2OCqBQ= 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:subject:in-reply-to:references:date:message-id:mime-version :content-type; s=default; bh=WYHtrSra6xy1suyIsf45sgbdGiU=; b=i0C SBzAaCX5Dfb/37QT/1NgDBtgwpyAjyQa1ohhSP18aLBvgiX3FCQ2tBScHkcXduXv xKEMMWIH8fnYuh+SG8IvOpd6jLfvS+jxGwyxozLJ9HL+9IETNOnlxJLBLua9Si1T 9NjZqwwCvMVsWgVgO2kSJjWcKCzl9oYFTozyNdHI= Received: (qmail 15481 invoked by alias); 6 Nov 2013 19:53:36 -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 15455 invoked by uid 89); 6 Nov 2013 19:53:36 -0000 Authentication-Results: sourceware.org; auth=none X-Virus-Found: No X-Spam-SWARE-Status: No, score=-0.5 required=5.0 tests=AWL, BAYES_50, RDNS_NONE, URIBL_BLOCKED autolearn=no version=3.3.2 X-HELO: relay1.mentorg.com Received: from Unknown (HELO relay1.mentorg.com) (192.94.38.131) by sourceware.org (qpsmtpd/0.93/v0.84-503-g423c35a) with ESMTP; Wed, 06 Nov 2013 19:53:31 +0000 Received: from svr-orw-exc-10.mgc.mentorg.com ([147.34.98.58]) by relay1.mentorg.com with esmtp id 1Ve9AB-0002GV-MQ from Thomas_Schwinge@mentor.com for gcc-patches@gcc.gnu.org; Wed, 06 Nov 2013 11:53:15 -0800 Received: from SVR-IES-FEM-01.mgc.mentorg.com ([137.202.0.104]) by SVR-ORW-EXC-10.mgc.mentorg.com with Microsoft SMTPSVC(6.0.3790.4675); Wed, 6 Nov 2013 11:53:15 -0800 Received: from feldtkeller.schwinge.homeip.net (137.202.0.76) by SVR-IES-FEM-01.mgc.mentorg.com (137.202.0.104) with Microsoft SMTP Server id 14.2.247.3; Wed, 6 Nov 2013 19:53:11 +0000 From: Thomas Schwinge To: Subject: Re: [gomp4 9/9] OpenACC: Basic support for #pragma acc parallel. In-Reply-To: <1383766943-8863-9-git-send-email-thomas@codesourcery.com> References: <878ux1jp2s.fsf@schwinge.name> <1383766943-8863-1-git-send-email-thomas@codesourcery.com> <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> User-Agent: Notmuch/0.9-101-g81dad07 (http://notmuchmail.org) Emacs/23.4.1 (x86_64-pc-linux-gnu) Date: Wed, 6 Nov 2013 20:53:00 +0100 Message-ID: <8761s5joir.fsf@schwinge.name> MIME-Version: 1.0 Hi! On Wed, 6 Nov 2013 20:42:23 +0100, I wrote: > diff --git gcc/omp-low.c gcc/omp-low.c > index 99811d0..84fe466 100644 > --- gcc/omp-low.c > +++ gcc/omp-low.c > [...] > @@ -4542,10 +4672,10 @@ expand_omp_build_assign (gimple_stmt_iterator *gsi_p, tree to, tree from) > } > } > > -/* Expand the OpenMP parallel or task directive starting at REGION. */ > +/* Expand the OpenACC parallel directive starting at REGION. */ > > static void > -expand_omp_taskreg (struct omp_region *region) > +expand_oacc_parallel (struct omp_region *region) > { > basic_block entry_bb, exit_bb, new_bb; > struct function *child_cfun; > [...] > +/* Expand the OpenMP parallel or task directive starting at REGION. */ > + > +static void > +expand_omp_taskreg (struct omp_region *region) Forgot to pass the --patience switch to Git, so the diff algorithm decided to first patch the existing expand_omp_taskreg into expand_oacc_parallel, and then later re-add expand_omp_taskreg. Here's a more readable version of that patch, avoiding all that back'n'forth: gcc/c-family/ * c-pragma.h (pragma_kind): Add PRAGMA_OACC_PARALLEL. * c-pragma.c (oacc_pragmas): Add "parallel". gcc/c/ * c-parser.c (c_parser_omp_structured_block): Update comment. (c_parser_oacc_parallel): New function. (c_parser_omp_construct): Handle PRAGMA_OACC_PARALLEL. gcc/ * tree.def (OACC_PARALLEL): New code. * doc/generic.texi (OpenMP): Document it. * tree.h (OMP_BODY, OMP_CLAUSES): Include it. (OACC_PARALLEL_BODY, OACC_PARALLEL_CLAUSES): New macros. * tree-pretty-print.c (dump_generic_node): Handle OACC_PARALLEL. gcc/c/ * c-tree.h (c_finish_oacc_parallel): New declaration. * c-typeck.c (c_finish_oacc_parallel): New function. gcc/c-family/ * c-omp.c (c_omp_split_clauses): Catch OACC_PARALLEL. gcc/ * gimple.def (GIMPLE_OACC_PARALLEL): New code. * doc/gimple.texi: Document it. * gimple.h (gimple_build_oacc_parallel): New declaration. (gimple_oacc_parallel_clauses, gimple_oacc_parallel_clauses_ptr) (gimple_oacc_parallel_set_clauses, gimple_oacc_parallel_child_fn) (gimple_oacc_parallel_child_fn_ptr) (gimple_oacc_parallel_set_child_fn, gimple_oacc_parallel_data_arg) (gimple_oacc_parallel_data_arg_ptr) (gimple_oacc_parallel_set_data_arg): New inline functions. (CASE_GIMPLE_OMP): Add GIMPLE_OACC_PARALLEL. * gimple.c (gimple_build_oacc_parallel): New function. (walk_gimple_op, walk_gimple_stmt, gimple_copy): Handle GIMPLE_OACC_PARALLEL. * gimplify.c (is_gimple_stmt): Handle GIMPLE_OACC_PARALLEL. (gimplify_oacc_parallel): New function. (gimplify_expr): Handle OACC_PARALLEL. * cgraphbuild.c (build_cgraph_edges): Handle GIMPLE_OACC_PARALLEL. * gimple-low.c (lower_stmt): Likewise. * gimple-pretty-print.c (pp_gimple_stmt_1): Likewise. (dump_gimple_oacc_parallel): New function. * oacc-builtins.def (BUILT_IN_GOACC_PARALLEL): New macro. * omp-low.c (scan_oacc_parallel, expand_oacc_parallel) (lower_oacc_parallel): New functions. (use_pointer_for_field, build_outer_var_ref, scan_sharing_clauses) (create_omp_child_function, check_omp_nesting_restrictions) (scan_omp_1_stmt, lower_rec_simd_input_clauses) (lower_lastprivate_clauses, lower_reduction_clauses) (lower_copyprivate_clauses, lower_send_clauses) (lower_send_shared_vars, expand_omp) (maybe_add_implicit_barrier_cancel, create_task_copyfn) (lower_omp_1, make_gimple_omp_edges): Handle GIMPLE_OACC_PARALLEL, or catch it. * tree-inline.c (remap_gimple_stmt): Likewise. * tree-nested.c (convert_nonlocal_reference_stmt) (convert_local_reference_stmt, convert_tramp_reference_stmt) (convert_gimple_call): Likewise. gcc/testsuite/ * c-c++-common/goacc-gomp/nesting-fail-1.c: New file. * c-c++-common/goacc/nesting-fail-1.c: Likewise. * c-c++-common/goacc/parallel-1.c: Likewise. * c-c++-common/goacc/parallel-fail-1.c: Likewise. libgomp/ * oacc-parallel.c: New file. * Makefile.am (libgomp_la_SOURCES): Add it. * Makefile.in: Regenerate. * libgomp.map (GOACC_2.0): Add GOACC_parallel. * libgomp_g.h (GOACC_parallel): New declaration. * testsuite/libgomp.oacc-c/goacc_parallel.c: New file. * testsuite/libgomp.oacc-c/parallel-1.c: New file. --- gcc/c-family/c-omp.c | 1 + gcc/c-family/c-pragma.c | 1 + gcc/c-family/c-pragma.h | 1 + gcc/c/c-parser.c | 42 +- gcc/c/c-tree.h | 1 + gcc/c/c-typeck.c | 19 + gcc/cgraphbuild.c | 12 +- gcc/doc/generic.texi | 5 + gcc/doc/gimple.texi | 8 + gcc/gimple-low.c | 1 + gcc/gimple-pretty-print.c | 58 ++ gcc/gimple.c | 36 ++ gcc/gimple.def | 10 +- gcc/gimple.h | 89 +++ gcc/gimplify.c | 38 ++ gcc/oacc-builtins.def | 3 + gcc/omp-low.c | 661 ++++++++++++++++++++- .../c-c++-common/goacc-gomp/nesting-fail-1.c | 121 ++++ gcc/testsuite/c-c++-common/goacc/nesting-fail-1.c | 11 + gcc/testsuite/c-c++-common/goacc/parallel-1.c | 6 + gcc/testsuite/c-c++-common/goacc/parallel-fail-1.c | 6 + gcc/tree-inline.c | 4 + gcc/tree-nested.c | 12 + gcc/tree-pretty-print.c | 5 + gcc/tree.def | 11 +- gcc/tree.h | 9 +- libgomp/Makefile.am | 2 +- libgomp/Makefile.in | 5 +- libgomp/libgomp.map | 2 + libgomp/libgomp_g.h | 5 + libgomp/oacc-parallel.c | 36 ++ libgomp/testsuite/libgomp.oacc-c/goacc_parallel.c | 25 + libgomp/testsuite/libgomp.oacc-c/parallel-1.c | 26 + 33 files changed, 1257 insertions(+), 15 deletions(-) create mode 100644 gcc/testsuite/c-c++-common/goacc-gomp/nesting-fail-1.c create mode 100644 gcc/testsuite/c-c++-common/goacc/nesting-fail-1.c create mode 100644 gcc/testsuite/c-c++-common/goacc/parallel-1.c create mode 100644 gcc/testsuite/c-c++-common/goacc/parallel-fail-1.c create mode 100644 libgomp/oacc-parallel.c create mode 100644 libgomp/testsuite/libgomp.oacc-c/goacc_parallel.c create mode 100644 libgomp/testsuite/libgomp.oacc-c/parallel-1.c Grüße, Thomas diff --git gcc/c-family/c-omp.c gcc/c-family/c-omp.c index f001a75..f7d2bd9 100644 --- gcc/c-family/c-omp.c +++ gcc/c-family/c-omp.c @@ -627,6 +627,7 @@ c_omp_split_clauses (location_t loc, enum tree_code code, enum c_omp_clause_split s; int i; + gcc_assert (code != OACC_PARALLEL); for (i = 0; i < C_OMP_CLAUSE_SPLIT_COUNT; i++) cclauses[i] = NULL; /* Add implicit nowait clause on diff --git gcc/c-family/c-pragma.c gcc/c-family/c-pragma.c index 98f98d0..c329f8d 100644 --- gcc/c-family/c-pragma.c +++ gcc/c-family/c-pragma.c @@ -1165,6 +1165,7 @@ static vec registered_pp_pragmas; struct omp_pragma_def { const char *name; unsigned int id; }; static const struct omp_pragma_def oacc_pragmas[] = { + { "parallel", PRAGMA_OACC_PARALLEL }, }; static const struct omp_pragma_def omp_pragmas[] = { { "atomic", PRAGMA_OMP_ATOMIC }, diff --git gcc/c-family/c-pragma.h gcc/c-family/c-pragma.h index 705bcb4..5c58e32 100644 --- gcc/c-family/c-pragma.h +++ gcc/c-family/c-pragma.h @@ -27,6 +27,7 @@ along with GCC; see the file COPYING3. If not see typedef enum pragma_kind { PRAGMA_NONE = 0, + PRAGMA_OACC_PARALLEL, PRAGMA_OMP_ATOMIC, PRAGMA_OMP_BARRIER, PRAGMA_OMP_CANCEL, diff --git gcc/c/c-parser.c gcc/c/c-parser.c index 8a1e988..297b6da7 100644 --- gcc/c/c-parser.c +++ gcc/c/c-parser.c @@ -4478,6 +4478,17 @@ c_parser_label (c_parser *parser) @throw expression ; @throw ; + OpenACC: + + statement: + openacc-construct + + openacc-construct: + parallel-construct + + parallel-construct: + parallel-directive structured-block + OpenMP: statement: @@ -10754,7 +10765,7 @@ c_parser_omp_all_clauses (c_parser *parser, omp_clause_mask mask, return clauses; } -/* OpenMP 2.5: +/* OpenACC 2.0, OpenMP 2.5: structured-block: statement @@ -10770,6 +10781,32 @@ c_parser_omp_structured_block (c_parser *parser) return pop_stmt_list (stmt); } +/* OpenACC 2.0: + # pragma acc parallel oacc-parallel-clause[optseq] new-line + + LOC is the location of the #pragma token. +*/ + +#define OACC_PARALLEL_CLAUSE_MASK \ + PRAGMA_OMP_CLAUSE_NONE + +static tree +c_parser_oacc_parallel (location_t loc, c_parser *parser) +{ + tree stmt, clauses, block; + + clauses = c_parser_omp_all_clauses (parser, OACC_PARALLEL_CLAUSE_MASK, + "#pragma acc parallel"); + gcc_assert (clauses == NULL); + + block = c_begin_omp_parallel (); + add_stmt (c_parser_omp_structured_block (parser)); + + stmt = c_finish_oacc_parallel (loc, clauses, block); + + return stmt; +} + /* OpenMP 2.5: # pragma omp atomic new-line expression-stmt @@ -12948,6 +12985,9 @@ c_parser_omp_construct (c_parser *parser) switch (p_kind) { + case PRAGMA_OACC_PARALLEL: + stmt = c_parser_oacc_parallel (loc, parser); + break; case PRAGMA_OMP_ATOMIC: c_parser_omp_atomic (loc, parser); return; diff --git gcc/c/c-tree.h gcc/c/c-tree.h index 2565ccb..f524e31 100644 --- gcc/c/c-tree.h +++ gcc/c/c-tree.h @@ -635,6 +635,7 @@ extern tree c_finish_bc_stmt (location_t, tree *, bool); extern tree c_finish_goto_label (location_t, tree); extern tree c_finish_goto_ptr (location_t, tree); extern tree c_expr_to_decl (tree, bool *, bool *); +extern tree c_finish_oacc_parallel (location_t, tree, tree); extern tree c_begin_omp_parallel (void); extern tree c_finish_omp_parallel (location_t, tree, tree); extern tree c_begin_omp_task (void); diff --git gcc/c/c-typeck.c gcc/c/c-typeck.c index 8f1d3a4..e7096e6 100644 --- gcc/c/c-typeck.c +++ gcc/c/c-typeck.c @@ -10644,6 +10644,25 @@ c_expr_to_decl (tree expr, bool *tc ATTRIBUTE_UNUSED, bool *se) return expr; } +/* Generate OACC_PARALLEL, with CLAUSES and BLOCK as its compound + statement. LOC is the location of the OACC_PARALLEL. */ + +tree +c_finish_oacc_parallel (location_t loc, tree clauses, tree block) +{ + tree stmt; + + block = c_end_compound_stmt (loc, block, true); + + stmt = make_node (OACC_PARALLEL); + TREE_TYPE (stmt) = void_type_node; + OACC_PARALLEL_CLAUSES (stmt) = clauses; + OACC_PARALLEL_BODY (stmt) = block; + SET_EXPR_LOCATION (stmt, loc); + + return add_stmt (stmt); +} + /* Like c_begin_compound_stmt, except force the retention of the BLOCK. */ tree diff --git gcc/cgraphbuild.c gcc/cgraphbuild.c index 87e06e3..efad3d9 100644 --- gcc/cgraphbuild.c +++ gcc/cgraphbuild.c @@ -333,7 +333,15 @@ build_cgraph_edges (void) bb->count, freq); } ipa_record_stmt_references (node, stmt); - if (gimple_code (stmt) == GIMPLE_OMP_PARALLEL + if (gimple_code (stmt) == GIMPLE_OACC_PARALLEL + && gimple_oacc_parallel_child_fn (stmt)) + { + tree fn = gimple_oacc_parallel_child_fn (stmt); + ipa_record_reference (node, + cgraph_get_create_real_symbol_node (fn), + IPA_REF_ADDR, stmt); + } + else if (gimple_code (stmt) == GIMPLE_OMP_PARALLEL && gimple_omp_parallel_child_fn (stmt)) { tree fn = gimple_omp_parallel_child_fn (stmt); @@ -341,7 +349,7 @@ build_cgraph_edges (void) cgraph_get_create_real_symbol_node (fn), IPA_REF_ADDR, stmt); } - if (gimple_code (stmt) == GIMPLE_OMP_TASK) + else if (gimple_code (stmt) == GIMPLE_OMP_TASK) { tree fn = gimple_omp_task_child_fn (stmt); if (fn) diff --git gcc/doc/generic.texi gcc/doc/generic.texi index 73dd123..812f5a9 100644 --- gcc/doc/generic.texi +++ gcc/doc/generic.texi @@ -2049,6 +2049,7 @@ edge. Rethrowing the exception is represented using @code{RESX_EXPR}. @node OpenMP @subsection OpenMP +@tindex OACC_PARALLEL @tindex OMP_PARALLEL @tindex OMP_FOR @tindex OMP_SECTIONS @@ -2066,6 +2067,10 @@ All the statements starting with @code{OMP_} represent directives and clauses used by the OpenMP API @w{@uref{http://www.openmp.org/}}. @table @code +@item OACC_PARALLEL + +Represents @code{#pragma acc parallel [clause1 @dots{} clauseN]}. + @item OMP_PARALLEL Represents @code{#pragma omp parallel [clause1 @dots{} clauseN]}. It diff --git gcc/doc/gimple.texi gcc/doc/gimple.texi index 7bd9fd5..0f1bbe6 100644 --- gcc/doc/gimple.texi +++ gcc/doc/gimple.texi @@ -338,6 +338,7 @@ 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_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 @@ -905,6 +906,7 @@ Return a deep copy of statement @code{STMT}. * @code{GIMPLE_EH_FILTER}:: * @code{GIMPLE_LABEL}:: * @code{GIMPLE_NOP}:: +* @code{GIMPLE_OACC_PARALLEL}:: * @code{GIMPLE_OMP_ATOMIC_LOAD}:: * @code{GIMPLE_OMP_ATOMIC_STORE}:: * @code{GIMPLE_OMP_CONTINUE}:: @@ -1554,6 +1556,12 @@ Build a @code{GIMPLE_NOP} statement. Returns @code{TRUE} if statement @code{G} is a @code{GIMPLE_NOP}. @end deftypefn + +@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 d527d86..74c9925 100644 --- gcc/gimple-low.c +++ gcc/gimple-low.c @@ -368,6 +368,7 @@ lower_stmt (gimple_stmt_iterator *gsi, struct lower_data *data) } break; + 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 6842213..59cb5bb 100644 --- gcc/gimple-pretty-print.c +++ gcc/gimple-pretty-print.c @@ -1823,6 +1823,60 @@ dump_gimple_phi (pretty_printer *buffer, gimple phi, int spc, bool comment, } +/* Dump a GIMPLE_OACC_PARALLEL 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_parallel (pretty_printer *buffer, gimple gs, int spc, + int flags) +{ + if (flags & TDF_RAW) + { + dump_gimple_fmt (buffer, spc, flags, "%G <%+BODY <%S>%nCLAUSES <", gs, + gimple_omp_body (gs)); + dump_omp_clauses (buffer, gimple_oacc_parallel_clauses (gs), spc, flags); + dump_gimple_fmt (buffer, spc, flags, " >, %T, %T%n>", + gimple_oacc_parallel_child_fn (gs), + gimple_oacc_parallel_data_arg (gs)); + } + else + { + gimple_seq body; + pp_string (buffer, "#pragma acc parallel"); + dump_omp_clauses (buffer, gimple_oacc_parallel_clauses (gs), spc, flags); + if (gimple_oacc_parallel_child_fn (gs)) + { + pp_string (buffer, " [child fn: "); + dump_generic_node (buffer, gimple_oacc_parallel_child_fn (gs), + spc, flags, false); + pp_string (buffer, " ("); + if (gimple_oacc_parallel_data_arg (gs)) + dump_generic_node (buffer, gimple_oacc_parallel_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). */ @@ -2123,6 +2177,10 @@ 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_PARALLEL: + dump_gimple_oacc_parallel (buffer, gs, spc, flags); + break; + case GIMPLE_OMP_PARALLEL: dump_gimple_omp_parallel (buffer, gs, spc, flags); break; diff --git gcc/gimple.c gcc/gimple.c index 20f6010..ea96d26 100644 --- gcc/gimple.c +++ gcc/gimple.c @@ -898,6 +898,23 @@ gimple_build_debug_source_bind_stat (tree var, tree value, } +/* 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. @@ -1571,6 +1588,21 @@ walk_gimple_op (gimple stmt, walk_tree_fn callback_op, 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); @@ -1866,6 +1898,7 @@ walk_gimple_stmt (gimple_stmt_iterator *gsi, walk_stmt_fn callback_stmt, return wi->callback_result; /* FALL THROUGH. */ + case GIMPLE_OACC_PARALLEL: case GIMPLE_OMP_CRITICAL: case GIMPLE_OMP_MASTER: case GIMPLE_OMP_TASKGROUP: @@ -2306,6 +2339,9 @@ gimple_copy (gimple stmt) gimple_try_set_cleanup (copy, new_seq); break; + case GIMPLE_OACC_PARALLEL: + abort (); + case GIMPLE_OMP_FOR: new_seq = gimple_seq_copy (gimple_omp_for_pre_body (stmt)); gimple_omp_for_set_pre_body (copy, new_seq); diff --git gcc/gimple.def gcc/gimple.def index 07370ae..9ff9ab3 100644 --- gcc/gimple.def +++ gcc/gimple.def @@ -205,10 +205,16 @@ DEFGSCODE(GIMPLE_NOP, "gimple_nop", GSS_BASE) /* IMPORTANT. - Do not rearrange any of the GIMPLE_OMP_* codes. This ordering is - exposed by the range check in gimple_omp_subcode(). */ + 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_PARALLEL represents + + #pragma acc parallel [CLAUSES] + BODY */ +DEFGSCODE(GIMPLE_OACC_PARALLEL, "gimple_oacc_parallel", GSS_OMP_PARALLEL) + /* 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 conversions may cause the expression to become more complex, so that it does diff --git gcc/gimple.h gcc/gimple.h index b34424c..c9be1c9 100644 --- gcc/gimple.h +++ gcc/gimple.h @@ -786,6 +786,7 @@ gimple gimple_build_resx (int); gimple gimple_build_eh_dispatch (int); gimple gimple_build_switch_nlabels (unsigned, tree, tree); gimple gimple_build_switch (tree, tree, vec ); +gimple gimple_build_oacc_parallel (gimple_seq, tree); gimple gimple_build_omp_parallel (gimple_seq, tree, tree, tree); gimple gimple_build_omp_task (gimple_seq, tree, tree, tree, tree, tree, tree); gimple gimple_build_omp_for (gimple_seq, int, tree, size_t, gimple_seq); @@ -1256,6 +1257,7 @@ gimple_has_substatements (gimple g) case GIMPLE_EH_FILTER: case GIMPLE_EH_ELSE: case GIMPLE_TRY: + case GIMPLE_OACC_PARALLEL: case GIMPLE_OMP_FOR: case GIMPLE_OMP_MASTER: case GIMPLE_OMP_TASKGROUP: @@ -4061,6 +4063,92 @@ gimple_omp_set_body (gimple gs, gimple_seq body) } +/* Return the clauses associated with OACC_PARALLEL statement GS. */ + +static inline tree +gimple_oacc_parallel_clauses (const_gimple gs) +{ + GIMPLE_CHECK (gs, GIMPLE_OACC_PARALLEL); + return gs->gimple_omp_parallel.clauses; +} + +/* Return a pointer to the clauses associated with OACC_PARALLEL statement + GS. */ + +static inline tree * +gimple_oacc_parallel_clauses_ptr (gimple gs) +{ + GIMPLE_CHECK (gs, GIMPLE_OACC_PARALLEL); + return &gs->gimple_omp_parallel.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_CHECK (gs, GIMPLE_OACC_PARALLEL); + gs->gimple_omp_parallel.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) +{ + GIMPLE_CHECK (gs, GIMPLE_OACC_PARALLEL); + return gs->gimple_omp_parallel.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_CHECK (gs, GIMPLE_OACC_PARALLEL); + return &gs->gimple_omp_parallel.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_CHECK (gs, GIMPLE_OACC_PARALLEL); + gs->gimple_omp_parallel.child_fn = child_fn; +} + +/* Return the data argument for OACC_PARALLEL statement GS. */ + +static inline tree +gimple_oacc_parallel_data_arg (const_gimple gs) +{ + GIMPLE_CHECK (gs, GIMPLE_OACC_PARALLEL); + return gs->gimple_omp_parallel.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_CHECK (gs, GIMPLE_OACC_PARALLEL); + return &gs->gimple_omp_parallel.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_CHECK (gs, GIMPLE_OACC_PARALLEL); + gs->gimple_omp_parallel.data_arg = data_arg; +} + + /* Return the name associated with OMP_CRITICAL statement GS. */ static inline tree @@ -5269,6 +5357,7 @@ 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_PARALLEL: \ case GIMPLE_OMP_PARALLEL: \ case GIMPLE_OMP_TASK: \ case GIMPLE_OMP_FOR: \ diff --git gcc/gimplify.c gcc/gimplify.c index 30c2b45..0c45729 100644 --- gcc/gimplify.c +++ gcc/gimplify.c @@ -4641,6 +4641,7 @@ is_gimple_stmt (tree t) case CATCH_EXPR: case ASM_EXPR: case STATEMENT_LIST: + case OACC_PARALLEL: case OMP_PARALLEL: case OMP_FOR: case OMP_SIMD: @@ -6745,6 +6746,37 @@ gimplify_adjust_omp_clauses (tree *list_p) delete_omp_context (ctx); } +/* Gimplify the contents of an OACC_PARALLEL statement. This involves + gimplification of the body, as well as scanning the body for used + variables. We need to do this scan now, because variable-sized + decls will be decomposed during gimplification. */ + +static void +gimplify_oacc_parallel (tree *expr_p, gimple_seq *pre_p) +{ + tree expr = *expr_p; + gimple g; + gimple_seq body = NULL; + struct gimplify_ctx gctx; + + gimplify_scan_omp_clauses (&OACC_PARALLEL_CLAUSES (expr), pre_p, + ORT_TARGET); + + push_gimplify_context (&gctx); + + g = gimplify_and_return_first (OACC_PARALLEL_BODY (expr), &body); + if (gimple_code (g) == GIMPLE_BIND) + pop_gimplify_context (g); + else + pop_gimplify_context (NULL); + + gimplify_adjust_omp_clauses (&OACC_PARALLEL_CLAUSES (expr)); + + g = gimple_build_oacc_parallel (body, OACC_PARALLEL_CLAUSES (expr)); + gimplify_seq_add_stmt (pre_p, g); + *expr_p = NULL_TREE; +} + /* Gimplify the contents of an OMP_PARALLEL statement. This involves gimplification of the body, as well as scanning the body for used variables. We need to do this scan now, because variable-sized @@ -8169,6 +8201,11 @@ gimplify_expr (tree *expr_p, gimple_seq *pre_p, gimple_seq *post_p, ret = GS_ALL_DONE; break; + case OACC_PARALLEL: + gimplify_oacc_parallel (expr_p, pre_p); + ret = GS_ALL_DONE; + break; + case OMP_PARALLEL: gimplify_omp_parallel (expr_p, pre_p); ret = GS_ALL_DONE; @@ -8575,6 +8612,7 @@ gimplify_expr (tree *expr_p, gimple_seq *pre_p, gimple_seq *post_p, && code != LOOP_EXPR && code != SWITCH_EXPR && code != TRY_FINALLY_EXPR + && code != OACC_PARALLEL && code != OMP_CRITICAL && code != OMP_FOR && code != OMP_MASTER diff --git gcc/oacc-builtins.def gcc/oacc-builtins.def index fd630e0..a75e42d 100644 --- gcc/oacc-builtins.def +++ gcc/oacc-builtins.def @@ -26,3 +26,6 @@ along with GCC; see the file COPYING3. If not see DEF_GOACC_BUILTIN (ENUM, NAME, TYPE, ATTRS) See builtins.def for details. */ + +DEF_GOACC_BUILTIN (BUILT_IN_GOACC_PARALLEL, "GOACC_parallel", + BT_FN_VOID_INT_OMPFN_PTR_SIZE_PTR_PTR_PTR, ATTR_NOTHROW_LIST) diff --git gcc/omp-low.c gcc/omp-low.c index 99811d0..84fe466 100644 --- gcc/omp-low.c +++ gcc/omp-low.c @@ -844,6 +844,8 @@ use_pointer_for_field (tree decl, omp_context *shared_ctx) when we know the value is not accessible from an outer scope. */ if (shared_ctx) { + gcc_assert (gimple_code (shared_ctx->stmt) != GIMPLE_OACC_PARALLEL); + /* ??? Trivially accessible from anywhere. But why would we even be passing an address in this case? Should we simply assert this to be false, or should we have a cleanup pass that removes @@ -985,6 +987,8 @@ build_receiver_ref (tree var, bool by_ref, omp_context *ctx) static tree build_outer_var_ref (tree var, omp_context *ctx) { + gcc_assert (gimple_code (ctx->stmt) != GIMPLE_OACC_PARALLEL); + tree x; if (is_global_var (maybe_lookup_decl_in_outer_ctx (var, ctx))) @@ -1484,6 +1488,7 @@ scan_sharing_clauses (tree clauses, omp_context *ctx) switch (OMP_CLAUSE_CODE (c)) { case OMP_CLAUSE_PRIVATE: + gcc_assert (gimple_code (ctx->stmt) != GIMPLE_OACC_PARALLEL); decl = OMP_CLAUSE_DECL (c); if (OMP_CLAUSE_PRIVATE_OUTER_REF (c)) goto do_private; @@ -1492,6 +1497,7 @@ scan_sharing_clauses (tree clauses, omp_context *ctx) break; case OMP_CLAUSE_SHARED: + gcc_assert (gimple_code (ctx->stmt) != GIMPLE_OACC_PARALLEL); /* Ignore shared directives in teams construct. */ if (gimple_code (ctx->stmt) == GIMPLE_OMP_TEAMS) break; @@ -1518,6 +1524,7 @@ scan_sharing_clauses (tree clauses, omp_context *ctx) goto do_private; case OMP_CLAUSE_LASTPRIVATE: + gcc_assert (gimple_code (ctx->stmt) != GIMPLE_OACC_PARALLEL); /* Let the corresponding firstprivate clause create the variable. */ if (OMP_CLAUSE_LASTPRIVATE_FIRSTPRIVATE (c)) @@ -1527,6 +1534,7 @@ scan_sharing_clauses (tree clauses, omp_context *ctx) case OMP_CLAUSE_FIRSTPRIVATE: case OMP_CLAUSE_REDUCTION: case OMP_CLAUSE_LINEAR: + gcc_assert (gimple_code (ctx->stmt) != GIMPLE_OACC_PARALLEL); decl = OMP_CLAUSE_DECL (c); do_private: if (is_variable_sized (decl)) @@ -1555,6 +1563,7 @@ scan_sharing_clauses (tree clauses, omp_context *ctx) break; case OMP_CLAUSE__LOOPTEMP_: + gcc_assert (gimple_code (ctx->stmt) != GIMPLE_OACC_PARALLEL); gcc_assert (is_parallel_ctx (ctx)); decl = OMP_CLAUSE_DECL (c); install_var_field (decl, false, 3, ctx); @@ -1563,12 +1572,14 @@ scan_sharing_clauses (tree clauses, omp_context *ctx) case OMP_CLAUSE_COPYPRIVATE: case OMP_CLAUSE_COPYIN: + gcc_assert (gimple_code (ctx->stmt) != GIMPLE_OACC_PARALLEL); decl = OMP_CLAUSE_DECL (c); by_ref = use_pointer_for_field (decl, NULL); install_var_field (decl, by_ref, 3, ctx); break; case OMP_CLAUSE_DEFAULT: + gcc_assert (gimple_code (ctx->stmt) != GIMPLE_OACC_PARALLEL); ctx->default_kind = OMP_CLAUSE_DEFAULT_KIND (c); break; @@ -1581,6 +1592,7 @@ scan_sharing_clauses (tree clauses, omp_context *ctx) case OMP_CLAUSE_SCHEDULE: case OMP_CLAUSE_DIST_SCHEDULE: case OMP_CLAUSE_DEPEND: + gcc_assert (gimple_code (ctx->stmt) != GIMPLE_OACC_PARALLEL); if (ctx->outer) scan_omp_op (&OMP_CLAUSE_OPERAND (c, 0), ctx->outer); break; @@ -1599,10 +1611,14 @@ scan_sharing_clauses (tree clauses, omp_context *ctx) && is_global_var (maybe_lookup_decl_in_outer_ctx (decl, ctx)) && lookup_attribute ("omp declare target", DECL_ATTRIBUTES (decl))) + { + gcc_assert (gimple_code (ctx->stmt) != GIMPLE_OACC_PARALLEL); break; + } if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP && OMP_CLAUSE_MAP_KIND (c) == OMP_CLAUSE_MAP_POINTER) { + gcc_assert (gimple_code (ctx->stmt) != GIMPLE_OACC_PARALLEL); /* Ignore OMP_CLAUSE_MAP_POINTER kind for arrays in #pragma omp target data, there is nothing to map for those. */ @@ -1632,8 +1648,9 @@ scan_sharing_clauses (tree clauses, omp_context *ctx) install_var_field (decl, true, 7, ctx); else install_var_field (decl, true, 3, ctx); - if (gimple_omp_target_kind (ctx->stmt) - == GF_OMP_TARGET_KIND_REGION) + if (gimple_code (ctx->stmt) == GIMPLE_OACC_PARALLEL + || (gimple_omp_target_kind (ctx->stmt) + == GF_OMP_TARGET_KIND_REGION)) install_var_local (decl, ctx); } } @@ -1673,9 +1690,11 @@ scan_sharing_clauses (tree clauses, omp_context *ctx) case OMP_CLAUSE_MERGEABLE: case OMP_CLAUSE_PROC_BIND: case OMP_CLAUSE_SAFELEN: + gcc_assert (gimple_code (ctx->stmt) != GIMPLE_OACC_PARALLEL); break; case OMP_CLAUSE_ALIGNED: + gcc_assert (gimple_code (ctx->stmt) != GIMPLE_OACC_PARALLEL); decl = OMP_CLAUSE_DECL (c); if (is_global_var (decl) && TREE_CODE (TREE_TYPE (decl)) == ARRAY_TYPE) @@ -1692,6 +1711,7 @@ scan_sharing_clauses (tree clauses, omp_context *ctx) switch (OMP_CLAUSE_CODE (c)) { case OMP_CLAUSE_LASTPRIVATE: + gcc_assert (gimple_code (ctx->stmt) != GIMPLE_OACC_PARALLEL); /* Let the corresponding firstprivate clause create the variable. */ if (OMP_CLAUSE_LASTPRIVATE_GIMPLE_SEQ (c)) @@ -1704,6 +1724,7 @@ scan_sharing_clauses (tree clauses, omp_context *ctx) case OMP_CLAUSE_FIRSTPRIVATE: case OMP_CLAUSE_REDUCTION: case OMP_CLAUSE_LINEAR: + gcc_assert (gimple_code (ctx->stmt) != GIMPLE_OACC_PARALLEL); decl = OMP_CLAUSE_DECL (c); if (is_variable_sized (decl)) install_var_local (decl, ctx); @@ -1716,6 +1737,7 @@ scan_sharing_clauses (tree clauses, omp_context *ctx) break; case OMP_CLAUSE_SHARED: + gcc_assert (gimple_code (ctx->stmt) != GIMPLE_OACC_PARALLEL); /* Ignore shared directives in teams construct. */ if (gimple_code (ctx->stmt) == GIMPLE_OMP_TEAMS) break; @@ -1725,14 +1747,18 @@ scan_sharing_clauses (tree clauses, omp_context *ctx) break; case OMP_CLAUSE_MAP: - if (gimple_omp_target_kind (ctx->stmt) == GF_OMP_TARGET_KIND_DATA) + if (gimple_code (ctx->stmt) != GIMPLE_OACC_PARALLEL + && gimple_omp_target_kind (ctx->stmt) == GF_OMP_TARGET_KIND_DATA) break; decl = OMP_CLAUSE_DECL (c); if (DECL_P (decl) && is_global_var (maybe_lookup_decl_in_outer_ctx (decl, ctx)) && lookup_attribute ("omp declare target", DECL_ATTRIBUTES (decl))) + { + gcc_assert (gimple_code (ctx->stmt) != GIMPLE_OACC_PARALLEL); break; + } if (DECL_P (decl)) { if (OMP_CLAUSE_MAP_KIND (c) == OMP_CLAUSE_MAP_POINTER @@ -1781,6 +1807,7 @@ scan_sharing_clauses (tree clauses, omp_context *ctx) case OMP_CLAUSE__LOOPTEMP_: case OMP_CLAUSE_TO: case OMP_CLAUSE_FROM: + gcc_assert (gimple_code (ctx->stmt) != GIMPLE_OACC_PARALLEL); break; default: @@ -1789,6 +1816,8 @@ scan_sharing_clauses (tree clauses, omp_context *ctx) } if (scan_array_reductions) + { + gcc_assert (gimple_code (ctx->stmt) != GIMPLE_OACC_PARALLEL); for (c = clauses; c; c = OMP_CLAUSE_CHAIN (c)) if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_REDUCTION && OMP_CLAUSE_REDUCTION_PLACEHOLDER (c)) @@ -1799,6 +1828,7 @@ scan_sharing_clauses (tree clauses, omp_context *ctx) else if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_LASTPRIVATE && OMP_CLAUSE_LASTPRIVATE_GIMPLE_SEQ (c)) scan_omp (&OMP_CLAUSE_LASTPRIVATE_GIMPLE_SEQ (c), ctx); + } } /* Create a new name for omp child function. Returns an identifier. */ @@ -1830,6 +1860,8 @@ create_omp_child_function (omp_context *ctx, bool task_copy) decl = build_decl (gimple_location (ctx->stmt), FUNCTION_DECL, name, type); + gcc_assert (gimple_code (ctx->stmt) != GIMPLE_OACC_PARALLEL + || !task_copy); if (!task_copy) ctx->cb.dst_fn = decl; else @@ -1861,6 +1893,8 @@ create_omp_child_function (omp_context *ctx, bool task_copy) break; } } + gcc_assert (gimple_code (ctx->stmt) != GIMPLE_OACC_PARALLEL + || !target_p); if (target_p) DECL_ATTRIBUTES (decl) = tree_cons (get_identifier ("omp declare target"), @@ -1935,6 +1969,52 @@ find_combined_for (gimple_stmt_iterator *gsi_p, return NULL; } +/* Scan an OpenACC parallel directive. */ + +static void +scan_oacc_parallel (gimple stmt, omp_context *outer_ctx) +{ + omp_context *ctx; + tree name; + + gcc_assert (taskreg_nesting_level == 0); + gcc_assert (target_nesting_level == 0); + + ctx = new_omp_context (stmt, outer_ctx); + ctx->field_map = splay_tree_new (splay_tree_compare_pointers, 0, 0); + ctx->default_kind = OMP_CLAUSE_DEFAULT_SHARED; + ctx->record_type = lang_hooks.types.make_type (RECORD_TYPE); + name = create_tmp_var_name (".omp_data_t"); + name = build_decl (gimple_location (stmt), + TYPE_DECL, name, ctx->record_type); + DECL_ARTIFICIAL (name) = 1; + DECL_NAMELESS (name) = 1; + TYPE_NAME (ctx->record_type) = name; + create_omp_child_function (ctx, false); + gimple_oacc_parallel_set_child_fn (stmt, ctx->cb.dst_fn); + + scan_sharing_clauses (gimple_oacc_parallel_clauses (stmt), ctx); + scan_omp (gimple_omp_body_ptr (stmt), ctx); + + if (TYPE_FIELDS (ctx->record_type) == NULL) + ctx->record_type = ctx->receiver_decl = NULL; + else + { + TYPE_FIELDS (ctx->record_type) + = nreverse (TYPE_FIELDS (ctx->record_type)); +#ifdef ENABLE_CHECKING + tree field; + unsigned int align = DECL_ALIGN (TYPE_FIELDS (ctx->record_type)); + for (field = TYPE_FIELDS (ctx->record_type); + field; + field = DECL_CHAIN (field)) + gcc_assert (DECL_ALIGN (field) == align); +#endif + layout_type (ctx->record_type); + fixup_child_record_type (ctx); + } +} + /* Scan an OpenMP parallel directive. */ static void @@ -2225,6 +2305,38 @@ scan_omp_teams (gimple stmt, omp_context *outer_ctx) static bool check_omp_nesting_restrictions (gimple stmt, omp_context *ctx) { + omp_context *ctx_; + + /* TODO: While the OpenACC specification does allow for certain kinds of + nesting, we don't support that yet. */ + /* No nesting of STMT (which is an OpenACC or OpenMP one, or a GOMP builtin) + inside any OpenACC CTX. */ + for (ctx_ = ctx; ctx_ != NULL; ctx_ = ctx_->outer) + switch (gimple_code (ctx_->stmt)) + { + case GIMPLE_OACC_PARALLEL: + error_at (gimple_location (stmt), + "may not be nested"); + return false; + default: + break; + } + /* No nesting of OpenACC STMT inside any OpenACC or OpenMP CTX. */ + switch (gimple_code (stmt)) + { + case GIMPLE_OACC_PARALLEL: + for (ctx_ = ctx; ctx_ != NULL; ctx_ = ctx_->outer) + if (is_gimple_omp (ctx_->stmt)) + { + error_at (gimple_location (stmt), + "may not be nested"); + return false; + } + break; + default: + break; + } + if (ctx != NULL) { if (gimple_code (ctx->stmt) == GIMPLE_OMP_FOR @@ -2584,6 +2696,10 @@ scan_omp_1_stmt (gimple_stmt_iterator *gsi, bool *handled_ops_p, switch (gimple_code (stmt)) { + case GIMPLE_OACC_PARALLEL: + scan_oacc_parallel (stmt, ctx); + break; + case GIMPLE_OMP_PARALLEL: taskreg_nesting_level++; scan_omp_parallel (gsi, ctx); @@ -2910,6 +3026,8 @@ static bool lower_rec_simd_input_clauses (tree new_var, omp_context *ctx, int &max_vf, tree &idx, tree &lane, tree &ivar, tree &lvar) { + gcc_assert (gimple_code (ctx->stmt) != GIMPLE_OACC_PARALLEL); + if (max_vf == 0) { max_vf = omp_max_vf (); @@ -2959,6 +3077,8 @@ static void lower_rec_input_clauses (tree clauses, gimple_seq *ilist, gimple_seq *dlist, omp_context *ctx, struct omp_for_data *fd) { + gcc_assert (gimple_code (ctx->stmt) != GIMPLE_OACC_PARALLEL); + tree c, dtor, copyin_seq, x, ptr; bool copyin_by_ref = false; bool lastprivate_firstprivate = false; @@ -3617,6 +3737,8 @@ static void lower_lastprivate_clauses (tree clauses, tree predicate, gimple_seq *stmt_list, omp_context *ctx) { + gcc_assert (gimple_code (ctx->stmt) != GIMPLE_OACC_PARALLEL); + tree x, c, label = NULL, orig_clauses = clauses; bool par_clauses = false; tree simduid = NULL, lastlane = NULL; @@ -3752,6 +3874,8 @@ lower_lastprivate_clauses (tree clauses, tree predicate, gimple_seq *stmt_list, static void lower_reduction_clauses (tree clauses, gimple_seq *stmt_seqp, omp_context *ctx) { + gcc_assert (gimple_code (ctx->stmt) != GIMPLE_OACC_PARALLEL); + gimple_seq sub_seq = NULL; gimple stmt; tree x, c; @@ -3853,6 +3977,8 @@ static void lower_copyprivate_clauses (tree clauses, gimple_seq *slist, gimple_seq *rlist, omp_context *ctx) { + gcc_assert (gimple_code (ctx->stmt) != GIMPLE_OACC_PARALLEL); + tree c; for (c = clauses; c ; c = OMP_CLAUSE_CHAIN (c)) @@ -3903,6 +4029,8 @@ static void lower_send_clauses (tree clauses, gimple_seq *ilist, gimple_seq *olist, omp_context *ctx) { + gcc_assert (gimple_code (ctx->stmt) != GIMPLE_OACC_PARALLEL); + tree c; for (c = clauses; c ; c = OMP_CLAUSE_CHAIN (c)) @@ -3994,6 +4122,8 @@ lower_send_clauses (tree clauses, gimple_seq *ilist, gimple_seq *olist, static void lower_send_shared_vars (gimple_seq *ilist, gimple_seq *olist, omp_context *ctx) { + gcc_assert (gimple_code (ctx->stmt) != GIMPLE_OACC_PARALLEL); + tree var, ovar, nvar, f, x, record_type; if (ctx->record_type == NULL) @@ -4542,6 +4672,234 @@ expand_omp_build_assign (gimple_stmt_iterator *gsi_p, tree to, tree from) } } +/* Expand the OpenACC parallel directive starting at REGION. */ + +static void +expand_oacc_parallel (struct omp_region *region) +{ + basic_block entry_bb, exit_bb, new_bb; + struct function *child_cfun; + tree child_fn, block, t; + gimple_stmt_iterator gsi; + gimple entry_stmt, stmt; + edge e; + + entry_stmt = last_stmt (region->entry); + child_fn = gimple_oacc_parallel_child_fn (entry_stmt); + child_cfun = DECL_STRUCT_FUNCTION (child_fn); + + /* Supported by expand_omp_taskreg, but not here. */ + gcc_assert (!child_cfun->cfg); + gcc_assert (!gimple_in_ssa_p (cfun)); + + entry_bb = region->entry; + exit_bb = region->exit; + + /* Preserve indentation of expand_omp_target and expand_omp_taskreg. */ + if (1) + { + unsigned srcidx, dstidx, num; + + /* If the parallel region needs data sent from the parent + function, then the very first statement (except possible + tree profile counter updates) of the parallel body + is a copy assignment .OMP_DATA_I = &.OMP_DATA_O. Since + &.OMP_DATA_O is passed as an argument to the child function, + we need to replace it with the argument as seen by the child + function. + + In most cases, this will end up being the identity assignment + .OMP_DATA_I = .OMP_DATA_I. However, if the parallel body had + 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_oacc_parallel_data_arg (entry_stmt)) + { + basic_block entry_succ_bb = single_succ (entry_bb); + gimple_stmt_iterator gsi; + tree arg; + gimple parcopy_stmt = NULL; + tree sender + = TREE_VEC_ELT (gimple_oacc_parallel_data_arg (entry_stmt), 0); + + for (gsi = gsi_start_bb (entry_succ_bb); ; gsi_next (&gsi)) + { + gcc_assert (!gsi_end_p (gsi)); + stmt = gsi_stmt (gsi); + if (gimple_code (stmt) != GIMPLE_ASSIGN) + continue; + + if (gimple_num_ops (stmt) == 2) + { + tree arg = gimple_assign_rhs1 (stmt); + + /* We're ignore the subcode because we're + effectively doing a STRIP_NOPS. */ + + if (TREE_CODE (arg) == ADDR_EXPR + && TREE_OPERAND (arg, 0) == sender) + { + parcopy_stmt = stmt; + break; + } + } + } + + gcc_assert (parcopy_stmt != NULL); + arg = DECL_ARGUMENTS (child_fn); + + gcc_assert (gimple_assign_lhs (parcopy_stmt) == arg); + gsi_remove (&gsi, true); + } + + /* Declare local variables needed in CHILD_CFUN. */ + block = DECL_INITIAL (child_fn); + BLOCK_VARS (block) = vec2chain (child_cfun->local_decls); + /* The gimplifier could record temporaries in the block + rather than in containing function's local_decls chain, + which would mean cgraph missed finalizing them. Do it now. */ + for (t = BLOCK_VARS (block); t; t = DECL_CHAIN (t)) + if (TREE_CODE (t) == VAR_DECL + && TREE_STATIC (t) + && !DECL_EXTERNAL (t)) + varpool_finalize_decl (t); + DECL_SAVED_TREE (child_fn) = NULL; + /* We'll create a CFG for child_fn, so no gimple body is needed. */ + gimple_set_body (child_fn, NULL); + TREE_USED (block) = 1; + + /* Reset DECL_CONTEXT on function arguments. */ + for (t = DECL_ARGUMENTS (child_fn); t; t = DECL_CHAIN (t)) + DECL_CONTEXT (t) = child_fn; + + /* Split ENTRY_BB at GIMPLE_OACC_PARALLEL, + so that it can be moved to the child function. */ + gsi = gsi_last_bb (entry_bb); + stmt = gsi_stmt (gsi); + gcc_assert (stmt && (gimple_code (stmt) == GIMPLE_OACC_PARALLEL)); + gsi_remove (&gsi, true); + e = split_block (entry_bb, stmt); + entry_bb = e->dest; + single_succ_edge (entry_bb)->flags = EDGE_FALLTHRU; + + /* Convert GIMPLE_OMP_RETURN into a RETURN_EXPR. */ + if (exit_bb) + { + gsi = gsi_last_bb (exit_bb); + gcc_assert (!gsi_end_p (gsi) + && gimple_code (gsi_stmt (gsi)) == GIMPLE_OMP_RETURN); + stmt = gimple_build_return (NULL); + gsi_insert_after (&gsi, stmt, GSI_SAME_STMT); + gsi_remove (&gsi, true); + } + + /* Move the region into CHILD_CFUN. */ + + block = gimple_block (entry_stmt); + + new_bb = move_sese_region_to_fn (child_cfun, entry_bb, exit_bb, block); + if (exit_bb) + single_succ_edge (new_bb)->flags = EDGE_FALLTHRU; + /* When the expansion process cannot guarantee an up-to-date + loop tree arrange for the child function to fixup loops. */ + if (loops_state_satisfies_p (LOOPS_NEED_FIXUP)) + child_cfun->x_current_loops->state |= LOOPS_NEED_FIXUP; + + /* Remove non-local VAR_DECLs from child_cfun->local_decls list. */ + num = vec_safe_length (child_cfun->local_decls); + for (srcidx = 0, dstidx = 0; srcidx < num; srcidx++) + { + t = (*child_cfun->local_decls)[srcidx]; + if (DECL_CONTEXT (t) == cfun->decl) + continue; + if (srcidx != dstidx) + (*child_cfun->local_decls)[dstidx] = t; + dstidx++; + } + if (dstidx != num) + vec_safe_truncate (child_cfun->local_decls, dstidx); + + /* Inform the callgraph about the new function. */ + DECL_STRUCT_FUNCTION (child_fn)->curr_properties = cfun->curr_properties; + cgraph_add_new_function (child_fn, true); + + /* Fix the callgraph edges for child_cfun. Those for cfun will be + fixed in a following pass. */ + push_cfun (child_cfun); + rebuild_cgraph_edges (); + + /* Some EH regions might become dead, see PR34608. If + pass_cleanup_cfg isn't the first pass to happen with the + new child, these dead EH edges might cause problems. + Clean them up now. */ + if (flag_exceptions) + { + basic_block bb; + bool changed = false; + + FOR_EACH_BB (bb) + changed |= gimple_purge_dead_eh_edges (bb); + if (changed) + cleanup_tree_cfg (); + } + pop_cfun (); + } + + /* Emit a library call to launch CHILD_FN. */ + tree t1, t2, t3, t4, device, c, clauses; + enum built_in_function start_ix; + location_t clause_loc; + + clauses = gimple_oacc_parallel_clauses (entry_stmt); + + start_ix = BUILT_IN_GOACC_PARALLEL; + + /* By default, the value of DEVICE is -1 (let runtime library choose). */ + device = build_int_cst (integer_type_node, -1); + + c = find_omp_clause (clauses, OMP_CLAUSE_DEVICE); + gcc_assert (c == NULL); + if (c) + { + device = OMP_CLAUSE_DEVICE_ID (c); + clause_loc = OMP_CLAUSE_LOCATION (c); + } + else + clause_loc = gimple_location (entry_stmt); + + /* Ensure 'device' is of the correct type. */ + device = fold_convert_loc (clause_loc, integer_type_node, device); + + gsi = gsi_last_bb (new_bb); + t = gimple_oacc_parallel_data_arg (entry_stmt); + if (t == NULL) + { + t1 = size_zero_node; + t2 = build_zero_cst (ptr_type_node); + t3 = t2; + t4 = t2; + } + else + { + t1 = TYPE_MAX_VALUE (TYPE_DOMAIN (TREE_TYPE (TREE_VEC_ELT (t, 1)))); + t1 = size_binop (PLUS_EXPR, t1, size_int (1)); + t2 = build_fold_addr_expr (TREE_VEC_ELT (t, 0)); + t3 = build_fold_addr_expr (TREE_VEC_ELT (t, 1)); + t4 = build_fold_addr_expr (TREE_VEC_ELT (t, 2)); + } + + gimple g; + /* FIXME: This will be address of + extern char __OPENMP_TARGET__[] __attribute__((visibility ("hidden"))) + symbol, as soon as the linker plugin is able to create it for us. */ + tree openmp_target = build_zero_cst (ptr_type_node); + tree fnaddr = build_fold_addr_expr (child_fn); + g = gimple_build_call (builtin_decl_explicit (start_ix), + 7, device, fnaddr, openmp_target, t1, t2, t3, t4); + gimple_set_location (g, gimple_location (entry_stmt)); + gsi_insert_before (&gsi, g, GSI_SAME_STMT); +} + /* Expand the OpenMP parallel or task directive starting at REGION. */ static void @@ -8037,6 +8395,10 @@ expand_omp (struct omp_region *region) switch (region->type) { + case GIMPLE_OACC_PARALLEL: + expand_oacc_parallel (region); + break; + case GIMPLE_OMP_PARALLEL: case GIMPLE_OMP_TASK: expand_omp_taskreg (region); @@ -8278,6 +8640,288 @@ make_pass_expand_omp (gcc::context *ctxt) /* Routines to lower OpenMP directives into OMP-GIMPLE. */ +/* Lower the OpenACC parallel directive in the current statement + in GSI_P. CTX holds context information for the directive. */ + +static void +lower_oacc_parallel (gimple_stmt_iterator *gsi_p, omp_context *ctx) +{ + tree clauses; + tree child_fn, t, c; + gimple stmt = gsi_stmt (*gsi_p); + gimple par_bind, bind; + gimple_seq par_body, olist, ilist, new_body; + struct gimplify_ctx gctx; + location_t loc = gimple_location (stmt); + unsigned int map_cnt = 0; + + clauses = gimple_oacc_parallel_clauses (stmt); + par_bind = gimple_seq_first_stmt (gimple_omp_body (stmt)); + par_body = gimple_bind_body (par_bind); + child_fn = ctx->cb.dst_fn; + + push_gimplify_context (&gctx); + + for (c = clauses; c ; c = OMP_CLAUSE_CHAIN (c)) + switch (OMP_CLAUSE_CODE (c)) + { + tree var, x; + + default: + break; + case OMP_CLAUSE_MAP: + case OMP_CLAUSE_TO: + case OMP_CLAUSE_FROM: + var = OMP_CLAUSE_DECL (c); + if (!DECL_P (var)) + { + if (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_MAP + || !OMP_CLAUSE_MAP_ZERO_BIAS_ARRAY_SECTION (c)) + map_cnt++; + continue; + } + + if (DECL_SIZE (var) + && TREE_CODE (DECL_SIZE (var)) != INTEGER_CST) + { + tree var2 = DECL_VALUE_EXPR (var); + gcc_assert (TREE_CODE (var2) == INDIRECT_REF); + var2 = TREE_OPERAND (var2, 0); + gcc_assert (DECL_P (var2)); + var = var2; + } + + if (!maybe_lookup_field (var, ctx)) + continue; + + /* Preserve indentation of lower_omp_target. */ + if (1) + { + x = build_receiver_ref (var, true, ctx); + tree new_var = lookup_decl (var, ctx); + if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP + && OMP_CLAUSE_MAP_KIND (c) == OMP_CLAUSE_MAP_POINTER + && !OMP_CLAUSE_MAP_ZERO_BIAS_ARRAY_SECTION (c) + && TREE_CODE (TREE_TYPE (var)) == ARRAY_TYPE) + x = build_simple_mem_ref (x); + SET_DECL_VALUE_EXPR (new_var, x); + DECL_HAS_VALUE_EXPR_P (new_var) = 1; + } + map_cnt++; + } + + target_nesting_level++; + lower_omp (&par_body, ctx); + target_nesting_level--; + + /* Declare all the variables created by mapping and the variables + declared in the scope of the body. */ + record_vars_into (ctx->block_vars, child_fn); + record_vars_into (gimple_bind_vars (par_bind), child_fn); + + olist = NULL; + ilist = NULL; + if (ctx->record_type) + { + ctx->sender_decl + = create_tmp_var (ctx->record_type, ".omp_data_arr"); + DECL_NAMELESS (ctx->sender_decl) = 1; + TREE_ADDRESSABLE (ctx->sender_decl) = 1; + t = make_tree_vec (3); + TREE_VEC_ELT (t, 0) = ctx->sender_decl; + TREE_VEC_ELT (t, 1) + = create_tmp_var (build_array_type_nelts (size_type_node, map_cnt), + ".omp_data_sizes"); + DECL_NAMELESS (TREE_VEC_ELT (t, 1)) = 1; + TREE_ADDRESSABLE (TREE_VEC_ELT (t, 1)) = 1; + TREE_STATIC (TREE_VEC_ELT (t, 1)) = 1; + TREE_VEC_ELT (t, 2) + = create_tmp_var (build_array_type_nelts (unsigned_char_type_node, + map_cnt), + ".omp_data_kinds"); + 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_oacc_parallel_set_data_arg (stmt, t); + + vec *vsize; + vec *vkind; + vec_alloc (vsize, map_cnt); + vec_alloc (vkind, map_cnt); + unsigned int map_idx = 0; + + for (c = clauses; c ; c = OMP_CLAUSE_CHAIN (c)) + switch (OMP_CLAUSE_CODE (c)) + { + tree ovar, nc; + + default: + break; + case OMP_CLAUSE_MAP: + case OMP_CLAUSE_TO: + case OMP_CLAUSE_FROM: + nc = c; + ovar = OMP_CLAUSE_DECL (c); + if (!DECL_P (ovar)) + { + if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP + && OMP_CLAUSE_MAP_ZERO_BIAS_ARRAY_SECTION (c)) + { + gcc_checking_assert (OMP_CLAUSE_DECL (OMP_CLAUSE_CHAIN (c)) + == get_base_address (ovar)); + nc = OMP_CLAUSE_CHAIN (c); + ovar = OMP_CLAUSE_DECL (nc); + } + else + { + tree x = build_sender_ref (ovar, ctx); + tree v + = build_fold_addr_expr_with_type (ovar, ptr_type_node); + gimplify_assign (x, v, &ilist); + nc = NULL_TREE; + } + } + else + { + if (DECL_SIZE (ovar) + && TREE_CODE (DECL_SIZE (ovar)) != INTEGER_CST) + { + tree ovar2 = DECL_VALUE_EXPR (ovar); + gcc_assert (TREE_CODE (ovar2) == INDIRECT_REF); + ovar2 = TREE_OPERAND (ovar2, 0); + gcc_assert (DECL_P (ovar2)); + ovar = ovar2; + } + if (!maybe_lookup_field (ovar, ctx)) + continue; + } + + if (nc) + { + tree var = lookup_decl_in_outer_ctx (ovar, ctx); + tree x = build_sender_ref (ovar, ctx); + if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP + && OMP_CLAUSE_MAP_KIND (c) == OMP_CLAUSE_MAP_POINTER + && !OMP_CLAUSE_MAP_ZERO_BIAS_ARRAY_SECTION (c) + && TREE_CODE (TREE_TYPE (ovar)) == ARRAY_TYPE) + { + tree avar + = create_tmp_var (TREE_TYPE (TREE_TYPE (x)), NULL); + mark_addressable (avar); + gimplify_assign (avar, build_fold_addr_expr (var), &ilist); + avar = build_fold_addr_expr (avar); + gimplify_assign (x, avar, &ilist); + } + else if (is_gimple_reg (var)) + { + tree avar = create_tmp_var (TREE_TYPE (var), NULL); + mark_addressable (avar); + if (OMP_CLAUSE_MAP_KIND (c) != OMP_CLAUSE_MAP_ALLOC + && OMP_CLAUSE_MAP_KIND (c) != OMP_CLAUSE_MAP_FROM) + gimplify_assign (avar, var, &ilist); + avar = build_fold_addr_expr (avar); + gimplify_assign (x, avar, &ilist); + if ((OMP_CLAUSE_MAP_KIND (c) == OMP_CLAUSE_MAP_FROM + || OMP_CLAUSE_MAP_KIND (c) == OMP_CLAUSE_MAP_TOFROM) + && !TYPE_READONLY (TREE_TYPE (var))) + { + x = build_sender_ref (ovar, ctx); + x = build_simple_mem_ref (x); + gimplify_assign (var, x, &olist); + } + } + else + { + var = build_fold_addr_expr (var); + gimplify_assign (x, var, &ilist); + } + } + tree s = OMP_CLAUSE_SIZE (c); + if (s == NULL_TREE) + s = TYPE_SIZE_UNIT (TREE_TYPE (ovar)); + s = fold_convert (size_type_node, s); + tree purpose = size_int (map_idx++); + CONSTRUCTOR_APPEND_ELT (vsize, purpose, s); + if (TREE_CODE (s) != INTEGER_CST) + TREE_STATIC (TREE_VEC_ELT (t, 1)) = 0; + + unsigned char tkind = 0; + switch (OMP_CLAUSE_CODE (c)) + { + case OMP_CLAUSE_MAP: + tkind = OMP_CLAUSE_MAP_KIND (c); + break; + case OMP_CLAUSE_TO: + tkind = OMP_CLAUSE_MAP_TO; + break; + case OMP_CLAUSE_FROM: + tkind = OMP_CLAUSE_MAP_FROM; + break; + default: + gcc_unreachable (); + } + unsigned int talign = TYPE_ALIGN_UNIT (TREE_TYPE (ovar)); + if (DECL_P (ovar) && DECL_ALIGN_UNIT (ovar) > talign) + talign = DECL_ALIGN_UNIT (ovar); + talign = ceil_log2 (talign); + tkind |= talign << 3; + CONSTRUCTOR_APPEND_ELT (vkind, purpose, + build_int_cst (unsigned_char_type_node, + tkind)); + if (nc && nc != c) + c = nc; + } + + gcc_assert (map_idx == map_cnt); + + DECL_INITIAL (TREE_VEC_ELT (t, 1)) + = build_constructor (TREE_TYPE (TREE_VEC_ELT (t, 1)), vsize); + DECL_INITIAL (TREE_VEC_ELT (t, 2)) + = build_constructor (TREE_TYPE (TREE_VEC_ELT (t, 2)), vkind); + if (!TREE_STATIC (TREE_VEC_ELT (t, 1))) + { + gimple_seq initlist = NULL; + force_gimple_operand (build1 (DECL_EXPR, void_type_node, + TREE_VEC_ELT (t, 1)), + &initlist, true, NULL_TREE); + gimple_seq_add_seq (&ilist, initlist); + } + + tree clobber = build_constructor (ctx->record_type, NULL); + TREE_THIS_VOLATILE (clobber) = 1; + gimple_seq_add_stmt (&olist, gimple_build_assign (ctx->sender_decl, + clobber)); + } + + /* Once all the expansions are done, sequence all the different + fragments inside gimple_omp_body. */ + + new_body = NULL; + + if (ctx->record_type) + { + t = build_fold_addr_expr_loc (loc, ctx->sender_decl); + /* fixup_child_record_type might have changed receiver_decl's type. */ + t = fold_convert_loc (loc, TREE_TYPE (ctx->receiver_decl), t); + gimple_seq_add_stmt (&new_body, + gimple_build_assign (ctx->receiver_decl, t)); + } + + gimple_seq_add_seq (&new_body, par_body); + gcc_assert (!ctx->cancellable); + new_body = maybe_catch_exception (new_body); + gimple_seq_add_stmt (&new_body, gimple_build_omp_return (false)); + gimple_omp_set_body (stmt, new_body); + + bind = gimple_build_bind (NULL, NULL, gimple_bind_block (par_bind)); + gsi_replace (gsi_p, bind, true); + gimple_bind_add_seq (bind, ilist); + gimple_bind_add_stmt (bind, stmt); + gimple_bind_add_seq (bind, olist); + + pop_gimplify_context (NULL); +} + /* If ctx is a worksharing context inside of a cancellable parallel region and it isn't nowait, add lhs to its GIMPLE_OMP_RETURN and conditional branch to parallel's cancel_label to handle @@ -8286,6 +8930,8 @@ make_pass_expand_omp (gcc::context *ctxt) static void maybe_add_implicit_barrier_cancel (omp_context *ctx, gimple_seq *body) { + gcc_assert (gimple_code (ctx->stmt) != GIMPLE_OACC_PARALLEL); + gimple omp_return = gimple_seq_last_stmt (*body); gcc_assert (gimple_code (omp_return) == GIMPLE_OMP_RETURN); if (gimple_omp_return_nowait_p (omp_return)) @@ -9051,6 +9697,8 @@ task_copyfn_remap_type (struct omp_taskcopy_context *tcctx, tree orig_type) static void create_task_copyfn (gimple task_stmt, omp_context *ctx) { + gcc_assert (gimple_code (ctx->stmt) != GIMPLE_OACC_PARALLEL); + struct function *child_cfun; tree child_fn, t, c, src, dst, f, sf, arg, sarg, decl; tree record_type, srecord_type, bind, list; @@ -9909,6 +10557,12 @@ lower_omp_1 (gimple_stmt_iterator *gsi_p, omp_context *ctx) case GIMPLE_BIND: lower_omp (gimple_bind_body_ptr (stmt), ctx); break; + case GIMPLE_OACC_PARALLEL: + ctx = maybe_lookup_ctx (stmt); + gcc_assert (ctx); + gcc_assert (!ctx->cancellable); + lower_oacc_parallel (gsi_p, ctx); + break; case GIMPLE_OMP_PARALLEL: case GIMPLE_OMP_TASK: ctx = maybe_lookup_ctx (stmt); @@ -10357,6 +11011,7 @@ make_gimple_omp_edges (basic_block bb, struct omp_region **region) switch (code) { + case GIMPLE_OACC_PARALLEL: case GIMPLE_OMP_PARALLEL: case GIMPLE_OMP_TASK: case GIMPLE_OMP_FOR: diff --git gcc/testsuite/c-c++-common/goacc-gomp/nesting-fail-1.c gcc/testsuite/c-c++-common/goacc-gomp/nesting-fail-1.c new file mode 100644 index 0000000..875ec66 --- /dev/null +++ gcc/testsuite/c-c++-common/goacc-gomp/nesting-fail-1.c @@ -0,0 +1,121 @@ +/* TODO: Some of these should either be allowed or fail with a more sensible + error message. */ +void +f1 (void) +{ + int i; + +#pragma omp parallel + { +#pragma acc parallel /* { dg-error "may not be nested" } */ + ; + } + +#pragma omp for + for (i = 0; i < 3; i++) + { +#pragma acc parallel /* { dg-error "may not be nested" } */ + ; + } + +#pragma omp sections + { +#pragma acc parallel /* { dg-error "may not be nested" } */ + ; + } + +#pragma omp single + { +#pragma acc parallel /* { dg-error "may not be nested" } */ + ; + } + +#pragma omp task + { +#pragma acc parallel /* { dg-error "may not be nested" } */ + ; + } + +#pragma omp master + { +#pragma acc parallel /* { dg-error "may not be nested" } */ + ; + } + +#pragma omp critical + { +#pragma acc parallel /* { dg-error "may not be nested" } */ + ; + } + +#pragma omp ordered + { +#pragma acc parallel /* { dg-error "may not be nested" } */ + ; + } +} + +/* TODO: Some of these should either be allowed or fail with a more sensible + error message. */ +void +f2 (void) +{ +#pragma acc parallel + { +#pragma omp parallel /* { dg-error "may not be nested" } */ + ; + } + +#pragma acc parallel + { + int i; +#pragma omp for /* { dg-error "may not be nested" } */ + for (i = 0; i < 3; i++) + ; + } + +#pragma acc parallel + { +#pragma omp sections /* { dg-error "may not be nested" } */ + { + ; + } + } + +#pragma acc parallel + { +#pragma omp single /* { dg-error "may not be nested" } */ + ; + } + +#pragma acc parallel + { +#pragma omp task /* { dg-error "may not be nested" } */ + ; + } + +#pragma acc parallel + { +#pragma omp master /* { dg-error "may not be nested" } */ + ; + } + +#pragma acc parallel + { +#pragma omp critical /* { dg-error "may not be nested" } */ + ; + } + +#pragma acc parallel + { + int i; +#pragma omp atomic write + i = 0; /* { dg-error "may not be nested" } */ + } + +#pragma acc parallel + { +#pragma omp ordered /* { dg-error "may not be nested" } */ + ; + } +} diff --git gcc/testsuite/c-c++-common/goacc/nesting-fail-1.c gcc/testsuite/c-c++-common/goacc/nesting-fail-1.c new file mode 100644 index 0000000..6501397 --- /dev/null +++ gcc/testsuite/c-c++-common/goacc/nesting-fail-1.c @@ -0,0 +1,11 @@ +/* TODO: While the OpenACC specification does allow for certain kinds of + nesting, we don't support that yet. */ +void +f1 (void) +{ +#pragma acc parallel + { +#pragma acc parallel /* { dg-error "may not be nested" } */ + ; + } +} diff --git gcc/testsuite/c-c++-common/goacc/parallel-1.c gcc/testsuite/c-c++-common/goacc/parallel-1.c new file mode 100644 index 0000000..cd19527 --- /dev/null +++ gcc/testsuite/c-c++-common/goacc/parallel-1.c @@ -0,0 +1,6 @@ +void +foo (void) +{ +#pragma acc parallel + foo (); +} diff --git gcc/testsuite/c-c++-common/goacc/parallel-fail-1.c gcc/testsuite/c-c++-common/goacc/parallel-fail-1.c new file mode 100644 index 0000000..efc6f14 --- /dev/null +++ gcc/testsuite/c-c++-common/goacc/parallel-fail-1.c @@ -0,0 +1,6 @@ +void +foo (void) +{ +#pragma acc parallel foo /* { dg-error "expected clause before 'foo'" } */ + foo (); +} diff --git gcc/tree-inline.c gcc/tree-inline.c index 74f333b..eeb4992 100644 --- gcc/tree-inline.c +++ gcc/tree-inline.c @@ -1299,6 +1299,9 @@ remap_gimple_stmt (gimple stmt, copy_body_data *id) copy = gimple_build_wce (s1); break; + case GIMPLE_OACC_PARALLEL: + abort (); + case GIMPLE_OMP_PARALLEL: s1 = remap_gimple_seq (gimple_omp_body (stmt), id); copy = gimple_build_omp_parallel @@ -3849,6 +3852,7 @@ 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_PARALLEL: case GIMPLE_OMP_PARALLEL: case GIMPLE_OMP_TASK: case GIMPLE_OMP_CRITICAL: diff --git gcc/tree-nested.c gcc/tree-nested.c index dc63ef6..8aba4f4 100644 --- gcc/tree-nested.c +++ gcc/tree-nested.c @@ -1238,6 +1238,9 @@ convert_nonlocal_reference_stmt (gimple_stmt_iterator *gsi, bool *handled_ops_p, } break; + case GIMPLE_OACC_PARALLEL: + abort (); + case GIMPLE_OMP_PARALLEL: case GIMPLE_OMP_TASK: save_suppress = info->suppress_expansion; @@ -1679,6 +1682,9 @@ convert_local_reference_stmt (gimple_stmt_iterator *gsi, bool *handled_ops_p, switch (gimple_code (stmt)) { + case GIMPLE_OACC_PARALLEL: + abort (); + case GIMPLE_OMP_PARALLEL: case GIMPLE_OMP_TASK: save_suppress = info->suppress_expansion; @@ -2008,6 +2014,9 @@ convert_tramp_reference_stmt (gimple_stmt_iterator *gsi, bool *handled_ops_p, break; } + case GIMPLE_OACC_PARALLEL: + abort (); + case GIMPLE_OMP_PARALLEL: case GIMPLE_OMP_TASK: { @@ -2068,6 +2077,9 @@ convert_gimple_call (gimple_stmt_iterator *gsi, bool *handled_ops_p, } break; + case GIMPLE_OACC_PARALLEL: + abort (); + case GIMPLE_OMP_PARALLEL: case GIMPLE_OMP_TASK: save_static_chain_added = info->static_chain_added; diff --git gcc/tree-pretty-print.c gcc/tree-pretty-print.c index fe75633..153d01f 100644 --- gcc/tree-pretty-print.c +++ gcc/tree-pretty-print.c @@ -2346,6 +2346,11 @@ dump_generic_node (pretty_printer *buffer, tree node, int spc, int flags, pp_string (buffer, " > "); break; + case OACC_PARALLEL: + pp_string (buffer, "#pragma acc parallel"); + dump_omp_clauses (buffer, OACC_PARALLEL_CLAUSES (node), spc, flags); + goto dump_omp_body; + case OMP_PARALLEL: pp_string (buffer, "#pragma omp parallel"); dump_omp_clauses (buffer, OMP_PARALLEL_CLAUSES (node), spc, flags); diff --git gcc/tree.def gcc/tree.def index 399b5af..87fec57 100644 --- gcc/tree.def +++ gcc/tree.def @@ -1000,8 +1000,15 @@ DEFTREECODE (TARGET_MEM_REF, "target_mem_ref", tcc_reference, 5) chain of component references offsetting p by c. */ DEFTREECODE (MEM_REF, "mem_ref", tcc_reference, 2) -/* The ordering of the codes between OMP_PARALLEL and OMP_CRITICAL is - exposed to TREE_RANGE_CHECK. */ +/* OpenACC and OpenMP. As it is exposed in TREE_RANGE_CHECK invocations, do + not change the ordering of these codes. */ + +/* OpenACC - #pragma acc parallel [clause1 ... clauseN] + Operand 0: OACC_PARALLEL_BODY: Code to be executed in parallel. + Operand 1: OACC_PARALLEL_CLAUSES: List of clauses. */ + +DEFTREECODE (OACC_PARALLEL, "oacc_parallel", tcc_statement, 2) + /* OpenMP - #pragma omp parallel [clause1 ... clauseN] Operand 0: OMP_PARALLEL_BODY: Code to be executed by all threads. Operand 1: OMP_PARALLEL_CLAUSES: List of clauses. */ diff --git gcc/tree.h gcc/tree.h index 22a576f..06d94cf 100644 --- gcc/tree.h +++ gcc/tree.h @@ -1171,9 +1171,14 @@ extern void protected_set_expr_location (tree, location_t); /* OpenMP directive and clause accessors. */ #define OMP_BODY(NODE) \ - TREE_OPERAND (TREE_RANGE_CHECK (NODE, OMP_PARALLEL, OMP_CRITICAL), 0) + TREE_OPERAND (TREE_RANGE_CHECK (NODE, OACC_PARALLEL, OMP_CRITICAL), 0) #define OMP_CLAUSES(NODE) \ - TREE_OPERAND (TREE_RANGE_CHECK (NODE, OMP_PARALLEL, OMP_SINGLE), 1) + TREE_OPERAND (TREE_RANGE_CHECK (NODE, OACC_PARALLEL, OMP_SINGLE), 1) + +#define OACC_PARALLEL_BODY(NODE) \ + TREE_OPERAND (OACC_PARALLEL_CHECK (NODE), 0) +#define OACC_PARALLEL_CLAUSES(NODE) \ + TREE_OPERAND (OACC_PARALLEL_CHECK (NODE), 1) #define OMP_PARALLEL_BODY(NODE) TREE_OPERAND (OMP_PARALLEL_CHECK (NODE), 0) #define OMP_PARALLEL_CLAUSES(NODE) TREE_OPERAND (OMP_PARALLEL_CHECK (NODE), 1) diff --git libgomp/Makefile.am libgomp/Makefile.am index 0b5c097..37b36bd 100644 --- libgomp/Makefile.am +++ libgomp/Makefile.am @@ -60,7 +60,7 @@ libgomp_la_LINK = $(LINK) $(libgomp_la_LDFLAGS) libgomp_la_SOURCES = alloc.c barrier.c critical.c env.c error.c iter.c \ iter_ull.c loop.c loop_ull.c ordered.c parallel.c sections.c single.c \ task.c team.c work.c lock.c mutex.c proc.c sem.c bar.c ptrlock.c \ - time.c fortran.c affinity.c target.c + time.c fortran.c affinity.c target.c oacc-parallel.c nodist_noinst_HEADERS = libgomp_f.h nodist_libsubinclude_HEADERS = omp.h openacc.h diff --git libgomp/Makefile.in libgomp/Makefile.in index 9ee1bec..bc60253d 100644 --- libgomp/Makefile.in +++ libgomp/Makefile.in @@ -96,7 +96,7 @@ am_libgomp_la_OBJECTS = alloc.lo barrier.lo critical.lo env.lo \ error.lo iter.lo iter_ull.lo loop.lo loop_ull.lo ordered.lo \ parallel.lo sections.lo single.lo task.lo team.lo work.lo \ lock.lo mutex.lo proc.lo sem.lo bar.lo ptrlock.lo time.lo \ - fortran.lo affinity.lo target.lo + fortran.lo affinity.lo target.lo oacc-parallel.lo libgomp_la_OBJECTS = $(am_libgomp_la_OBJECTS) DEFAULT_INCLUDES = -I.@am__isrc@ depcomp = $(SHELL) $(top_srcdir)/../depcomp @@ -317,7 +317,7 @@ libgomp_la_LINK = $(LINK) $(libgomp_la_LDFLAGS) libgomp_la_SOURCES = alloc.c barrier.c critical.c env.c error.c iter.c \ iter_ull.c loop.c loop_ull.c ordered.c parallel.c sections.c single.c \ task.c team.c work.c lock.c mutex.c proc.c sem.c bar.c ptrlock.c \ - time.c fortran.c affinity.c target.c + time.c fortran.c affinity.c target.c oacc-parallel.c nodist_noinst_HEADERS = libgomp_f.h nodist_libsubinclude_HEADERS = omp.h openacc.h @@ -469,6 +469,7 @@ distclean-compile: @AMDEP_TRUE@@am__include@ @am__quote@./$(DEPDIR)/loop.Plo@am__quote@ @AMDEP_TRUE@@am__include@ @am__quote@./$(DEPDIR)/loop_ull.Plo@am__quote@ @AMDEP_TRUE@@am__include@ @am__quote@./$(DEPDIR)/mutex.Plo@am__quote@ +@AMDEP_TRUE@@am__include@ @am__quote@./$(DEPDIR)/oacc-parallel.Plo@am__quote@ @AMDEP_TRUE@@am__include@ @am__quote@./$(DEPDIR)/ordered.Plo@am__quote@ @AMDEP_TRUE@@am__include@ @am__quote@./$(DEPDIR)/parallel.Plo@am__quote@ @AMDEP_TRUE@@am__include@ @am__quote@./$(DEPDIR)/proc.Plo@am__quote@ diff --git libgomp/libgomp.map libgomp/libgomp.map index f094ed2..2b64d05 100644 --- libgomp/libgomp.map +++ libgomp/libgomp.map @@ -232,4 +232,6 @@ OACC_2.0 { }; GOACC_2.0 { + global: + GOACC_parallel; }; diff --git libgomp/libgomp_g.h libgomp/libgomp_g.h index 577956a..394f3a8 100644 --- libgomp/libgomp_g.h +++ libgomp/libgomp_g.h @@ -214,4 +214,9 @@ extern void GOMP_target_update (int, const void *, size_t, void **, size_t *, unsigned char *); extern void GOMP_teams (unsigned int, unsigned int); +/* oacc-parallel.c */ + +extern void GOACC_parallel (int, void (*) (void *), const void *, + size_t, void **, size_t *, unsigned char *); + #endif /* LIBGOMP_G_H */ diff --git libgomp/oacc-parallel.c libgomp/oacc-parallel.c new file mode 100644 index 0000000..730b83b --- /dev/null +++ libgomp/oacc-parallel.c @@ -0,0 +1,36 @@ +/* Copyright (C) 2013 Free Software Foundation, Inc. + + Contributed by Thomas Schwinge . + + This file is part of the GNU OpenMP Library (libgomp). + + Libgomp is free software; you can redistribute it and/or modify it + under the terms of the GNU General Public License as published by + the Free Software Foundation; either version 3, or (at your option) + any later version. + + Libgomp is distributed in the hope that it will be useful, but WITHOUT ANY + WARRANTY; without even the implied warranty of MERCHANTABILITY or FITNESS + FOR A PARTICULAR PURPOSE. See the GNU General Public License for + more details. + + Under Section 7 of GPL version 3, you are granted additional + permissions described in the GCC Runtime Library Exception, version + 3.1, as published by the Free Software Foundation. + + You should have received a copy of the GNU General Public License and + a copy of the GCC Runtime Library Exception along with this program; + see the files COPYING3 and COPYING.RUNTIME respectively. If not, see + . */ + +/* This file handles the OpenACC parallel construct. */ + +#include "libgomp_g.h" + +void +GOACC_parallel (int device, void (*fn) (void *), const void *openmp_target, + size_t mapnum, void **hostaddrs, size_t *sizes, + unsigned char *kinds) +{ + GOMP_target (device, fn, openmp_target, mapnum, hostaddrs, sizes, kinds); +} diff --git libgomp/testsuite/libgomp.oacc-c/goacc_parallel.c libgomp/testsuite/libgomp.oacc-c/goacc_parallel.c new file mode 100644 index 0000000..b9bdffa --- /dev/null +++ libgomp/testsuite/libgomp.oacc-c/goacc_parallel.c @@ -0,0 +1,25 @@ +/* { dg-do run } */ + +#include "libgomp_g.h" + +extern void abort (); + +volatile int i; + +void +f (void *data) +{ + if (i != -1) + abort (); + i = 42; +} + +int main(void) +{ + i = -1; + GOACC_parallel (0, f, (const void *) 0, 0, (void *) 0, (void *) 0, (void *) 0); + if (i != 42) + abort (); + + return 0; +} diff --git libgomp/testsuite/libgomp.oacc-c/parallel-1.c libgomp/testsuite/libgomp.oacc-c/parallel-1.c new file mode 100644 index 0000000..b40545d --- /dev/null +++ libgomp/testsuite/libgomp.oacc-c/parallel-1.c @@ -0,0 +1,26 @@ +/* { dg-do run } */ + +extern void abort (); + +volatile int i; + +int main(void) +{ + volatile int j; + + i = -0x42; + j = -42; +#pragma acc parallel + { + if (i != -0x42 || j != -42) + abort (); + i = 42; + j = 0x42; + if (i != 42 || j != 0x42) + abort (); + } + if (i != 42 || j != 0x42) + abort (); + + return 0; +}