This is the mail archive of the
gcc-patches@gcc.gnu.org
mailing list for the GCC project.
[gomp4 9/9] OpenACC: Basic support for #pragma acc parallel.
- From: <thomas at codesourcery dot com>
- To: <gcc-patches at gcc dot gnu dot org>
- Cc: Thomas Schwinge <thomas at codesourcery dot com>
- Date: Wed, 6 Nov 2013 20:42:23 +0100
- Subject: [gomp4 9/9] OpenACC: Basic support for #pragma acc parallel.
- Authentication-results: sourceware.org; auth=none
- References: <878ux1jp2s dot fsf at schwinge dot name> <1383766943-8863-1-git-send-email-thomas at codesourcery dot com> <1383766943-8863-2-git-send-email-thomas at codesourcery dot com> <1383766943-8863-3-git-send-email-thomas at codesourcery dot com> <1383766943-8863-4-git-send-email-thomas at codesourcery dot com> <1383766943-8863-5-git-send-email-thomas at codesourcery dot com> <1383766943-8863-6-git-send-email-thomas at codesourcery dot com> <1383766943-8863-7-git-send-email-thomas at codesourcery dot com> <1383766943-8863-8-git-send-email-thomas at codesourcery dot com>
From: Thomas Schwinge <thomas@codesourcery.com>
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<pragma_ns_name> 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 <BODY, CLAUSES, CHILD_FN, DATA_ARG> 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<tree> );
+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<tree, va_gc> *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<tree, va_gc> *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<constructor_elt, va_gc> *vsize;
+ vec<constructor_elt, va_gc> *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 <thomas@codesourcery.com>.
+
+ 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
+ <http://www.gnu.org/licenses/>. */
+
+/* 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