@@ -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
@@ -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 },
@@ -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,
@@ -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;
@@ -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);
@@ -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
@@ -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)
@@ -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
@@ -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}
@@ -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:
@@ -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;
@@ -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);
@@ -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
@@ -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: \
@@ -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
@@ -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)
@@ -844,6 +844,8 @@ use_pointer_for_field (tree decl, omp_context *shared_ctx)
when we know the value is not accessible from an outer scope. */
if (shared_ctx)
{
+ gcc_assert (gimple_code (shared_ctx->stmt) != GIMPLE_OACC_PARALLEL);
+
/* ??? Trivially accessible from anywhere. But why would we even
be passing an address in this case? Should we simply assert
this to be false, or should we have a cleanup pass that removes
@@ -985,6 +987,8 @@ build_receiver_ref (tree var, bool by_ref, omp_context *ctx)
static tree
build_outer_var_ref (tree var, omp_context *ctx)
{
+ gcc_assert (gimple_code (ctx->stmt) != GIMPLE_OACC_PARALLEL);
+
tree x;
if (is_global_var (maybe_lookup_decl_in_outer_ctx (var, ctx)))
@@ -1484,6 +1488,7 @@ scan_sharing_clauses (tree clauses, omp_context *ctx)
switch (OMP_CLAUSE_CODE (c))
{
case OMP_CLAUSE_PRIVATE:
+ gcc_assert (gimple_code (ctx->stmt) != GIMPLE_OACC_PARALLEL);
decl = OMP_CLAUSE_DECL (c);
if (OMP_CLAUSE_PRIVATE_OUTER_REF (c))
goto do_private;
@@ -1492,6 +1497,7 @@ scan_sharing_clauses (tree clauses, omp_context *ctx)
break;
case OMP_CLAUSE_SHARED:
+ gcc_assert (gimple_code (ctx->stmt) != GIMPLE_OACC_PARALLEL);
/* Ignore shared directives in teams construct. */
if (gimple_code (ctx->stmt) == GIMPLE_OMP_TEAMS)
break;
@@ -1518,6 +1524,7 @@ scan_sharing_clauses (tree clauses, omp_context *ctx)
goto do_private;
case OMP_CLAUSE_LASTPRIVATE:
+ gcc_assert (gimple_code (ctx->stmt) != GIMPLE_OACC_PARALLEL);
/* Let the corresponding firstprivate clause create
the variable. */
if (OMP_CLAUSE_LASTPRIVATE_FIRSTPRIVATE (c))
@@ -1527,6 +1534,7 @@ scan_sharing_clauses (tree clauses, omp_context *ctx)
case OMP_CLAUSE_FIRSTPRIVATE:
case OMP_CLAUSE_REDUCTION:
case OMP_CLAUSE_LINEAR:
+ gcc_assert (gimple_code (ctx->stmt) != GIMPLE_OACC_PARALLEL);
decl = OMP_CLAUSE_DECL (c);
do_private:
if (is_variable_sized (decl))
@@ -1555,6 +1563,7 @@ scan_sharing_clauses (tree clauses, omp_context *ctx)
break;
case OMP_CLAUSE__LOOPTEMP_:
+ gcc_assert (gimple_code (ctx->stmt) != GIMPLE_OACC_PARALLEL);
gcc_assert (is_parallel_ctx (ctx));
decl = OMP_CLAUSE_DECL (c);
install_var_field (decl, false, 3, ctx);
@@ -1563,12 +1572,14 @@ scan_sharing_clauses (tree clauses, omp_context *ctx)
case OMP_CLAUSE_COPYPRIVATE:
case OMP_CLAUSE_COPYIN:
+ gcc_assert (gimple_code (ctx->stmt) != GIMPLE_OACC_PARALLEL);
decl = OMP_CLAUSE_DECL (c);
by_ref = use_pointer_for_field (decl, NULL);
install_var_field (decl, by_ref, 3, ctx);
break;
case OMP_CLAUSE_DEFAULT:
+ gcc_assert (gimple_code (ctx->stmt) != GIMPLE_OACC_PARALLEL);
ctx->default_kind = OMP_CLAUSE_DEFAULT_KIND (c);
break;
@@ -1581,6 +1592,7 @@ scan_sharing_clauses (tree clauses, omp_context *ctx)
case OMP_CLAUSE_SCHEDULE:
case OMP_CLAUSE_DIST_SCHEDULE:
case OMP_CLAUSE_DEPEND:
+ gcc_assert (gimple_code (ctx->stmt) != GIMPLE_OACC_PARALLEL);
if (ctx->outer)
scan_omp_op (&OMP_CLAUSE_OPERAND (c, 0), ctx->outer);
break;
@@ -1599,10 +1611,14 @@ scan_sharing_clauses (tree clauses, omp_context *ctx)
&& is_global_var (maybe_lookup_decl_in_outer_ctx (decl, ctx))
&& lookup_attribute ("omp declare target",
DECL_ATTRIBUTES (decl)))
+ {
+ gcc_assert (gimple_code (ctx->stmt) != GIMPLE_OACC_PARALLEL);
break;
+ }
if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
&& OMP_CLAUSE_MAP_KIND (c) == OMP_CLAUSE_MAP_POINTER)
{
+ gcc_assert (gimple_code (ctx->stmt) != GIMPLE_OACC_PARALLEL);
/* Ignore OMP_CLAUSE_MAP_POINTER kind for arrays in
#pragma omp target data, there is nothing to map for
those. */
@@ -1632,8 +1648,9 @@ scan_sharing_clauses (tree clauses, omp_context *ctx)
install_var_field (decl, true, 7, ctx);
else
install_var_field (decl, true, 3, ctx);
- if (gimple_omp_target_kind (ctx->stmt)
- == GF_OMP_TARGET_KIND_REGION)
+ if (gimple_code (ctx->stmt) == GIMPLE_OACC_PARALLEL
+ || (gimple_omp_target_kind (ctx->stmt)
+ == GF_OMP_TARGET_KIND_REGION))
install_var_local (decl, ctx);
}
}
@@ -1673,9 +1690,11 @@ scan_sharing_clauses (tree clauses, omp_context *ctx)
case OMP_CLAUSE_MERGEABLE:
case OMP_CLAUSE_PROC_BIND:
case OMP_CLAUSE_SAFELEN:
+ gcc_assert (gimple_code (ctx->stmt) != GIMPLE_OACC_PARALLEL);
break;
case OMP_CLAUSE_ALIGNED:
+ gcc_assert (gimple_code (ctx->stmt) != GIMPLE_OACC_PARALLEL);
decl = OMP_CLAUSE_DECL (c);
if (is_global_var (decl)
&& TREE_CODE (TREE_TYPE (decl)) == ARRAY_TYPE)
@@ -1692,6 +1711,7 @@ scan_sharing_clauses (tree clauses, omp_context *ctx)
switch (OMP_CLAUSE_CODE (c))
{
case OMP_CLAUSE_LASTPRIVATE:
+ gcc_assert (gimple_code (ctx->stmt) != GIMPLE_OACC_PARALLEL);
/* Let the corresponding firstprivate clause create
the variable. */
if (OMP_CLAUSE_LASTPRIVATE_GIMPLE_SEQ (c))
@@ -1704,6 +1724,7 @@ scan_sharing_clauses (tree clauses, omp_context *ctx)
case OMP_CLAUSE_FIRSTPRIVATE:
case OMP_CLAUSE_REDUCTION:
case OMP_CLAUSE_LINEAR:
+ gcc_assert (gimple_code (ctx->stmt) != GIMPLE_OACC_PARALLEL);
decl = OMP_CLAUSE_DECL (c);
if (is_variable_sized (decl))
install_var_local (decl, ctx);
@@ -1716,6 +1737,7 @@ scan_sharing_clauses (tree clauses, omp_context *ctx)
break;
case OMP_CLAUSE_SHARED:
+ gcc_assert (gimple_code (ctx->stmt) != GIMPLE_OACC_PARALLEL);
/* Ignore shared directives in teams construct. */
if (gimple_code (ctx->stmt) == GIMPLE_OMP_TEAMS)
break;
@@ -1725,14 +1747,18 @@ scan_sharing_clauses (tree clauses, omp_context *ctx)
break;
case OMP_CLAUSE_MAP:
- if (gimple_omp_target_kind (ctx->stmt) == GF_OMP_TARGET_KIND_DATA)
+ if (gimple_code (ctx->stmt) != GIMPLE_OACC_PARALLEL
+ && gimple_omp_target_kind (ctx->stmt) == GF_OMP_TARGET_KIND_DATA)
break;
decl = OMP_CLAUSE_DECL (c);
if (DECL_P (decl)
&& is_global_var (maybe_lookup_decl_in_outer_ctx (decl, ctx))
&& lookup_attribute ("omp declare target",
DECL_ATTRIBUTES (decl)))
+ {
+ gcc_assert (gimple_code (ctx->stmt) != GIMPLE_OACC_PARALLEL);
break;
+ }
if (DECL_P (decl))
{
if (OMP_CLAUSE_MAP_KIND (c) == OMP_CLAUSE_MAP_POINTER
@@ -1781,6 +1807,7 @@ scan_sharing_clauses (tree clauses, omp_context *ctx)
case OMP_CLAUSE__LOOPTEMP_:
case OMP_CLAUSE_TO:
case OMP_CLAUSE_FROM:
+ gcc_assert (gimple_code (ctx->stmt) != GIMPLE_OACC_PARALLEL);
break;
default:
@@ -1789,6 +1816,8 @@ scan_sharing_clauses (tree clauses, omp_context *ctx)
}
if (scan_array_reductions)
+ {
+ gcc_assert (gimple_code (ctx->stmt) != GIMPLE_OACC_PARALLEL);
for (c = clauses; c; c = OMP_CLAUSE_CHAIN (c))
if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_REDUCTION
&& OMP_CLAUSE_REDUCTION_PLACEHOLDER (c))
@@ -1799,6 +1828,7 @@ scan_sharing_clauses (tree clauses, omp_context *ctx)
else if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_LASTPRIVATE
&& OMP_CLAUSE_LASTPRIVATE_GIMPLE_SEQ (c))
scan_omp (&OMP_CLAUSE_LASTPRIVATE_GIMPLE_SEQ (c), ctx);
+ }
}
/* Create a new name for omp child function. Returns an identifier. */
@@ -1830,6 +1860,8 @@ create_omp_child_function (omp_context *ctx, bool task_copy)
decl = build_decl (gimple_location (ctx->stmt),
FUNCTION_DECL, name, type);
+ gcc_assert (gimple_code (ctx->stmt) != GIMPLE_OACC_PARALLEL
+ || !task_copy);
if (!task_copy)
ctx->cb.dst_fn = decl;
else
@@ -1861,6 +1893,8 @@ create_omp_child_function (omp_context *ctx, bool task_copy)
break;
}
}
+ gcc_assert (gimple_code (ctx->stmt) != GIMPLE_OACC_PARALLEL
+ || !target_p);
if (target_p)
DECL_ATTRIBUTES (decl)
= tree_cons (get_identifier ("omp declare target"),
@@ -1935,6 +1969,52 @@ find_combined_for (gimple_stmt_iterator *gsi_p,
return NULL;
}
+/* Scan an OpenACC parallel directive. */
+
+static void
+scan_oacc_parallel (gimple stmt, omp_context *outer_ctx)
+{
+ omp_context *ctx;
+ tree name;
+
+ gcc_assert (taskreg_nesting_level == 0);
+ gcc_assert (target_nesting_level == 0);
+
+ ctx = new_omp_context (stmt, outer_ctx);
+ ctx->field_map = splay_tree_new (splay_tree_compare_pointers, 0, 0);
+ ctx->default_kind = OMP_CLAUSE_DEFAULT_SHARED;
+ ctx->record_type = lang_hooks.types.make_type (RECORD_TYPE);
+ name = create_tmp_var_name (".omp_data_t");
+ name = build_decl (gimple_location (stmt),
+ TYPE_DECL, name, ctx->record_type);
+ DECL_ARTIFICIAL (name) = 1;
+ DECL_NAMELESS (name) = 1;
+ TYPE_NAME (ctx->record_type) = name;
+ create_omp_child_function (ctx, false);
+ gimple_oacc_parallel_set_child_fn (stmt, ctx->cb.dst_fn);
+
+ scan_sharing_clauses (gimple_oacc_parallel_clauses (stmt), ctx);
+ scan_omp (gimple_omp_body_ptr (stmt), ctx);
+
+ if (TYPE_FIELDS (ctx->record_type) == NULL)
+ ctx->record_type = ctx->receiver_decl = NULL;
+ else
+ {
+ TYPE_FIELDS (ctx->record_type)
+ = nreverse (TYPE_FIELDS (ctx->record_type));
+#ifdef ENABLE_CHECKING
+ tree field;
+ unsigned int align = DECL_ALIGN (TYPE_FIELDS (ctx->record_type));
+ for (field = TYPE_FIELDS (ctx->record_type);
+ field;
+ field = DECL_CHAIN (field))
+ gcc_assert (DECL_ALIGN (field) == align);
+#endif
+ layout_type (ctx->record_type);
+ fixup_child_record_type (ctx);
+ }
+}
+
/* Scan an OpenMP parallel directive. */
static void
@@ -2225,6 +2305,38 @@ scan_omp_teams (gimple stmt, omp_context *outer_ctx)
static bool
check_omp_nesting_restrictions (gimple stmt, omp_context *ctx)
{
+ omp_context *ctx_;
+
+ /* TODO: While the OpenACC specification does allow for certain kinds of
+ nesting, we don't support that yet. */
+ /* No nesting of STMT (which is an OpenACC or OpenMP one, or a GOMP builtin)
+ inside any OpenACC CTX. */
+ for (ctx_ = ctx; ctx_ != NULL; ctx_ = ctx_->outer)
+ switch (gimple_code (ctx_->stmt))
+ {
+ case GIMPLE_OACC_PARALLEL:
+ error_at (gimple_location (stmt),
+ "may not be nested");
+ return false;
+ default:
+ break;
+ }
+ /* No nesting of OpenACC STMT inside any OpenACC or OpenMP CTX. */
+ switch (gimple_code (stmt))
+ {
+ case GIMPLE_OACC_PARALLEL:
+ for (ctx_ = ctx; ctx_ != NULL; ctx_ = ctx_->outer)
+ if (is_gimple_omp (ctx_->stmt))
+ {
+ error_at (gimple_location (stmt),
+ "may not be nested");
+ return false;
+ }
+ break;
+ default:
+ break;
+ }
+
if (ctx != NULL)
{
if (gimple_code (ctx->stmt) == GIMPLE_OMP_FOR
@@ -2584,6 +2696,10 @@ scan_omp_1_stmt (gimple_stmt_iterator *gsi, bool *handled_ops_p,
switch (gimple_code (stmt))
{
+ case GIMPLE_OACC_PARALLEL:
+ scan_oacc_parallel (stmt, ctx);
+ break;
+
case GIMPLE_OMP_PARALLEL:
taskreg_nesting_level++;
scan_omp_parallel (gsi, ctx);
@@ -2910,6 +3026,8 @@ static bool
lower_rec_simd_input_clauses (tree new_var, omp_context *ctx, int &max_vf,
tree &idx, tree &lane, tree &ivar, tree &lvar)
{
+ gcc_assert (gimple_code (ctx->stmt) != GIMPLE_OACC_PARALLEL);
+
if (max_vf == 0)
{
max_vf = omp_max_vf ();
@@ -2959,6 +3077,8 @@ static void
lower_rec_input_clauses (tree clauses, gimple_seq *ilist, gimple_seq *dlist,
omp_context *ctx, struct omp_for_data *fd)
{
+ gcc_assert (gimple_code (ctx->stmt) != GIMPLE_OACC_PARALLEL);
+
tree c, dtor, copyin_seq, x, ptr;
bool copyin_by_ref = false;
bool lastprivate_firstprivate = false;
@@ -3617,6 +3737,8 @@ static void
lower_lastprivate_clauses (tree clauses, tree predicate, gimple_seq *stmt_list,
omp_context *ctx)
{
+ gcc_assert (gimple_code (ctx->stmt) != GIMPLE_OACC_PARALLEL);
+
tree x, c, label = NULL, orig_clauses = clauses;
bool par_clauses = false;
tree simduid = NULL, lastlane = NULL;
@@ -3752,6 +3874,8 @@ lower_lastprivate_clauses (tree clauses, tree predicate, gimple_seq *stmt_list,
static void
lower_reduction_clauses (tree clauses, gimple_seq *stmt_seqp, omp_context *ctx)
{
+ gcc_assert (gimple_code (ctx->stmt) != GIMPLE_OACC_PARALLEL);
+
gimple_seq sub_seq = NULL;
gimple stmt;
tree x, c;
@@ -3853,6 +3977,8 @@ static void
lower_copyprivate_clauses (tree clauses, gimple_seq *slist, gimple_seq *rlist,
omp_context *ctx)
{
+ gcc_assert (gimple_code (ctx->stmt) != GIMPLE_OACC_PARALLEL);
+
tree c;
for (c = clauses; c ; c = OMP_CLAUSE_CHAIN (c))
@@ -3903,6 +4029,8 @@ static void
lower_send_clauses (tree clauses, gimple_seq *ilist, gimple_seq *olist,
omp_context *ctx)
{
+ gcc_assert (gimple_code (ctx->stmt) != GIMPLE_OACC_PARALLEL);
+
tree c;
for (c = clauses; c ; c = OMP_CLAUSE_CHAIN (c))
@@ -3994,6 +4122,8 @@ lower_send_clauses (tree clauses, gimple_seq *ilist, gimple_seq *olist,
static void
lower_send_shared_vars (gimple_seq *ilist, gimple_seq *olist, omp_context *ctx)
{
+ gcc_assert (gimple_code (ctx->stmt) != GIMPLE_OACC_PARALLEL);
+
tree var, ovar, nvar, f, x, record_type;
if (ctx->record_type == NULL)
@@ -4542,6 +4672,234 @@ expand_omp_build_assign (gimple_stmt_iterator *gsi_p, tree to, tree from)
}
}
+/* Expand the OpenACC parallel directive starting at REGION. */
+
+static void
+expand_oacc_parallel (struct omp_region *region)
+{
+ basic_block entry_bb, exit_bb, new_bb;
+ struct function *child_cfun;
+ tree child_fn, block, t;
+ gimple_stmt_iterator gsi;
+ gimple entry_stmt, stmt;
+ edge e;
+
+ entry_stmt = last_stmt (region->entry);
+ child_fn = gimple_oacc_parallel_child_fn (entry_stmt);
+ child_cfun = DECL_STRUCT_FUNCTION (child_fn);
+
+ /* Supported by expand_omp_taskreg, but not here. */
+ gcc_assert (!child_cfun->cfg);
+ gcc_assert (!gimple_in_ssa_p (cfun));
+
+ entry_bb = region->entry;
+ exit_bb = region->exit;
+
+ /* Preserve indentation of expand_omp_target and expand_omp_taskreg. */
+ if (1)
+ {
+ unsigned srcidx, dstidx, num;
+
+ /* If the parallel region needs data sent from the parent
+ function, then the very first statement (except possible
+ tree profile counter updates) of the parallel body
+ is a copy assignment .OMP_DATA_I = &.OMP_DATA_O. Since
+ &.OMP_DATA_O is passed as an argument to the child function,
+ we need to replace it with the argument as seen by the child
+ function.
+
+ In most cases, this will end up being the identity assignment
+ .OMP_DATA_I = .OMP_DATA_I. However, if the parallel body had
+ a function call that has been inlined, the original PARM_DECL
+ .OMP_DATA_I may have been converted into a different local
+ variable. In which case, we need to keep the assignment. */
+ if (gimple_oacc_parallel_data_arg (entry_stmt))
+ {
+ basic_block entry_succ_bb = single_succ (entry_bb);
+ gimple_stmt_iterator gsi;
+ tree arg;
+ gimple parcopy_stmt = NULL;
+ tree sender
+ = TREE_VEC_ELT (gimple_oacc_parallel_data_arg (entry_stmt), 0);
+
+ for (gsi = gsi_start_bb (entry_succ_bb); ; gsi_next (&gsi))
+ {
+ gcc_assert (!gsi_end_p (gsi));
+ stmt = gsi_stmt (gsi);
+ if (gimple_code (stmt) != GIMPLE_ASSIGN)
+ continue;
+
+ if (gimple_num_ops (stmt) == 2)
+ {
+ tree arg = gimple_assign_rhs1 (stmt);
+
+ /* We're ignore the subcode because we're
+ effectively doing a STRIP_NOPS. */
+
+ if (TREE_CODE (arg) == ADDR_EXPR
+ && TREE_OPERAND (arg, 0) == sender)
+ {
+ parcopy_stmt = stmt;
+ break;
+ }
+ }
+ }
+
+ gcc_assert (parcopy_stmt != NULL);
+ arg = DECL_ARGUMENTS (child_fn);
+
+ gcc_assert (gimple_assign_lhs (parcopy_stmt) == arg);
+ gsi_remove (&gsi, true);
+ }
+
+ /* Declare local variables needed in CHILD_CFUN. */
+ block = DECL_INITIAL (child_fn);
+ BLOCK_VARS (block) = vec2chain (child_cfun->local_decls);
+ /* The gimplifier could record temporaries in the block
+ rather than in containing function's local_decls chain,
+ which would mean cgraph missed finalizing them. Do it now. */
+ for (t = BLOCK_VARS (block); t; t = DECL_CHAIN (t))
+ if (TREE_CODE (t) == VAR_DECL
+ && TREE_STATIC (t)
+ && !DECL_EXTERNAL (t))
+ varpool_finalize_decl (t);
+ DECL_SAVED_TREE (child_fn) = NULL;
+ /* We'll create a CFG for child_fn, so no gimple body is needed. */
+ gimple_set_body (child_fn, NULL);
+ TREE_USED (block) = 1;
+
+ /* Reset DECL_CONTEXT on function arguments. */
+ for (t = DECL_ARGUMENTS (child_fn); t; t = DECL_CHAIN (t))
+ DECL_CONTEXT (t) = child_fn;
+
+ /* Split ENTRY_BB at GIMPLE_OACC_PARALLEL,
+ so that it can be moved to the child function. */
+ gsi = gsi_last_bb (entry_bb);
+ stmt = gsi_stmt (gsi);
+ gcc_assert (stmt && (gimple_code (stmt) == GIMPLE_OACC_PARALLEL));
+ gsi_remove (&gsi, true);
+ e = split_block (entry_bb, stmt);
+ entry_bb = e->dest;
+ single_succ_edge (entry_bb)->flags = EDGE_FALLTHRU;
+
+ /* Convert GIMPLE_OMP_RETURN into a RETURN_EXPR. */
+ if (exit_bb)
+ {
+ gsi = gsi_last_bb (exit_bb);
+ gcc_assert (!gsi_end_p (gsi)
+ && gimple_code (gsi_stmt (gsi)) == GIMPLE_OMP_RETURN);
+ stmt = gimple_build_return (NULL);
+ gsi_insert_after (&gsi, stmt, GSI_SAME_STMT);
+ gsi_remove (&gsi, true);
+ }
+
+ /* Move the region into CHILD_CFUN. */
+
+ block = gimple_block (entry_stmt);
+
+ new_bb = move_sese_region_to_fn (child_cfun, entry_bb, exit_bb, block);
+ if (exit_bb)
+ single_succ_edge (new_bb)->flags = EDGE_FALLTHRU;
+ /* When the expansion process cannot guarantee an up-to-date
+ loop tree arrange for the child function to fixup loops. */
+ if (loops_state_satisfies_p (LOOPS_NEED_FIXUP))
+ child_cfun->x_current_loops->state |= LOOPS_NEED_FIXUP;
+
+ /* Remove non-local VAR_DECLs from child_cfun->local_decls list. */
+ num = vec_safe_length (child_cfun->local_decls);
+ for (srcidx = 0, dstidx = 0; srcidx < num; srcidx++)
+ {
+ t = (*child_cfun->local_decls)[srcidx];
+ if (DECL_CONTEXT (t) == cfun->decl)
+ continue;
+ if (srcidx != dstidx)
+ (*child_cfun->local_decls)[dstidx] = t;
+ dstidx++;
+ }
+ if (dstidx != num)
+ vec_safe_truncate (child_cfun->local_decls, dstidx);
+
+ /* Inform the callgraph about the new function. */
+ DECL_STRUCT_FUNCTION (child_fn)->curr_properties = cfun->curr_properties;
+ cgraph_add_new_function (child_fn, true);
+
+ /* Fix the callgraph edges for child_cfun. Those for cfun will be
+ fixed in a following pass. */
+ push_cfun (child_cfun);
+ rebuild_cgraph_edges ();
+
+ /* Some EH regions might become dead, see PR34608. If
+ pass_cleanup_cfg isn't the first pass to happen with the
+ new child, these dead EH edges might cause problems.
+ Clean them up now. */
+ if (flag_exceptions)
+ {
+ basic_block bb;
+ bool changed = false;
+
+ FOR_EACH_BB (bb)
+ changed |= gimple_purge_dead_eh_edges (bb);
+ if (changed)
+ cleanup_tree_cfg ();
+ }
+ pop_cfun ();
+ }
+
+ /* Emit a library call to launch CHILD_FN. */
+ tree t1, t2, t3, t4, device, c, clauses;
+ enum built_in_function start_ix;
+ location_t clause_loc;
+
+ clauses = gimple_oacc_parallel_clauses (entry_stmt);
+
+ start_ix = BUILT_IN_GOACC_PARALLEL;
+
+ /* By default, the value of DEVICE is -1 (let runtime library choose). */
+ device = build_int_cst (integer_type_node, -1);
+
+ c = find_omp_clause (clauses, OMP_CLAUSE_DEVICE);
+ gcc_assert (c == NULL);
+ if (c)
+ {
+ device = OMP_CLAUSE_DEVICE_ID (c);
+ clause_loc = OMP_CLAUSE_LOCATION (c);
+ }
+ else
+ clause_loc = gimple_location (entry_stmt);
+
+ /* Ensure 'device' is of the correct type. */
+ device = fold_convert_loc (clause_loc, integer_type_node, device);
+
+ gsi = gsi_last_bb (new_bb);
+ t = gimple_oacc_parallel_data_arg (entry_stmt);
+ if (t == NULL)
+ {
+ t1 = size_zero_node;
+ t2 = build_zero_cst (ptr_type_node);
+ t3 = t2;
+ t4 = t2;
+ }
+ else
+ {
+ t1 = TYPE_MAX_VALUE (TYPE_DOMAIN (TREE_TYPE (TREE_VEC_ELT (t, 1))));
+ t1 = size_binop (PLUS_EXPR, t1, size_int (1));
+ t2 = build_fold_addr_expr (TREE_VEC_ELT (t, 0));
+ t3 = build_fold_addr_expr (TREE_VEC_ELT (t, 1));
+ t4 = build_fold_addr_expr (TREE_VEC_ELT (t, 2));
+ }
+
+ gimple g;
+ /* FIXME: This will be address of
+ extern char __OPENMP_TARGET__[] __attribute__((visibility ("hidden")))
+ symbol, as soon as the linker plugin is able to create it for us. */
+ tree openmp_target = build_zero_cst (ptr_type_node);
+ tree fnaddr = build_fold_addr_expr (child_fn);
+ g = gimple_build_call (builtin_decl_explicit (start_ix),
+ 7, device, fnaddr, openmp_target, t1, t2, t3, t4);
+ gimple_set_location (g, gimple_location (entry_stmt));
+ gsi_insert_before (&gsi, g, GSI_SAME_STMT);
+}
+
/* Expand the OpenMP parallel or task directive starting at REGION. */
static void
@@ -8037,6 +8395,10 @@ expand_omp (struct omp_region *region)
switch (region->type)
{
+ case GIMPLE_OACC_PARALLEL:
+ expand_oacc_parallel (region);
+ break;
+
case GIMPLE_OMP_PARALLEL:
case GIMPLE_OMP_TASK:
expand_omp_taskreg (region);
@@ -8278,6 +8640,288 @@ make_pass_expand_omp (gcc::context *ctxt)
/* Routines to lower OpenMP directives into OMP-GIMPLE. */
+/* Lower the OpenACC parallel directive in the current statement
+ in GSI_P. CTX holds context information for the directive. */
+
+static void
+lower_oacc_parallel (gimple_stmt_iterator *gsi_p, omp_context *ctx)
+{
+ tree clauses;
+ tree child_fn, t, c;
+ gimple stmt = gsi_stmt (*gsi_p);
+ gimple par_bind, bind;
+ gimple_seq par_body, olist, ilist, new_body;
+ struct gimplify_ctx gctx;
+ location_t loc = gimple_location (stmt);
+ unsigned int map_cnt = 0;
+
+ clauses = gimple_oacc_parallel_clauses (stmt);
+ par_bind = gimple_seq_first_stmt (gimple_omp_body (stmt));
+ par_body = gimple_bind_body (par_bind);
+ child_fn = ctx->cb.dst_fn;
+
+ push_gimplify_context (&gctx);
+
+ for (c = clauses; c ; c = OMP_CLAUSE_CHAIN (c))
+ switch (OMP_CLAUSE_CODE (c))
+ {
+ tree var, x;
+
+ default:
+ break;
+ case OMP_CLAUSE_MAP:
+ case OMP_CLAUSE_TO:
+ case OMP_CLAUSE_FROM:
+ var = OMP_CLAUSE_DECL (c);
+ if (!DECL_P (var))
+ {
+ if (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_MAP
+ || !OMP_CLAUSE_MAP_ZERO_BIAS_ARRAY_SECTION (c))
+ map_cnt++;
+ continue;
+ }
+
+ if (DECL_SIZE (var)
+ && TREE_CODE (DECL_SIZE (var)) != INTEGER_CST)
+ {
+ tree var2 = DECL_VALUE_EXPR (var);
+ gcc_assert (TREE_CODE (var2) == INDIRECT_REF);
+ var2 = TREE_OPERAND (var2, 0);
+ gcc_assert (DECL_P (var2));
+ var = var2;
+ }
+
+ if (!maybe_lookup_field (var, ctx))
+ continue;
+
+ /* Preserve indentation of lower_omp_target. */
+ if (1)
+ {
+ x = build_receiver_ref (var, true, ctx);
+ tree new_var = lookup_decl (var, ctx);
+ if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
+ && OMP_CLAUSE_MAP_KIND (c) == OMP_CLAUSE_MAP_POINTER
+ && !OMP_CLAUSE_MAP_ZERO_BIAS_ARRAY_SECTION (c)
+ && TREE_CODE (TREE_TYPE (var)) == ARRAY_TYPE)
+ x = build_simple_mem_ref (x);
+ SET_DECL_VALUE_EXPR (new_var, x);
+ DECL_HAS_VALUE_EXPR_P (new_var) = 1;
+ }
+ map_cnt++;
+ }
+
+ target_nesting_level++;
+ lower_omp (&par_body, ctx);
+ target_nesting_level--;
+
+ /* Declare all the variables created by mapping and the variables
+ declared in the scope of the body. */
+ record_vars_into (ctx->block_vars, child_fn);
+ record_vars_into (gimple_bind_vars (par_bind), child_fn);
+
+ olist = NULL;
+ ilist = NULL;
+ if (ctx->record_type)
+ {
+ ctx->sender_decl
+ = create_tmp_var (ctx->record_type, ".omp_data_arr");
+ DECL_NAMELESS (ctx->sender_decl) = 1;
+ TREE_ADDRESSABLE (ctx->sender_decl) = 1;
+ t = make_tree_vec (3);
+ TREE_VEC_ELT (t, 0) = ctx->sender_decl;
+ TREE_VEC_ELT (t, 1)
+ = create_tmp_var (build_array_type_nelts (size_type_node, map_cnt),
+ ".omp_data_sizes");
+ DECL_NAMELESS (TREE_VEC_ELT (t, 1)) = 1;
+ TREE_ADDRESSABLE (TREE_VEC_ELT (t, 1)) = 1;
+ TREE_STATIC (TREE_VEC_ELT (t, 1)) = 1;
+ TREE_VEC_ELT (t, 2)
+ = create_tmp_var (build_array_type_nelts (unsigned_char_type_node,
+ map_cnt),
+ ".omp_data_kinds");
+ DECL_NAMELESS (TREE_VEC_ELT (t, 2)) = 1;
+ TREE_ADDRESSABLE (TREE_VEC_ELT (t, 2)) = 1;
+ TREE_STATIC (TREE_VEC_ELT (t, 2)) = 1;
+ gimple_oacc_parallel_set_data_arg (stmt, t);
+
+ vec<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;
+
+ for (c = clauses; c ; c = OMP_CLAUSE_CHAIN (c))
+ switch (OMP_CLAUSE_CODE (c))
+ {
+ tree ovar, nc;
+
+ default:
+ break;
+ case OMP_CLAUSE_MAP:
+ case OMP_CLAUSE_TO:
+ case OMP_CLAUSE_FROM:
+ nc = c;
+ ovar = OMP_CLAUSE_DECL (c);
+ if (!DECL_P (ovar))
+ {
+ if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
+ && OMP_CLAUSE_MAP_ZERO_BIAS_ARRAY_SECTION (c))
+ {
+ gcc_checking_assert (OMP_CLAUSE_DECL (OMP_CLAUSE_CHAIN (c))
+ == get_base_address (ovar));
+ nc = OMP_CLAUSE_CHAIN (c);
+ ovar = OMP_CLAUSE_DECL (nc);
+ }
+ else
+ {
+ tree x = build_sender_ref (ovar, ctx);
+ tree v
+ = build_fold_addr_expr_with_type (ovar, ptr_type_node);
+ gimplify_assign (x, v, &ilist);
+ nc = NULL_TREE;
+ }
+ }
+ else
+ {
+ if (DECL_SIZE (ovar)
+ && TREE_CODE (DECL_SIZE (ovar)) != INTEGER_CST)
+ {
+ tree ovar2 = DECL_VALUE_EXPR (ovar);
+ gcc_assert (TREE_CODE (ovar2) == INDIRECT_REF);
+ ovar2 = TREE_OPERAND (ovar2, 0);
+ gcc_assert (DECL_P (ovar2));
+ ovar = ovar2;
+ }
+ if (!maybe_lookup_field (ovar, ctx))
+ continue;
+ }
+
+ if (nc)
+ {
+ tree var = lookup_decl_in_outer_ctx (ovar, ctx);
+ tree x = build_sender_ref (ovar, ctx);
+ if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
+ && OMP_CLAUSE_MAP_KIND (c) == OMP_CLAUSE_MAP_POINTER
+ && !OMP_CLAUSE_MAP_ZERO_BIAS_ARRAY_SECTION (c)
+ && TREE_CODE (TREE_TYPE (ovar)) == ARRAY_TYPE)
+ {
+ tree avar
+ = create_tmp_var (TREE_TYPE (TREE_TYPE (x)), NULL);
+ mark_addressable (avar);
+ gimplify_assign (avar, build_fold_addr_expr (var), &ilist);
+ avar = build_fold_addr_expr (avar);
+ gimplify_assign (x, avar, &ilist);
+ }
+ else if (is_gimple_reg (var))
+ {
+ tree avar = create_tmp_var (TREE_TYPE (var), NULL);
+ mark_addressable (avar);
+ if (OMP_CLAUSE_MAP_KIND (c) != OMP_CLAUSE_MAP_ALLOC
+ && OMP_CLAUSE_MAP_KIND (c) != OMP_CLAUSE_MAP_FROM)
+ gimplify_assign (avar, var, &ilist);
+ avar = build_fold_addr_expr (avar);
+ gimplify_assign (x, avar, &ilist);
+ if ((OMP_CLAUSE_MAP_KIND (c) == OMP_CLAUSE_MAP_FROM
+ || OMP_CLAUSE_MAP_KIND (c) == OMP_CLAUSE_MAP_TOFROM)
+ && !TYPE_READONLY (TREE_TYPE (var)))
+ {
+ x = build_sender_ref (ovar, ctx);
+ x = build_simple_mem_ref (x);
+ gimplify_assign (var, x, &olist);
+ }
+ }
+ else
+ {
+ var = build_fold_addr_expr (var);
+ gimplify_assign (x, var, &ilist);
+ }
+ }
+ tree s = OMP_CLAUSE_SIZE (c);
+ if (s == NULL_TREE)
+ s = TYPE_SIZE_UNIT (TREE_TYPE (ovar));
+ s = fold_convert (size_type_node, s);
+ tree purpose = size_int (map_idx++);
+ CONSTRUCTOR_APPEND_ELT (vsize, purpose, s);
+ if (TREE_CODE (s) != INTEGER_CST)
+ TREE_STATIC (TREE_VEC_ELT (t, 1)) = 0;
+
+ unsigned char tkind = 0;
+ switch (OMP_CLAUSE_CODE (c))
+ {
+ case OMP_CLAUSE_MAP:
+ tkind = OMP_CLAUSE_MAP_KIND (c);
+ break;
+ case OMP_CLAUSE_TO:
+ tkind = OMP_CLAUSE_MAP_TO;
+ break;
+ case OMP_CLAUSE_FROM:
+ tkind = OMP_CLAUSE_MAP_FROM;
+ break;
+ default:
+ gcc_unreachable ();
+ }
+ unsigned int talign = TYPE_ALIGN_UNIT (TREE_TYPE (ovar));
+ if (DECL_P (ovar) && DECL_ALIGN_UNIT (ovar) > talign)
+ talign = DECL_ALIGN_UNIT (ovar);
+ talign = ceil_log2 (talign);
+ tkind |= talign << 3;
+ CONSTRUCTOR_APPEND_ELT (vkind, purpose,
+ build_int_cst (unsigned_char_type_node,
+ tkind));
+ if (nc && nc != c)
+ c = nc;
+ }
+
+ gcc_assert (map_idx == map_cnt);
+
+ DECL_INITIAL (TREE_VEC_ELT (t, 1))
+ = build_constructor (TREE_TYPE (TREE_VEC_ELT (t, 1)), vsize);
+ DECL_INITIAL (TREE_VEC_ELT (t, 2))
+ = build_constructor (TREE_TYPE (TREE_VEC_ELT (t, 2)), vkind);
+ if (!TREE_STATIC (TREE_VEC_ELT (t, 1)))
+ {
+ gimple_seq initlist = NULL;
+ force_gimple_operand (build1 (DECL_EXPR, void_type_node,
+ TREE_VEC_ELT (t, 1)),
+ &initlist, true, NULL_TREE);
+ gimple_seq_add_seq (&ilist, initlist);
+ }
+
+ tree clobber = build_constructor (ctx->record_type, NULL);
+ TREE_THIS_VOLATILE (clobber) = 1;
+ gimple_seq_add_stmt (&olist, gimple_build_assign (ctx->sender_decl,
+ clobber));
+ }
+
+ /* Once all the expansions are done, sequence all the different
+ fragments inside gimple_omp_body. */
+
+ new_body = NULL;
+
+ if (ctx->record_type)
+ {
+ t = build_fold_addr_expr_loc (loc, ctx->sender_decl);
+ /* fixup_child_record_type might have changed receiver_decl's type. */
+ t = fold_convert_loc (loc, TREE_TYPE (ctx->receiver_decl), t);
+ gimple_seq_add_stmt (&new_body,
+ gimple_build_assign (ctx->receiver_decl, t));
+ }
+
+ gimple_seq_add_seq (&new_body, par_body);
+ gcc_assert (!ctx->cancellable);
+ new_body = maybe_catch_exception (new_body);
+ gimple_seq_add_stmt (&new_body, gimple_build_omp_return (false));
+ gimple_omp_set_body (stmt, new_body);
+
+ bind = gimple_build_bind (NULL, NULL, gimple_bind_block (par_bind));
+ gsi_replace (gsi_p, bind, true);
+ gimple_bind_add_seq (bind, ilist);
+ gimple_bind_add_stmt (bind, stmt);
+ gimple_bind_add_seq (bind, olist);
+
+ pop_gimplify_context (NULL);
+}
+
/* If ctx is a worksharing context inside of a cancellable parallel
region and it isn't nowait, add lhs to its GIMPLE_OMP_RETURN
and conditional branch to parallel's cancel_label to handle
@@ -8286,6 +8930,8 @@ make_pass_expand_omp (gcc::context *ctxt)
static void
maybe_add_implicit_barrier_cancel (omp_context *ctx, gimple_seq *body)
{
+ gcc_assert (gimple_code (ctx->stmt) != GIMPLE_OACC_PARALLEL);
+
gimple omp_return = gimple_seq_last_stmt (*body);
gcc_assert (gimple_code (omp_return) == GIMPLE_OMP_RETURN);
if (gimple_omp_return_nowait_p (omp_return))
@@ -9051,6 +9697,8 @@ task_copyfn_remap_type (struct omp_taskcopy_context *tcctx, tree orig_type)
static void
create_task_copyfn (gimple task_stmt, omp_context *ctx)
{
+ gcc_assert (gimple_code (ctx->stmt) != GIMPLE_OACC_PARALLEL);
+
struct function *child_cfun;
tree child_fn, t, c, src, dst, f, sf, arg, sarg, decl;
tree record_type, srecord_type, bind, list;
@@ -9909,6 +10557,12 @@ lower_omp_1 (gimple_stmt_iterator *gsi_p, omp_context *ctx)
case GIMPLE_BIND:
lower_omp (gimple_bind_body_ptr (stmt), ctx);
break;
+ case GIMPLE_OACC_PARALLEL:
+ ctx = maybe_lookup_ctx (stmt);
+ gcc_assert (ctx);
+ gcc_assert (!ctx->cancellable);
+ lower_oacc_parallel (gsi_p, ctx);
+ break;
case GIMPLE_OMP_PARALLEL:
case GIMPLE_OMP_TASK:
ctx = maybe_lookup_ctx (stmt);
@@ -10357,6 +11011,7 @@ make_gimple_omp_edges (basic_block bb, struct omp_region **region)
switch (code)
{
+ case GIMPLE_OACC_PARALLEL:
case GIMPLE_OMP_PARALLEL:
case GIMPLE_OMP_TASK:
case GIMPLE_OMP_FOR:
new file mode 100644
@@ -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" } */
+ ;
+ }
+}
new file mode 100644
@@ -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" } */
+ ;
+ }
+}
new file mode 100644
@@ -0,0 +1,6 @@
+void
+foo (void)
+{
+#pragma acc parallel
+ foo ();
+}
new file mode 100644
@@ -0,0 +1,6 @@
+void
+foo (void)
+{
+#pragma acc parallel foo /* { dg-error "expected clause before 'foo'" } */
+ foo ();
+}
@@ -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:
@@ -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;
@@ -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);
@@ -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. */
@@ -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)
@@ -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
@@ -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@
@@ -232,4 +232,6 @@ OACC_2.0 {
};
GOACC_2.0 {
+ global:
+ GOACC_parallel;
};
@@ -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 */
new file mode 100644
@@ -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);
+}
new file mode 100644
@@ -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;
+}
new file mode 100644
@@ -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;
+}