@@ -1,3 +1,39 @@
+2014-02-28 Thomas Schwinge <thomas@codesourcery.com>
+
+ * gimple.def (GIMPLE_OACC_KERNELS): New code.
+ * doc/gimple.texi: Document it.
+ * gimple.h (gimple_has_substatements, CASE_GIMPLE_OMP)
+ (is_gimple_omp_oacc_specifically): Handle it.
+ (gimple_statement_oacc_kernels): New struct.
+ (gimple_build_oacc_kernels): New prototype.
+ (gimple_oacc_kernels_clauses, gimple_oacc_kernels_clauses_ptr)
+ (gimple_oacc_kernels_set_clauses, gimple_oacc_kernels_child_fn)
+ (gimple_oacc_kernels_child_fn_ptr)
+ (gimple_oacc_kernels_set_child_fn, gimple_oacc_kernels_data_arg)
+ (gimple_oacc_kernels_data_arg_ptr)
+ (gimple_oacc_kernels_set_data_arg): New inline functions.
+ * gimple.c (gimple_build_oacc_kernels): New function.
+ (gimple_copy): Handle GIMPLE_OACC_KERNELS.
+ * gimple-low.c (lower_stmt): Likewise.
+ * gimple-walk.c (walk_gimple_op, walk_gimple_stmt): Likewise.
+ * gimple-pretty-print.c (pp_gimple_stmt_1): Likewise.
+ (dump_gimple_oacc_parallel): Rename to dump_gimple_oacc_offload.
+ Also handle GIMPLE_OACC_KERNELS. Update all callers.
+ * gimplify.c (gimplify_omp_workshare, gimplify_expr): Handle
+ OACC_KERNELS.
+ * oacc-builtins.def (BUILT_IN_GOACC_KERNELS): New builtin.
+ * omp-low.c (scan_oacc_parallel, expand_oacc_parallel)
+ (lower_oacc_parallel): Rename to scan_oacc_offload,
+ expand_oacc_offload, and lower_oacc_offload. Also handle
+ GIMPLE_OACC_KERNELS. Update all callers.
+ (scan_sharing_clauses, scan_omp_1_stmt, expand_omp, lower_omp_1)
+ (diagnose_sb_0, diagnose_sb_1, diagnose_sb_2)
+ (make_gimple_omp_edges): Handle GIMPLE_OACC_KERNELS.
+ * tree-inline.c (remap_gimple_stmt, estimate_num_insns): Likewise.
+ * tree-nested.c (convert_nonlocal_reference_stmt)
+ (convert_local_reference_stmt, convert_tramp_reference_stmt)
+ (convert_gimple_call): Likewise.
+
2014-02-27 Thomas Schwinge <thomas@codesourcery.com>
* gimplify.c (gimplify_oacc_parallel): Merge into
@@ -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_KERNELS} @tab x @tab x
@item @code{GIMPLE_OACC_PARALLEL} @tab x @tab x
@item @code{GIMPLE_OMP_ATOMIC_LOAD} @tab x @tab x
@item @code{GIMPLE_OMP_ATOMIC_STORE} @tab x @tab x
@@ -906,6 +907,7 @@ Return a deep copy of statement @code{STMT}.
* @code{GIMPLE_EH_FILTER}::
* @code{GIMPLE_LABEL}::
* @code{GIMPLE_NOP}::
+* @code{GIMPLE_OACC_KERNELS}::
* @code{GIMPLE_OACC_PARALLEL}::
* @code{GIMPLE_OMP_ATOMIC_LOAD}::
* @code{GIMPLE_OMP_ATOMIC_STORE}::
@@ -1553,6 +1555,11 @@ Returns @code{TRUE} if statement @code{G} is a @code{GIMPLE_NOP}.
@end deftypefn
+@node @code{GIMPLE_OACC_KERNELS}
+@subsection @code{GIMPLE_OACC_KERNELS}
+@cindex @code{GIMPLE_OACC_KERNELS}
+
+
@node @code{GIMPLE_OACC_PARALLEL}
@subsection @code{GIMPLE_OACC_PARALLEL}
@cindex @code{GIMPLE_OACC_PARALLEL}
@@ -353,6 +353,7 @@ lower_stmt (gimple_stmt_iterator *gsi, struct lower_data *data)
}
break;
+ case GIMPLE_OACC_KERNELS:
case GIMPLE_OACC_PARALLEL:
case GIMPLE_OMP_PARALLEL:
case GIMPLE_OMP_TASK:
@@ -1843,36 +1843,57 @@ 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
+/* Dump an OpenACC offload tuple on the pretty_printer BUFFER, SPC spaces
of indent. FLAGS specifies details to show in the dump (see TDF_* in
dumpfile.h). */
static void
-dump_gimple_oacc_parallel (pretty_printer *buffer, gimple gs, int spc,
- int flags)
+dump_gimple_oacc_offload (pretty_printer *buffer, gimple gs, int spc,
+ int flags)
{
+ tree (*gimple_omp_clauses) (const_gimple);
+ tree (*gimple_omp_child_fn) (const_gimple);
+ tree (*gimple_omp_data_arg) (const_gimple);
+ const char *kind;
+ switch (gimple_code (gs))
+ {
+ case GIMPLE_OACC_KERNELS:
+ gimple_omp_clauses = gimple_oacc_kernels_clauses;
+ gimple_omp_child_fn = gimple_oacc_kernels_child_fn;
+ gimple_omp_data_arg = gimple_oacc_kernels_data_arg;
+ kind = "kernels";
+ break;
+ case GIMPLE_OACC_PARALLEL:
+ gimple_omp_clauses = gimple_oacc_parallel_clauses;
+ gimple_omp_child_fn = gimple_oacc_parallel_child_fn;
+ gimple_omp_data_arg = gimple_oacc_parallel_data_arg;
+ kind = "parallel";
+ break;
+ default:
+ gcc_unreachable ();
+ }
if (flags & TDF_RAW)
{
dump_gimple_fmt (buffer, spc, flags, "%G <%+BODY <%S>%nCLAUSES <", gs,
gimple_omp_body (gs));
- dump_omp_clauses (buffer, gimple_oacc_parallel_clauses (gs), spc, flags);
+ dump_omp_clauses (buffer, gimple_omp_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));
+ gimple_omp_child_fn (gs), gimple_omp_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, "#pragma acc ");
+ pp_string (buffer, kind);
+ dump_omp_clauses (buffer, gimple_omp_clauses (gs), spc, flags);
+ if (gimple_omp_child_fn (gs))
{
pp_string (buffer, " [child fn: ");
- dump_generic_node (buffer, gimple_oacc_parallel_child_fn (gs),
+ dump_generic_node (buffer, gimple_omp_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),
+ if (gimple_omp_data_arg (gs))
+ dump_generic_node (buffer, gimple_omp_data_arg (gs),
spc, flags, false);
else
pp_string (buffer, "???");
@@ -2193,8 +2214,9 @@ pp_gimple_stmt_1 (pretty_printer *buffer, gimple gs, int spc, int flags)
dump_gimple_phi (buffer, gs, spc, false, flags);
break;
+ case GIMPLE_OACC_KERNELS:
case GIMPLE_OACC_PARALLEL:
- dump_gimple_oacc_parallel (buffer, gs, spc, flags);
+ dump_gimple_oacc_offload (buffer, gs, spc, flags);
break;
case GIMPLE_OMP_PARALLEL:
@@ -296,6 +296,21 @@ walk_gimple_op (gimple stmt, walk_tree_fn callback_op,
return ret;
break;
+ case GIMPLE_OACC_KERNELS:
+ ret = walk_tree (gimple_oacc_kernels_clauses_ptr (stmt), callback_op,
+ wi, pset);
+ if (ret)
+ return ret;
+ ret = walk_tree (gimple_oacc_kernels_child_fn_ptr (stmt), callback_op,
+ wi, pset);
+ if (ret)
+ return ret;
+ ret = walk_tree (gimple_oacc_kernels_data_arg_ptr (stmt), callback_op,
+ wi, pset);
+ if (ret)
+ return ret;
+ break;
+
case GIMPLE_OACC_PARALLEL:
ret = walk_tree (gimple_oacc_parallel_clauses_ptr (stmt), callback_op,
wi, pset);
@@ -606,6 +621,7 @@ walk_gimple_stmt (gimple_stmt_iterator *gsi, walk_stmt_fn callback_stmt,
return wi->callback_result;
/* FALL THROUGH. */
+ case GIMPLE_OACC_KERNELS:
case GIMPLE_OACC_PARALLEL:
case GIMPLE_OMP_CRITICAL:
case GIMPLE_OMP_MASTER:
@@ -799,6 +799,23 @@ gimple_build_debug_source_bind_stat (tree var, tree value,
}
+/* Build a GIMPLE_OACC_KERNELS statement.
+
+ BODY is sequence of statements which are executed as kernels.
+ CLAUSES are the OpenACC kernels construct's clauses. */
+
+gimple
+gimple_build_oacc_kernels (gimple_seq body, tree clauses)
+{
+ gimple p = gimple_alloc (GIMPLE_OACC_KERNELS, 0);
+ if (body)
+ gimple_omp_set_body (p, body);
+ gimple_oacc_kernels_set_clauses (p, clauses);
+
+ return p;
+}
+
+
/* Build a GIMPLE_OACC_PARALLEL statement.
BODY is sequence of statements which are executed in parallel.
@@ -1672,6 +1689,7 @@ gimple_copy (gimple stmt)
gimple_try_set_cleanup (copy, new_seq);
break;
+ case GIMPLE_OACC_KERNELS:
case GIMPLE_OACC_PARALLEL:
gcc_unreachable ();
@@ -209,10 +209,28 @@ DEFGSCODE(GIMPLE_NOP, "gimple_nop", GSS_BASE)
ordering is exposed by the range check in gimple_omp_subcode. */
+/* GIMPLE_OACC_KERNELS <BODY, CLAUSES, CHILD_FN, DATA_ARG> represents
+ #pragma acc kernels [CLAUSES]
+ BODY is the sequence of statements inside the kernels construct.
+ CLAUSES is an OMP_CLAUSE chain holding the associated clauses.
+ CHILD_FN is set when outlining the body of the kernels region.
+ All the statements in BODY are moved into this newly created
+ function when converting OMP constructs into low-GIMPLE.
+ DATA_ARG is a vec of 3 local variables in the parent function
+ containing data to be mapped to CHILD_FN. This is used to
+ implement the MAP clauses. */
+DEFGSCODE(GIMPLE_OACC_KERNELS, "gimple_oacc_kernels", GSS_OMP_PARALLEL_LAYOUT)
+
/* GIMPLE_OACC_PARALLEL <BODY, CLAUSES, CHILD_FN, DATA_ARG> represents
-
#pragma acc parallel [CLAUSES]
- BODY */
+ BODY is the sequence of statements inside the parallel construct.
+ CLAUSES is an OMP_CLAUSE chain holding the associated clauses.
+ CHILD_FN is set when outlining the body of the parallel region.
+ All the statements in BODY are moved into this newly created
+ function when converting OMP constructs into low-GIMPLE.
+ DATA_ARG is a vec of 3 local variables in the parent function
+ containing data to be mapped to CHILD_FN. This is used to
+ implement the MAP clauses. */
DEFGSCODE(GIMPLE_OACC_PARALLEL, "gimple_oacc_parallel", GSS_OMP_PARALLEL_LAYOUT)
/* Tuples used for lowering of OMP_ATOMIC. Although the form of the OMP_ATOMIC
@@ -549,8 +549,8 @@ struct GTY((tag("GSS_OMP_FOR")))
};
-/* GIMPLE_OACC_PARALLEL, GIMPLE_OMP_PARALLEL, GIMPLE_OMP_TARGET,
- GIMPLE_OMP_TASK */
+/* GIMPLE_OACC_KERNELS, GIMPLE_OACC_PARALLEL, GIMPLE_OMP_PARALLEL,
+ GIMPLE_OMP_TARGET, GIMPLE_OMP_TASK */
struct GTY((tag("GSS_OMP_PARALLEL_LAYOUT")))
gimple_statement_omp_parallel_layout : public gimple_statement_omp
{
@@ -569,6 +569,14 @@ struct GTY((tag("GSS_OMP_PARALLEL_LAYOUT")))
tree data_arg;
};
+/* GIMPLE_OACC_KERNELS */
+struct GTY((tag("GSS_OMP_PARALLEL_LAYOUT")))
+ gimple_statement_oacc_kernels : public gimple_statement_omp_parallel_layout
+{
+ /* No extra fields; adds invariant:
+ stmt->code == GIMPLE_OACC_KERNELS. */
+};
+
/* GIMPLE_OACC_PARALLEL */
struct GTY((tag("GSS_OMP_PARALLEL_LAYOUT")))
gimple_statement_oacc_parallel : public gimple_statement_omp_parallel_layout
@@ -894,6 +902,14 @@ is_a_helper <gimple_statement_omp_for>::test (gimple gs)
template <>
template <>
inline bool
+is_a_helper <gimple_statement_oacc_kernels>::test (gimple gs)
+{
+ return gs->code == GIMPLE_OACC_KERNELS;
+}
+
+template <>
+template <>
+inline bool
is_a_helper <gimple_statement_oacc_parallel>::test (gimple gs)
{
return gs->code == GIMPLE_OACC_PARALLEL;
@@ -1094,6 +1110,14 @@ is_a_helper <const gimple_statement_omp_for>::test (const_gimple gs)
template <>
template <>
inline bool
+is_a_helper <const gimple_statement_oacc_kernels>::test (const_gimple gs)
+{
+ return gs->code == GIMPLE_OACC_KERNELS;
+}
+
+template <>
+template <>
+inline bool
is_a_helper <const gimple_statement_oacc_parallel>::test (const_gimple gs)
{
return gs->code == GIMPLE_OACC_PARALLEL;
@@ -1225,6 +1249,7 @@ gimple gimple_build_debug_bind_stat (tree, tree, gimple MEM_STAT_DECL);
gimple gimple_build_debug_source_bind_stat (tree, tree, gimple MEM_STAT_DECL);
#define gimple_build_debug_source_bind(var,val,stmt) \
gimple_build_debug_source_bind_stat ((var), (val), (stmt) MEM_STAT_INFO)
+gimple gimple_build_oacc_kernels (gimple_seq, tree);
gimple gimple_build_oacc_parallel (gimple_seq, tree);
gimple gimple_build_omp_critical (gimple_seq, tree);
gimple gimple_build_omp_for (gimple_seq, int, tree, size_t, gimple_seq);
@@ -1462,6 +1487,7 @@ gimple_has_substatements (gimple g)
case GIMPLE_EH_FILTER:
case GIMPLE_EH_ELSE:
case GIMPLE_TRY:
+ case GIMPLE_OACC_KERNELS:
case GIMPLE_OACC_PARALLEL:
case GIMPLE_OMP_FOR:
case GIMPLE_OMP_MASTER:
@@ -4266,6 +4292,101 @@ gimple_omp_set_body (gimple gs, gimple_seq body)
}
+/* Return the clauses associated with OACC_KERNELS statement GS. */
+
+static inline tree
+gimple_oacc_kernels_clauses (const_gimple gs)
+{
+ const gimple_statement_oacc_kernels *oacc_kernels_stmt =
+ as_a <const gimple_statement_oacc_kernels> (gs);
+ return oacc_kernels_stmt->clauses;
+}
+
+/* Return a pointer to the clauses associated with OACC_KERNELS statement GS. */
+
+static inline tree *
+gimple_oacc_kernels_clauses_ptr (gimple gs)
+{
+ gimple_statement_oacc_kernels *oacc_kernels_stmt =
+ as_a <gimple_statement_oacc_kernels> (gs);
+ return &oacc_kernels_stmt->clauses;
+}
+
+/* Set CLAUSES to be the list of clauses associated with OACC_KERNELS statement
+ GS. */
+
+static inline void
+gimple_oacc_kernels_set_clauses (gimple gs, tree clauses)
+{
+ gimple_statement_oacc_kernels *oacc_kernels_stmt =
+ as_a <gimple_statement_oacc_kernels> (gs);
+ oacc_kernels_stmt->clauses = clauses;
+}
+
+/* Return the child function used to hold the body of OACC_KERNELS statement
+ GS. */
+
+static inline tree
+gimple_oacc_kernels_child_fn (const_gimple gs)
+{
+ const gimple_statement_oacc_kernels *oacc_kernels_stmt =
+ as_a <const gimple_statement_oacc_kernels> (gs);
+ return oacc_kernels_stmt->child_fn;
+}
+
+/* Return a pointer to the child function used to hold the body of OACC_KERNELS
+ statement GS. */
+
+static inline tree *
+gimple_oacc_kernels_child_fn_ptr (gimple gs)
+{
+ gimple_statement_oacc_kernels *oacc_kernels_stmt =
+ as_a <gimple_statement_oacc_kernels> (gs);
+ return &oacc_kernels_stmt->child_fn;
+}
+
+/* Set CHILD_FN to be the child function for OACC_KERNELS statement GS. */
+
+static inline void
+gimple_oacc_kernels_set_child_fn (gimple gs, tree child_fn)
+{
+ gimple_statement_oacc_kernels *oacc_kernels_stmt =
+ as_a <gimple_statement_oacc_kernels> (gs);
+ oacc_kernels_stmt->child_fn = child_fn;
+}
+
+/* Return the artificial argument used to send variables and values
+ from the parent to the children threads in OACC_KERNELS statement GS. */
+
+static inline tree
+gimple_oacc_kernels_data_arg (const_gimple gs)
+{
+ const gimple_statement_oacc_kernels *oacc_kernels_stmt =
+ as_a <const gimple_statement_oacc_kernels> (gs);
+ return oacc_kernels_stmt->data_arg;
+}
+
+/* Return a pointer to the data argument for OACC_KERNELS statement GS. */
+
+static inline tree *
+gimple_oacc_kernels_data_arg_ptr (gimple gs)
+{
+ gimple_statement_oacc_kernels *oacc_kernels_stmt =
+ as_a <gimple_statement_oacc_kernels> (gs);
+ return &oacc_kernels_stmt->data_arg;
+}
+
+/* Set DATA_ARG to be the data argument for OACC_KERNELS statement GS. */
+
+static inline void
+gimple_oacc_kernels_set_data_arg (gimple gs, tree data_arg)
+{
+ gimple_statement_oacc_kernels *oacc_kernels_stmt =
+ as_a <gimple_statement_oacc_kernels> (gs);
+ oacc_kernels_stmt->data_arg = data_arg;
+}
+
+
/* Return the clauses associated with OACC_PARALLEL statement GS. */
static inline tree
@@ -4330,7 +4451,8 @@ gimple_oacc_parallel_set_child_fn (gimple gs, tree child_fn)
oacc_parallel_stmt->child_fn = child_fn;
}
-/* Return the data argument for OACC_PARALLEL statement GS. */
+/* Return the artificial argument used to send variables and values
+ from the parent to the children threads in OACC_PARALLEL statement GS. */
static inline tree
gimple_oacc_parallel_data_arg (const_gimple gs)
@@ -5640,6 +5762,7 @@ gimple_return_set_retval (gimple gs, tree retval)
/* Returns true when the gimple statement STMT is any of the OpenMP types. */
#define CASE_GIMPLE_OMP \
+ case GIMPLE_OACC_KERNELS: \
case GIMPLE_OACC_PARALLEL: \
case GIMPLE_OMP_PARALLEL: \
case GIMPLE_OMP_TASK: \
@@ -5683,6 +5806,7 @@ is_gimple_omp_oacc_specifically (const_gimple stmt)
gcc_assert (is_gimple_omp (stmt));
switch (gimple_code (stmt))
{
+ case GIMPLE_OACC_KERNELS:
case GIMPLE_OACC_PARALLEL:
return true;
case GIMPLE_OMP_TARGET:
@@ -7007,6 +7007,7 @@ gimplify_omp_workshare (tree *expr_p, gimple_seq *pre_p)
ort = (enum omp_region_type) (ORT_TARGET
| ORT_TARGET_MAP_FORCE);
break;
+ case OACC_KERNELS:
case OACC_PARALLEL:
ort = (enum omp_region_type) (ORT_TARGET
| ORT_TARGET_OFFLOAD
@@ -7070,6 +7071,9 @@ gimplify_omp_workshare (tree *expr_p, gimple_seq *pre_p)
stmt = gimple_build_omp_target (body, GF_OMP_TARGET_KIND_OACC_DATA,
OACC_DATA_CLAUSES (expr));
break;
+ case OACC_KERNELS:
+ stmt = gimple_build_oacc_kernels (body, OACC_KERNELS_CLAUSES (expr));
+ break;
case OACC_PARALLEL:
stmt = gimple_build_oacc_parallel (body, OACC_PARALLEL_CLAUSES (expr));
break;
@@ -8036,7 +8040,6 @@ gimplify_expr (tree *expr_p, gimple_seq *pre_p, gimple_seq *post_p,
ret = GS_ALL_DONE;
break;
- case OACC_KERNELS:
case OACC_HOST_DATA:
case OACC_DECLARE:
case OACC_UPDATE:
@@ -8066,6 +8069,7 @@ gimplify_expr (tree *expr_p, gimple_seq *pre_p, gimple_seq *post_p,
break;
case OACC_DATA:
+ case OACC_KERNELS:
case OACC_PARALLEL:
case OMP_SECTIONS:
case OMP_SINGLE:
@@ -27,9 +27,11 @@ along with GCC; see the file COPYING3. If not see
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)
DEF_GOACC_BUILTIN (BUILT_IN_GOACC_DATA_START, "GOACC_data_start",
BT_FN_VOID_INT_PTR_SIZE_PTR_PTR_PTR, ATTR_NOTHROW_LIST)
DEF_GOACC_BUILTIN (BUILT_IN_GOACC_DATA_END, "GOACC_data_end",
BT_FN_VOID, ATTR_NOTHROW_LIST)
+DEF_GOACC_BUILTIN (BUILT_IN_GOACC_KERNELS, "GOACC_kernels",
+ BT_FN_VOID_INT_OMPFN_PTR_SIZE_PTR_PTR_PTR, ATTR_NOTHROW_LIST)
+DEF_GOACC_BUILTIN (BUILT_IN_GOACC_PARALLEL, "GOACC_parallel",
+ BT_FN_VOID_INT_OMPFN_PTR_SIZE_PTR_PTR_PTR, ATTR_NOTHROW_LIST)
@@ -1502,6 +1502,7 @@ scan_sharing_clauses (tree clauses, omp_context *ctx)
bool offloaded;
switch (gimple_code (ctx->stmt))
{
+ case GIMPLE_OACC_KERNELS:
case GIMPLE_OACC_PARALLEL:
offloaded = true;
break;
@@ -2085,13 +2086,28 @@ find_combined_for (gimple_stmt_iterator *gsi_p,
return NULL;
}
-/* Scan an OpenACC parallel directive. */
+/* Scan an OpenACC offload directive. */
static void
-scan_oacc_parallel (gimple stmt, omp_context *outer_ctx)
+scan_oacc_offload (gimple stmt, omp_context *outer_ctx)
{
omp_context *ctx;
tree name;
+ void (*gimple_omp_set_child_fn) (gimple, tree);
+ tree (*gimple_omp_clauses) (const_gimple);
+ switch (gimple_code (stmt))
+ {
+ case GIMPLE_OACC_KERNELS:
+ gimple_omp_set_child_fn = gimple_oacc_kernels_set_child_fn;
+ gimple_omp_clauses = gimple_oacc_kernels_clauses;
+ break;
+ case GIMPLE_OACC_PARALLEL:
+ gimple_omp_set_child_fn = gimple_oacc_parallel_set_child_fn;
+ gimple_omp_clauses = gimple_oacc_parallel_clauses;
+ break;
+ default:
+ gcc_unreachable ();
+ }
gcc_assert (taskreg_nesting_level == 0);
gcc_assert (target_nesting_level == 0);
@@ -2107,9 +2123,10 @@ scan_oacc_parallel (gimple stmt, omp_context *outer_ctx)
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);
+ gimple_omp_set_child_fn (stmt, ctx->cb.dst_fn);
+
+ scan_sharing_clauses (gimple_omp_clauses (stmt), ctx);
scan_omp (gimple_omp_body_ptr (stmt), ctx);
if (TYPE_FIELDS (ctx->record_type) == NULL)
@@ -2841,8 +2858,9 @@ scan_omp_1_stmt (gimple_stmt_iterator *gsi, bool *handled_ops_p,
switch (gimple_code (stmt))
{
+ case GIMPLE_OACC_KERNELS:
case GIMPLE_OACC_PARALLEL:
- scan_oacc_parallel (stmt, ctx);
+ scan_oacc_offload (stmt, ctx);
break;
case GIMPLE_OMP_PARALLEL:
@@ -4860,10 +4878,10 @@ expand_omp_build_assign (gimple_stmt_iterator *gsi_p, tree to, tree from)
}
}
-/* Expand the OpenACC parallel directive starting at REGION. */
+/* Expand the OpenACC offload directive starting at REGION. */
static void
-expand_oacc_parallel (struct omp_region *region)
+expand_oacc_offload (struct omp_region *region)
{
basic_block entry_bb, exit_bb, new_bb;
struct function *child_cfun;
@@ -4871,9 +4889,24 @@ expand_oacc_parallel (struct omp_region *region)
gimple_stmt_iterator gsi;
gimple entry_stmt, stmt;
edge e;
+ tree (*gimple_omp_child_fn) (const_gimple);
+ tree (*gimple_omp_data_arg) (const_gimple);
+ switch (region->type)
+ {
+ case GIMPLE_OACC_KERNELS:
+ gimple_omp_child_fn = gimple_oacc_kernels_child_fn;
+ gimple_omp_data_arg = gimple_oacc_kernels_data_arg;
+ break;
+ case GIMPLE_OACC_PARALLEL:
+ gimple_omp_child_fn = gimple_oacc_parallel_child_fn;
+ gimple_omp_data_arg = gimple_oacc_parallel_data_arg;
+ break;
+ default:
+ gcc_unreachable ();
+ }
entry_stmt = last_stmt (region->entry);
- child_fn = gimple_oacc_parallel_child_fn (entry_stmt);
+ child_fn = gimple_omp_child_fn (entry_stmt);
child_cfun = DECL_STRUCT_FUNCTION (child_fn);
/* Supported by expand_omp_taskreg, but not here. */
@@ -4901,14 +4934,13 @@ expand_oacc_parallel (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_oacc_parallel_data_arg (entry_stmt))
+ if (gimple_omp_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);
+ tree sender = TREE_VEC_ELT (gimple_omp_data_arg (entry_stmt), 0);
for (gsi = gsi_start_bb (entry_succ_bb); ; gsi_next (&gsi))
{
@@ -4964,7 +4996,8 @@ expand_oacc_parallel (struct omp_region *region)
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));
+ gcc_assert (stmt && (gimple_code (stmt) == GIMPLE_OACC_KERNELS
+ || gimple_code (stmt) == GIMPLE_OACC_PARALLEL));
gsi_remove (&gsi, true);
e = split_block (entry_bb, stmt);
entry_bb = e->dest;
@@ -5037,10 +5070,22 @@ expand_oacc_parallel (struct omp_region *region)
tree t1, t2, t3, t4, device, c, clauses;
enum built_in_function start_ix;
location_t clause_loc;
+ tree (*gimple_omp_clauses) (const_gimple);
+ switch (region->type)
+ {
+ case GIMPLE_OACC_KERNELS:
+ gimple_omp_clauses = gimple_oacc_kernels_clauses;
+ start_ix = BUILT_IN_GOACC_KERNELS;
+ break;
+ case GIMPLE_OACC_PARALLEL:
+ gimple_omp_clauses = gimple_oacc_parallel_clauses;
+ start_ix = BUILT_IN_GOACC_PARALLEL;
+ break;
+ default:
+ gcc_unreachable ();
+ }
- clauses = gimple_oacc_parallel_clauses (entry_stmt);
-
- start_ix = BUILT_IN_GOACC_PARALLEL;
+ clauses = gimple_omp_clauses (entry_stmt);
/* By default, the value of DEVICE is -1 (let runtime library choose). */
device = build_int_cst (integer_type_node, -1);
@@ -5059,7 +5104,7 @@ expand_oacc_parallel (struct omp_region *region)
device = fold_convert_loc (clause_loc, integer_type_node, device);
gsi = gsi_last_bb (new_bb);
- t = gimple_oacc_parallel_data_arg (entry_stmt);
+ t = gimple_omp_data_arg (entry_stmt);
if (t == NULL)
{
t1 = size_zero_node;
@@ -8606,8 +8651,9 @@ expand_omp (struct omp_region *region)
switch (region->type)
{
+ case GIMPLE_OACC_KERNELS:
case GIMPLE_OACC_PARALLEL:
- expand_oacc_parallel (region);
+ expand_oacc_offload (region);
break;
case GIMPLE_OMP_PARALLEL:
@@ -8851,11 +8897,11 @@ make_pass_expand_omp (gcc::context *ctxt)
/* Routines to lower OpenMP directives into OMP-GIMPLE. */
-/* Lower the OpenACC parallel directive in the current statement
+/* Lower the OpenACC offload 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)
+lower_oacc_offload (gimple_stmt_iterator *gsi_p, omp_context *ctx)
{
tree clauses;
tree child_fn, t, c;
@@ -8864,8 +8910,23 @@ lower_oacc_parallel (gimple_stmt_iterator *gsi_p, omp_context *ctx)
gimple_seq par_body, olist, ilist, new_body;
location_t loc = gimple_location (stmt);
unsigned int map_cnt = 0;
+ tree (*gimple_omp_clauses) (const_gimple);
+ void (*gimple_omp_set_data_arg) (gimple, tree);
+ switch (gimple_code (stmt))
+ {
+ case GIMPLE_OACC_KERNELS:
+ gimple_omp_clauses = gimple_oacc_kernels_clauses;
+ gimple_omp_set_data_arg = gimple_oacc_kernels_set_data_arg;
+ break;
+ case GIMPLE_OACC_PARALLEL:
+ gimple_omp_clauses = gimple_oacc_parallel_clauses;
+ gimple_omp_set_data_arg = gimple_oacc_parallel_set_data_arg;
+ break;
+ default:
+ gcc_unreachable ();
+ }
- clauses = gimple_oacc_parallel_clauses (stmt);
+ clauses = gimple_omp_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;
@@ -8950,7 +9011,7 @@ lower_oacc_parallel (gimple_stmt_iterator *gsi_p, omp_context *ctx)
DECL_NAMELESS (TREE_VEC_ELT (t, 2)) = 1;
TREE_ADDRESSABLE (TREE_VEC_ELT (t, 2)) = 1;
TREE_STATIC (TREE_VEC_ELT (t, 2)) = 1;
- gimple_oacc_parallel_set_data_arg (stmt, t);
+ gimple_omp_set_data_arg (stmt, t);
vec<constructor_elt, va_gc> *vsize;
vec<constructor_elt, va_gc> *vkind;
@@ -10820,11 +10881,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_KERNELS:
case GIMPLE_OACC_PARALLEL:
ctx = maybe_lookup_ctx (stmt);
gcc_assert (ctx);
gcc_assert (!ctx->cancellable);
- lower_oacc_parallel (gsi_p, ctx);
+ lower_oacc_offload (gsi_p, ctx);
break;
case GIMPLE_OMP_PARALLEL:
case GIMPLE_OMP_TASK:
@@ -11053,6 +11115,9 @@ static bool
diagnose_sb_0 (gimple_stmt_iterator *gsi_p,
gimple branch_ctx, gimple label_ctx)
{
+ gcc_assert (!branch_ctx || is_gimple_omp (branch_ctx));
+ gcc_assert (!label_ctx || is_gimple_omp (label_ctx));
+
if (label_ctx == branch_ctx)
return false;
@@ -11070,8 +11135,8 @@ diagnose_sb_0 (gimple_stmt_iterator *gsi_p,
}
if (flag_openacc)
{
- if ((branch_ctx && gimple_code (branch_ctx) == GIMPLE_OACC_PARALLEL)
- || (label_ctx && gimple_code (label_ctx) == GIMPLE_OACC_PARALLEL))
+ if ((branch_ctx && is_gimple_omp_oacc_specifically (branch_ctx))
+ || (label_ctx && is_gimple_omp_oacc_specifically (label_ctx)))
{
gcc_assert (kind == NULL);
kind = "OpenACC";
@@ -11149,6 +11214,7 @@ diagnose_sb_1 (gimple_stmt_iterator *gsi_p, bool *handled_ops_p,
{
WALK_SUBSTMTS;
+ case GIMPLE_OACC_KERNELS:
case GIMPLE_OACC_PARALLEL:
case GIMPLE_OMP_PARALLEL:
case GIMPLE_OMP_TASK:
@@ -11208,6 +11274,7 @@ diagnose_sb_2 (gimple_stmt_iterator *gsi_p, bool *handled_ops_p,
{
WALK_SUBSTMTS;
+ case GIMPLE_OACC_KERNELS:
case GIMPLE_OACC_PARALLEL:
case GIMPLE_OMP_PARALLEL:
case GIMPLE_OMP_TASK:
@@ -11304,6 +11371,7 @@ make_gimple_omp_edges (basic_block bb, struct omp_region **region,
switch (code)
{
+ case GIMPLE_OACC_KERNELS:
case GIMPLE_OACC_PARALLEL:
case GIMPLE_OMP_PARALLEL:
case GIMPLE_OMP_TASK:
@@ -1316,6 +1316,7 @@ remap_gimple_stmt (gimple stmt, copy_body_data *id)
copy = gimple_build_wce (s1);
break;
+ case GIMPLE_OACC_KERNELS:
case GIMPLE_OACC_PARALLEL:
gcc_unreachable ();
@@ -3940,6 +3941,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_KERNELS:
case GIMPLE_OACC_PARALLEL:
case GIMPLE_OMP_PARALLEL:
case GIMPLE_OMP_TASK:
@@ -1248,6 +1248,7 @@ convert_nonlocal_reference_stmt (gimple_stmt_iterator *gsi, bool *handled_ops_p,
}
break;
+ case GIMPLE_OACC_KERNELS:
case GIMPLE_OACC_PARALLEL:
gcc_unreachable ();
@@ -1712,6 +1713,7 @@ convert_local_reference_stmt (gimple_stmt_iterator *gsi, bool *handled_ops_p,
switch (gimple_code (stmt))
{
+ case GIMPLE_OACC_KERNELS:
case GIMPLE_OACC_PARALLEL:
gcc_unreachable ();
@@ -2075,6 +2077,7 @@ convert_tramp_reference_stmt (gimple_stmt_iterator *gsi, bool *handled_ops_p,
break;
}
+ case GIMPLE_OACC_KERNELS:
case GIMPLE_OACC_PARALLEL:
gcc_unreachable ();
@@ -2138,6 +2141,7 @@ convert_gimple_call (gimple_stmt_iterator *gsi, bool *handled_ops_p,
}
break;
+ case GIMPLE_OACC_KERNELS:
case GIMPLE_OACC_PARALLEL:
gcc_unreachable ();
@@ -1,3 +1,9 @@
+2014-02-28 Thomas Schwinge <thomas@codesourcery.com>
+
+ * libgomp.map (GOACC_2.0): Add GOACC_kernels.
+ * libgomp_g.h (GOACC_kernels): New prototype.
+ * oacc-parallel.c (GOACC_kernels): New function.
+
2014-02-21 Thomas Schwinge <thomas@codesourcery.com>
* testsuite/libgomp.oacc-c/data-1.c: New file.
@@ -235,5 +235,6 @@ GOACC_2.0 {
global:
GOACC_data_end;
GOACC_data_start;
+ GOACC_kernels;
GOACC_parallel;
};
@@ -216,10 +216,12 @@ 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 short *);
extern void GOACC_data_start (int, const void *,
size_t, void **, size_t *, unsigned short *);
extern void GOACC_data_end (void);
+extern void GOACC_kernels (int, void (*) (void *), const void *,
+ size_t, void **, size_t *, unsigned short *);
+extern void GOACC_parallel (int, void (*) (void *), const void *,
+ size_t, void **, size_t *, unsigned short *);
#endif /* LIBGOMP_G_H */
@@ -23,7 +23,7 @@
see the files COPYING3 and COPYING.RUNTIME respectively. If not, see
<http://www.gnu.org/licenses/>. */
-/* This file handles the OpenACC data and parallel constructs. */
+/* This file handles OpenACC constructs. */
#include "libgomp.h"
#include "libgomp_g.h"
@@ -81,3 +81,13 @@ GOACC_data_end (void)
{
GOMP_target_end_data ();
}
+
+
+void
+GOACC_kernels (int device, void (*fn) (void *), const void *openmp_target,
+ size_t mapnum, void **hostaddrs, size_t *sizes,
+ unsigned short *kinds)
+{
+ /* TODO. */
+ GOACC_parallel (device, fn, openmp_target, mapnum, hostaddrs, sizes, kinds);
+}
From: tschwinge <tschwinge@138bc75d-0d04-0410-961f-82ee72b054a4> gcc/ * gimple.def (GIMPLE_OACC_KERNELS): New code. * doc/gimple.texi: Document it. * gimple.h (gimple_has_substatements, CASE_GIMPLE_OMP) (is_gimple_omp_oacc_specifically): Handle it. (gimple_statement_oacc_kernels): New struct. (gimple_build_oacc_kernels): New prototype. (gimple_oacc_kernels_clauses, gimple_oacc_kernels_clauses_ptr) (gimple_oacc_kernels_set_clauses, gimple_oacc_kernels_child_fn) (gimple_oacc_kernels_child_fn_ptr) (gimple_oacc_kernels_set_child_fn, gimple_oacc_kernels_data_arg) (gimple_oacc_kernels_data_arg_ptr) (gimple_oacc_kernels_set_data_arg): New inline functions. * gimple.c (gimple_build_oacc_kernels): New function. (gimple_copy): Handle GIMPLE_OACC_KERNELS. * gimple-low.c (lower_stmt): Likewise. * gimple-walk.c (walk_gimple_op, walk_gimple_stmt): Likewise. * gimple-pretty-print.c (pp_gimple_stmt_1): Likewise. (dump_gimple_oacc_parallel): Rename to dump_gimple_oacc_offload. Also handle GIMPLE_OACC_KERNELS. Update all callers. * gimplify.c (gimplify_omp_workshare, gimplify_expr): Handle OACC_KERNELS. * oacc-builtins.def (BUILT_IN_GOACC_KERNELS): New builtin. * omp-low.c (scan_oacc_parallel, expand_oacc_parallel) (lower_oacc_parallel): Rename to scan_oacc_offload, expand_oacc_offload, and lower_oacc_offload. Also handle GIMPLE_OACC_KERNELS. Update all callers. (scan_sharing_clauses, scan_omp_1_stmt, expand_omp, lower_omp_1) (diagnose_sb_0, diagnose_sb_1, diagnose_sb_2) (make_gimple_omp_edges): Handle GIMPLE_OACC_KERNELS. * tree-inline.c (remap_gimple_stmt, estimate_num_insns): Likewise. * tree-nested.c (convert_nonlocal_reference_stmt) (convert_local_reference_stmt, convert_tramp_reference_stmt) (convert_gimple_call): Likewise. libgomp/ * libgomp.map (GOACC_2.0): Add GOACC_kernels. * libgomp_g.h (GOACC_kernels): New prototype. * oacc-parallel.c (GOACC_kernels): New function. git-svn-id: svn+ssh://gcc.gnu.org/svn/gcc/branches/gomp-4_0-branch@208215 138bc75d-0d04-0410-961f-82ee72b054a4 --- gcc/ChangeLog.gomp | 36 +++++++++++++ gcc/doc/gimple.texi | 7 +++ gcc/gimple-low.c | 1 + gcc/gimple-pretty-print.c | 48 ++++++++++++----- gcc/gimple-walk.c | 16 ++++++ gcc/gimple.c | 18 +++++++ gcc/gimple.def | 22 +++++++- gcc/gimple.h | 130 ++++++++++++++++++++++++++++++++++++++++++++-- gcc/gimplify.c | 6 ++- gcc/oacc-builtins.def | 6 ++- gcc/omp-low.c | 116 ++++++++++++++++++++++++++++++++--------- gcc/tree-inline.c | 2 + gcc/tree-nested.c | 4 ++ libgomp/ChangeLog.gomp | 6 +++ libgomp/libgomp.map | 1 + libgomp/libgomp_g.h | 6 ++- libgomp/oacc-parallel.c | 12 ++++- 17 files changed, 389 insertions(+), 48 deletions(-)