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(-) diff --git gcc/ChangeLog.gomp gcc/ChangeLog.gomp index 3d9b06d..79030d6 100644 --- gcc/ChangeLog.gomp +++ gcc/ChangeLog.gomp @@ -1,3 +1,39 @@ +2014-02-28 Thomas Schwinge <tho...@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 <tho...@codesourcery.com> * gimplify.c (gimplify_oacc_parallel): Merge into diff --git gcc/doc/gimple.texi gcc/doc/gimple.texi index 6136963..91748a6 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_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} diff --git gcc/gimple-low.c gcc/gimple-low.c index 7bf69bd..c7d9c1c 100644 --- gcc/gimple-low.c +++ gcc/gimple-low.c @@ -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: diff --git gcc/gimple-pretty-print.c gcc/gimple-pretty-print.c index ad9369c..1a31192 100644 --- gcc/gimple-pretty-print.c +++ gcc/gimple-pretty-print.c @@ -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: diff --git gcc/gimple-walk.c gcc/gimple-walk.c index a90ba55..b5b4095 100644 --- gcc/gimple-walk.c +++ gcc/gimple-walk.c @@ -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: diff --git gcc/gimple.c gcc/gimple.c index 30561b1..1862de2 100644 --- gcc/gimple.c +++ gcc/gimple.c @@ -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 (); diff --git gcc/gimple.def gcc/gimple.def index ce800bd..c9756b7 100644 --- gcc/gimple.def +++ gcc/gimple.def @@ -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 diff --git gcc/gimple.h gcc/gimple.h index b4ee9fa..514af32 100644 --- gcc/gimple.h +++ gcc/gimple.h @@ -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: diff --git gcc/gimplify.c gcc/gimplify.c index 6dbabfa..f3c34f9 100644 --- gcc/gimplify.c +++ gcc/gimplify.c @@ -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: diff --git gcc/oacc-builtins.def gcc/oacc-builtins.def index eaf3228..2d5c22c 100644 --- gcc/oacc-builtins.def +++ gcc/oacc-builtins.def @@ -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) diff --git gcc/omp-low.c gcc/omp-low.c index eec862e..2f13fb4 100644 --- gcc/omp-low.c +++ gcc/omp-low.c @@ -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: diff --git gcc/tree-inline.c gcc/tree-inline.c index 61c1cc8..8b22b86 100644 --- gcc/tree-inline.c +++ gcc/tree-inline.c @@ -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: diff --git gcc/tree-nested.c gcc/tree-nested.c index afa7abb..397f851 100644 --- gcc/tree-nested.c +++ gcc/tree-nested.c @@ -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 (); diff --git libgomp/ChangeLog.gomp libgomp/ChangeLog.gomp index b90b09b..3ea5901 100644 --- libgomp/ChangeLog.gomp +++ libgomp/ChangeLog.gomp @@ -1,3 +1,9 @@ +2014-02-28 Thomas Schwinge <tho...@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 <tho...@codesourcery.com> * testsuite/libgomp.oacc-c/data-1.c: New file. diff --git libgomp/libgomp.map libgomp/libgomp.map index cb52e45..e9f8b55 100644 --- libgomp/libgomp.map +++ libgomp/libgomp.map @@ -235,5 +235,6 @@ GOACC_2.0 { global: GOACC_data_end; GOACC_data_start; + GOACC_kernels; GOACC_parallel; }; diff --git libgomp/libgomp_g.h libgomp/libgomp_g.h index b9083a5..9681566 100644 --- libgomp/libgomp_g.h +++ libgomp/libgomp_g.h @@ -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 */ diff --git libgomp/oacc-parallel.c libgomp/oacc-parallel.c index 3ac7e39..cb883a8 100644 --- libgomp/oacc-parallel.c +++ libgomp/oacc-parallel.c @@ -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); +} -- 1.8.1.1