From mboxrd@z Thu Jan 1 00:00:00 1970 Return-Path: Received: (qmail 4383 invoked by alias); 6 Nov 2013 19:44:58 -0000 Mailing-List: contact gcc-patches-help@gcc.gnu.org; run by ezmlm Precedence: bulk List-Id: List-Archive: List-Post: List-Help: Sender: gcc-patches-owner@gcc.gnu.org Received: (qmail 4358 invoked by uid 89); 6 Nov 2013 19:44:58 -0000 Authentication-Results: sourceware.org; auth=none X-Virus-Found: No X-Spam-SWARE-Status: No, score=1.5 required=5.0 tests=AWL,BAYES_99,RDNS_NONE,URIBL_BLOCKED autolearn=no version=3.3.2 X-HELO: eggs.gnu.org Received: from Unknown (HELO eggs.gnu.org) (208.118.235.92) by sourceware.org (qpsmtpd/0.93/v0.84-503-g423c35a) with (AES256-SHA encrypted) ESMTPS; Wed, 06 Nov 2013 19:42:58 +0000 Received: from Debian-exim by eggs.gnu.org with spam-scanned (Exim 4.71) (envelope-from ) id 1Ve8zz-0005pM-ME for gcc-patches@gcc.gnu.org; Wed, 06 Nov 2013 14:42:50 -0500 Received: from relay1.mentorg.com ([192.94.38.131]:60290) by eggs.gnu.org with esmtp (Exim 4.71) (envelope-from ) id 1Ve8zz-0005p4-3i for gcc-patches@gcc.gnu.org; Wed, 06 Nov 2013 14:42:43 -0500 Received: from svr-orw-fem-01.mgc.mentorg.com ([147.34.98.93]) by relay1.mentorg.com with esmtp id 1Ve8zy-0000UJ-Ex from Thomas_Schwinge@mentor.com for gcc-patches@gcc.gnu.org; Wed, 06 Nov 2013 11:42:42 -0800 Received: from SVR-ORW-FEM-03.mgc.mentorg.com ([147.34.97.39]) by svr-orw-fem-01.mgc.mentorg.com over TLS secured channel with Microsoft SMTPSVC(6.0.3790.4675); Wed, 6 Nov 2013 11:42:42 -0800 Received: from build5-lucid-cs (147.34.91.1) by svr-orw-fem-03.mgc.mentorg.com (147.34.97.39) with Microsoft SMTP Server id 14.2.247.3; Wed, 6 Nov 2013 11:42:41 -0800 Received: by build5-lucid-cs (Postfix, from userid 49978) id A33FD3268AB; Wed, 6 Nov 2013 11:42:37 -0800 (PST) From: To: CC: Thomas Schwinge Subject: [gomp4 9/9] OpenACC: Basic support for #pragma acc parallel. Date: Wed, 06 Nov 2013 19:53:00 -0000 Message-ID: <1383766943-8863-9-git-send-email-thomas@codesourcery.com> In-Reply-To: <1383766943-8863-8-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> MIME-Version: 1.0 Content-Type: text/plain X-detected-operating-system: by eggs.gnu.org: Windows NT kernel [generic] [fuzzy] X-Received-From: 192.94.38.131 X-SW-Source: 2013-11/txt/msg00631.txt.bz2 From: Thomas Schwinge 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 | 1047 ++++++++++++++++---- .../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, 1450 insertions(+), 208 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 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,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; @@ -4553,44 +4683,20 @@ expand_omp_taskreg (struct omp_region *region) gimple_stmt_iterator gsi; gimple entry_stmt, stmt; edge e; - vec *ws_args; entry_stmt = last_stmt (region->entry); - child_fn = gimple_omp_taskreg_child_fn (entry_stmt); + 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; - if (is_combined_parallel (region)) - ws_args = region->ws_args; - else - ws_args = NULL; - - if (child_cfun->cfg) - { - /* Due to inlining, it may happen that we have already outlined - the region, in which case all we need to do is make the - sub-graph unreachable and emit the parallel call. */ - edge entry_succ_e, exit_succ_e; - gimple_stmt_iterator gsi; - - entry_succ_e = single_succ_edge (entry_bb); - - gsi = gsi_last_bb (entry_bb); - gcc_assert (gimple_code (gsi_stmt (gsi)) == GIMPLE_OMP_PARALLEL - || gimple_code (gsi_stmt (gsi)) == GIMPLE_OMP_TASK); - gsi_remove (&gsi, true); - - new_bb = entry_bb; - if (exit_bb) - { - exit_succ_e = single_succ_edge (exit_bb); - make_edge (new_bb, exit_succ_e->dest, EDGE_FALLTHRU); - } - remove_edge_and_dominated_blocks (entry_succ_e); - } - else + /* Preserve indentation of expand_omp_target and expand_omp_taskreg. */ + if (1) { unsigned srcidx, dstidx, num; @@ -4607,17 +4713,17 @@ expand_omp_taskreg (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_taskreg_data_arg (entry_stmt)) + if (gimple_oacc_parallel_data_arg (entry_stmt)) { basic_block entry_succ_bb = single_succ (entry_bb); gimple_stmt_iterator gsi; - tree arg, narg; + 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)) { - gimple stmt; - gcc_assert (!gsi_end_p (gsi)); stmt = gsi_stmt (gsi); if (gimple_code (stmt) != GIMPLE_ASSIGN) @@ -4631,8 +4737,7 @@ expand_omp_taskreg (struct omp_region *region) effectively doing a STRIP_NOPS. */ if (TREE_CODE (arg) == ADDR_EXPR - && TREE_OPERAND (arg, 0) - == gimple_omp_taskreg_data_arg (entry_stmt)) + && TREE_OPERAND (arg, 0) == sender) { parcopy_stmt = stmt; break; @@ -4643,36 +4748,14 @@ expand_omp_taskreg (struct omp_region *region) gcc_assert (parcopy_stmt != NULL); arg = DECL_ARGUMENTS (child_fn); - if (!gimple_in_ssa_p (cfun)) - { - if (gimple_assign_lhs (parcopy_stmt) == arg) - gsi_remove (&gsi, true); - else - { - /* ?? Is setting the subcode really necessary ?? */ - gimple_omp_set_subcode (parcopy_stmt, TREE_CODE (arg)); - gimple_assign_set_rhs1 (parcopy_stmt, arg); - } - } - else - { - /* If we are in ssa form, we must load the value from the default - definition of the argument. That should not be defined now, - since the argument is not used uninitialized. */ - gcc_assert (ssa_default_def (cfun, arg) == NULL); - narg = make_ssa_name (arg, gimple_build_nop ()); - set_ssa_default_def (cfun, arg, narg); - /* ?? Is setting the subcode really necessary ?? */ - gimple_omp_set_subcode (parcopy_stmt, TREE_CODE (narg)); - gimple_assign_set_rhs1 (parcopy_stmt, narg); - update_stmt (parcopy_stmt); - } + 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 parallel/task block + /* 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)) @@ -4689,12 +4772,11 @@ expand_omp_taskreg (struct omp_region *region) for (t = DECL_ARGUMENTS (child_fn); t; t = DECL_CHAIN (t)) DECL_CONTEXT (t) = child_fn; - /* Split ENTRY_BB at GIMPLE_OMP_PARALLEL or GIMPLE_OMP_TASK, + /* 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_OMP_PARALLEL - || gimple_code (stmt) == GIMPLE_OMP_TASK)); + gcc_assert (stmt && (gimple_code (stmt) == GIMPLE_OACC_PARALLEL)); gsi_remove (&gsi, true); e = split_block (entry_bb, stmt); entry_bb = e->dest; @@ -4711,22 +4793,14 @@ expand_omp_taskreg (struct omp_region *region) gsi_remove (&gsi, true); } - /* Move the parallel region into CHILD_CFUN. */ + /* Move the region into CHILD_CFUN. */ - if (gimple_in_ssa_p (cfun)) - { - init_tree_ssa (child_cfun); - init_ssa_operands (child_cfun); - child_cfun->gimple_df->in_ssa_p = true; - block = NULL_TREE; - } - else - block = gimple_block (entry_stmt); + 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 OMP expansion process cannot guarantee an up-to-date + /* 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; @@ -4752,8 +4826,6 @@ expand_omp_taskreg (struct omp_region *region) /* Fix the callgraph edges for child_cfun. Those for cfun will be fixed in a following pass. */ push_cfun (child_cfun); - if (optimize) - optimize_omp_library_calls (entry_stmt); rebuild_cgraph_edges (); /* Some EH regions might become dead, see PR34608. If @@ -4770,73 +4842,359 @@ expand_omp_taskreg (struct omp_region *region) if (changed) cleanup_tree_cfg (); } - if (gimple_in_ssa_p (cfun)) - update_ssa (TODO_update_ssa); pop_cfun (); } - /* Emit a library call to launch the children threads. */ - if (gimple_code (entry_stmt) == GIMPLE_OMP_PARALLEL) - expand_parallel_call (region, new_bb, entry_stmt, ws_args); - else - expand_task_call (new_bb, entry_stmt); - if (gimple_in_ssa_p (cfun)) - update_ssa (TODO_update_ssa_only_virtuals); -} + /* 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); -/* Helper function for expand_omp_{for_*,simd}. If this is the outermost - of the combined collapse > 1 loop constructs, generate code like: - if (__builtin_expect (N32 cond3 N31, 0)) goto ZERO_ITER_BB; - if (cond3 is <) - adj = STEP3 - 1; - else - adj = STEP3 + 1; - count3 = (adj + N32 - N31) / STEP3; - if (__builtin_expect (N22 cond2 N21, 0)) goto ZERO_ITER_BB; - if (cond2 is <) - adj = STEP2 - 1; - else - adj = STEP2 + 1; - count2 = (adj + N22 - N21) / STEP2; - if (__builtin_expect (N12 cond1 N11, 0)) goto ZERO_ITER_BB; - if (cond1 is <) - adj = STEP1 - 1; - else - adj = STEP1 + 1; - count1 = (adj + N12 - N11) / STEP1; - count = count1 * count2 * count3; - Furthermore, if ZERO_ITER_BB is NULL, create a BB which does: - count = 0; - and set ZERO_ITER_BB to that bb. If this isn't the outermost - of the combined loop constructs, just initialize COUNTS array - from the _looptemp_ clauses. */ + start_ix = BUILT_IN_GOACC_PARALLEL; -/* NOTE: It *could* be better to moosh all of the BBs together, - creating one larger BB with all the computation and the unexpected - jump at the end. I.e. + /* By default, the value of DEVICE is -1 (let runtime library choose). */ + device = build_int_cst (integer_type_node, -1); - bool zero3, zero2, zero1, zero; + 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); - zero3 = N32 c3 N31; - count3 = (N32 - N31) /[cl] STEP3; - zero2 = N22 c2 N21; - count2 = (N22 - N21) /[cl] STEP2; - zero1 = N12 c1 N11; - count1 = (N12 - N11) /[cl] STEP1; - zero = zero3 || zero2 || zero1; - count = count1 * count2 * count3; - if (__builtin_expect(zero, false)) goto zero_iter_bb; + /* Ensure 'device' is of the correct type. */ + device = fold_convert_loc (clause_loc, integer_type_node, device); - After all, we expect the zero=false, and thus we expect to have to - evaluate all of the comparison expressions, so short-circuiting - oughtn't be a win. Since the condition isn't protecting a - denominator, we're not concerned about divide-by-zero, so we can - fully evaluate count even if a numerator turned out to be wrong. + 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)); + } - It seems like putting this all together would create much better - scheduling opportunities, and less pressure on the chip's branch - predictor. */ + 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 +expand_omp_taskreg (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; + vec *ws_args; + + entry_stmt = last_stmt (region->entry); + child_fn = gimple_omp_taskreg_child_fn (entry_stmt); + child_cfun = DECL_STRUCT_FUNCTION (child_fn); + + entry_bb = region->entry; + exit_bb = region->exit; + + if (is_combined_parallel (region)) + ws_args = region->ws_args; + else + ws_args = NULL; + + if (child_cfun->cfg) + { + /* Due to inlining, it may happen that we have already outlined + the region, in which case all we need to do is make the + sub-graph unreachable and emit the parallel call. */ + edge entry_succ_e, exit_succ_e; + gimple_stmt_iterator gsi; + + entry_succ_e = single_succ_edge (entry_bb); + + gsi = gsi_last_bb (entry_bb); + gcc_assert (gimple_code (gsi_stmt (gsi)) == GIMPLE_OMP_PARALLEL + || gimple_code (gsi_stmt (gsi)) == GIMPLE_OMP_TASK); + gsi_remove (&gsi, true); + + new_bb = entry_bb; + if (exit_bb) + { + exit_succ_e = single_succ_edge (exit_bb); + make_edge (new_bb, exit_succ_e->dest, EDGE_FALLTHRU); + } + remove_edge_and_dominated_blocks (entry_succ_e); + } + else + { + 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_omp_taskreg_data_arg (entry_stmt)) + { + basic_block entry_succ_bb = single_succ (entry_bb); + gimple_stmt_iterator gsi; + tree arg, narg; + gimple parcopy_stmt = NULL; + + for (gsi = gsi_start_bb (entry_succ_bb); ; gsi_next (&gsi)) + { + gimple stmt; + + 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) + == gimple_omp_taskreg_data_arg (entry_stmt)) + { + parcopy_stmt = stmt; + break; + } + } + } + + gcc_assert (parcopy_stmt != NULL); + arg = DECL_ARGUMENTS (child_fn); + + if (!gimple_in_ssa_p (cfun)) + { + if (gimple_assign_lhs (parcopy_stmt) == arg) + gsi_remove (&gsi, true); + else + { + /* ?? Is setting the subcode really necessary ?? */ + gimple_omp_set_subcode (parcopy_stmt, TREE_CODE (arg)); + gimple_assign_set_rhs1 (parcopy_stmt, arg); + } + } + else + { + /* If we are in ssa form, we must load the value from the default + definition of the argument. That should not be defined now, + since the argument is not used uninitialized. */ + gcc_assert (ssa_default_def (cfun, arg) == NULL); + narg = make_ssa_name (arg, gimple_build_nop ()); + set_ssa_default_def (cfun, arg, narg); + /* ?? Is setting the subcode really necessary ?? */ + gimple_omp_set_subcode (parcopy_stmt, TREE_CODE (narg)); + gimple_assign_set_rhs1 (parcopy_stmt, narg); + update_stmt (parcopy_stmt); + } + } + + /* 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 parallel/task 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_OMP_PARALLEL or GIMPLE_OMP_TASK, + 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_OMP_PARALLEL + || gimple_code (stmt) == GIMPLE_OMP_TASK)); + 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 parallel region into CHILD_CFUN. */ + + if (gimple_in_ssa_p (cfun)) + { + init_tree_ssa (child_cfun); + init_ssa_operands (child_cfun); + child_cfun->gimple_df->in_ssa_p = true; + block = NULL_TREE; + } + else + 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 OMP 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); + if (optimize) + optimize_omp_library_calls (entry_stmt); + 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 (); + } + if (gimple_in_ssa_p (cfun)) + update_ssa (TODO_update_ssa); + pop_cfun (); + } + + /* Emit a library call to launch the children threads. */ + if (gimple_code (entry_stmt) == GIMPLE_OMP_PARALLEL) + expand_parallel_call (region, new_bb, entry_stmt, ws_args); + else + expand_task_call (new_bb, entry_stmt); + if (gimple_in_ssa_p (cfun)) + update_ssa (TODO_update_ssa_only_virtuals); +} + + +/* Helper function for expand_omp_{for_*,simd}. If this is the outermost + of the combined collapse > 1 loop constructs, generate code like: + if (__builtin_expect (N32 cond3 N31, 0)) goto ZERO_ITER_BB; + if (cond3 is <) + adj = STEP3 - 1; + else + adj = STEP3 + 1; + count3 = (adj + N32 - N31) / STEP3; + if (__builtin_expect (N22 cond2 N21, 0)) goto ZERO_ITER_BB; + if (cond2 is <) + adj = STEP2 - 1; + else + adj = STEP2 + 1; + count2 = (adj + N22 - N21) / STEP2; + if (__builtin_expect (N12 cond1 N11, 0)) goto ZERO_ITER_BB; + if (cond1 is <) + adj = STEP1 - 1; + else + adj = STEP1 + 1; + count1 = (adj + N12 - N11) / STEP1; + count = count1 * count2 * count3; + Furthermore, if ZERO_ITER_BB is NULL, create a BB which does: + count = 0; + and set ZERO_ITER_BB to that bb. If this isn't the outermost + of the combined loop constructs, just initialize COUNTS array + from the _looptemp_ clauses. */ + +/* NOTE: It *could* be better to moosh all of the BBs together, + creating one larger BB with all the computation and the unexpected + jump at the end. I.e. + + bool zero3, zero2, zero1, zero; + + zero3 = N32 c3 N31; + count3 = (N32 - N31) /[cl] STEP3; + zero2 = N22 c2 N21; + count2 = (N22 - N21) /[cl] STEP2; + zero1 = N12 c1 N11; + count1 = (N12 - N11) /[cl] STEP1; + zero = zero3 || zero2 || zero1; + count = count1 * count2 * count3; + if (__builtin_expect(zero, false)) goto zero_iter_bb; + + After all, we expect the zero=false, and thus we expect to have to + evaluate all of the comparison expressions, so short-circuiting + oughtn't be a win. Since the condition isn't protecting a + denominator, we're not concerned about divide-by-zero, so we can + fully evaluate count even if a numerator turned out to be wrong. + + It seems like putting this all together would create much better + scheduling opportunities, and less pressure on the chip's branch + predictor. */ static void expand_omp_for_init_counts (struct omp_for_data *fd, gimple_stmt_iterator *gsi, @@ -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); @@ -8203,80 +8565,362 @@ build_omp_regions (void) /* Main entry point for expanding OMP-GIMPLE into runtime calls. */ -static unsigned int -execute_expand_omp (void) -{ - build_omp_regions (); +static unsigned int +execute_expand_omp (void) +{ + build_omp_regions (); + + if (!root_omp_region) + return 0; + + if (dump_file) + { + fprintf (dump_file, "\nOMP region tree\n\n"); + dump_omp_region (dump_file, root_omp_region, 0); + fprintf (dump_file, "\n"); + } + + remove_exit_barriers (root_omp_region); + + expand_omp (root_omp_region); + + cleanup_tree_cfg (); + + free_omp_regions (); + + return 0; +} + +/* OMP expansion -- the default pass, run before creation of SSA form. */ + +static bool +gate_expand_omp (void) +{ + return ((flag_openacc || flag_openmp) + && !seen_error ()); +} + +namespace { + +const pass_data pass_data_expand_omp = +{ + GIMPLE_PASS, /* type */ + "ompexp", /* name */ + OPTGROUP_NONE, /* optinfo_flags */ + true, /* has_gate */ + true, /* has_execute */ + TV_NONE, /* tv_id */ + PROP_gimple_any, /* properties_required */ + 0, /* properties_provided */ + 0, /* properties_destroyed */ + 0, /* todo_flags_start */ + 0, /* todo_flags_finish */ +}; + +class pass_expand_omp : public gimple_opt_pass +{ +public: + pass_expand_omp (gcc::context *ctxt) + : gimple_opt_pass (pass_data_expand_omp, ctxt) + {} + + /* opt_pass methods: */ + bool gate () { return gate_expand_omp (); } + unsigned int execute () { return execute_expand_omp (); } + +}; // class pass_expand_omp + +} // anon namespace + +gimple_opt_pass * +make_pass_expand_omp (gcc::context *ctxt) +{ + return new pass_expand_omp (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--; - if (!root_omp_region) - return 0; + /* 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); - if (dump_file) + olist = NULL; + ilist = NULL; + if (ctx->record_type) { - fprintf (dump_file, "\nOMP region tree\n\n"); - dump_omp_region (dump_file, root_omp_region, 0); - fprintf (dump_file, "\n"); - } + 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); - remove_exit_barriers (root_omp_region); + vec *vsize; + vec *vkind; + vec_alloc (vsize, map_cnt); + vec_alloc (vkind, map_cnt); + unsigned int map_idx = 0; - expand_omp (root_omp_region); + for (c = clauses; c ; c = OMP_CLAUSE_CHAIN (c)) + switch (OMP_CLAUSE_CODE (c)) + { + tree ovar, nc; - cleanup_tree_cfg (); + 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; + } - free_omp_regions (); + 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; - return 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; + } -/* OMP expansion -- the default pass, run before creation of SSA form. */ + gcc_assert (map_idx == map_cnt); -static bool -gate_expand_omp (void) -{ - return ((flag_openacc || flag_openmp) - && !seen_error ()); -} + 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); + } -namespace { + 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)); + } -const pass_data pass_data_expand_omp = -{ - GIMPLE_PASS, /* type */ - "ompexp", /* name */ - OPTGROUP_NONE, /* optinfo_flags */ - true, /* has_gate */ - true, /* has_execute */ - TV_NONE, /* tv_id */ - PROP_gimple_any, /* properties_required */ - 0, /* properties_provided */ - 0, /* properties_destroyed */ - 0, /* todo_flags_start */ - 0, /* todo_flags_finish */ -}; + /* Once all the expansions are done, sequence all the different + fragments inside gimple_omp_body. */ -class pass_expand_omp : public gimple_opt_pass -{ -public: - pass_expand_omp (gcc::context *ctxt) - : gimple_opt_pass (pass_data_expand_omp, ctxt) - {} + new_body = NULL; - /* opt_pass methods: */ - bool gate () { return gate_expand_omp (); } - unsigned int execute () { return execute_expand_omp (); } + 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)); + } -}; // class pass_expand_omp + 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); -} // anon namespace + 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); -gimple_opt_pass * -make_pass_expand_omp (gcc::context *ctxt) -{ - return new pass_expand_omp (ctxt); + pop_gimplify_context (NULL); } - -/* Routines to lower OpenMP directives into OMP-GIMPLE. */ /* If ctx is a worksharing context inside of a cancellable parallel region and it isn't nowait, add lhs to its GIMPLE_OMP_RETURN @@ -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; +} -- 1.8.1.1