Hi! On Wed, 12 Nov 2014 14:45:02 +0100, Jakub Jelinek <ja...@redhat.com> wrote: > On Wed, Nov 12, 2014 at 02:33:43PM +0100, Thomas Schwinge wrote: > > Months later, with months' worth of GCC internals experience, I now came > > to realize that maybe this has not actually been a useful thing to do > > (and likewise for the GIMPLE_OACC_KERNELS also added later on, > > <http://news.gmane.org/find-root.php?message_id=%3C1393579386-11666-1-git-send-email-thomas%40codesourcery.com%3E>). > > All handling of GIMPLE_OACC_PARALLEL and GIMPLE_OACC_KERNELS closely > > follows that of GIMPLE_OMP_TARGET's GF_OMP_TARGET_KIND_REGION, with only > > minor divergence. What I did not understand back then, has not been > > obvious to me, was that the underlying structure of all those codes will > > in fact be the same (as already made apparent by using the one > > GIMPLE_OMP_TARGET for all of: OpenMP target offloading regions, OpenMP > > target data regions, OpenMP target data maintenenace "executable" > > statements), and any "customization" then happens via the clauses > > attached to GIMPLE_OMP_TARGET. > > I'm fine with merging them into kinds, [...] > > > So, sanity check: should we now merge GIMPLE_OACC_PARALLEL and > > GIMPLE_OACC_KERNELS into being "subtypes" of GIMPLE_OMP_TARGET (like > > GF_OMP_TARGET_KIND_REGION), as already done for > > GF_OMP_TARGET_KIND_OACC_DATA (like GF_OMP_TARGET_KIND_DATA), and > > GF_OMP_TARGET_KIND_OACC_UPDATE and > > GF_OMP_TARGET_KIND_OACC_ENTER_EXIT_DATA (like GF_OMP_TARGET_KIND_UPDATE). > > Yep.
In r218568, I applied the following to gomp-4_0-branch:
commit 28629d718a63a782170cfb06a4d0278de0779039
Author: tschwinge <tschwinge@138bc75d-0d04-0410-961f-82ee72b054a4>
Date: Wed Dec 10 09:52:28 2014 +0000
Merge GIMPLE_OACC_KERNELS and GIMPLE_OACC_PARALLEL into GIMPLE_OMP_TARGET.
gcc/
* gimple.def (GIMPLE_OACC_KERNELS, GIMPLE_OACC_PARALLEL): Merge
into GIMPLE_OMP_TARGET. Update all users.
gcc/
* cgraphbuild.c (pass_build_cgraph_edges::execute): Remove
handling of GIMPLE_OACC_PARALLEL.
* gimple-pretty-print.c (dump_gimple_omp_target): Dump a bit more
data, pretty-printing.
git-svn-id: svn+ssh://gcc.gnu.org/svn/gcc/branches/gomp-4_0-branch@218568
138bc75d-0d04-0410-961f-82ee72b054a4
---
gcc/ChangeLog.gomp | 11 ++
gcc/cgraphbuild.c | 13 +-
gcc/doc/gimple.texi | 15 --
gcc/gimple-low.c | 2 -
gcc/gimple-pretty-print.c | 118 +++--------
gcc/gimple-walk.c | 32 ---
gcc/gimple.c | 40 +---
gcc/gimple.def | 35 +---
gcc/gimple.h | 273 ++------------------------
gcc/gimplify.c | 6 +-
gcc/omp-low.c | 225 +++++++--------------
gcc/testsuite/gfortran.dg/goacc/private-1.f95 | 2 +-
gcc/tree-inline.c | 6 -
gcc/tree-nested.c | 16 --
14 files changed, 136 insertions(+), 658 deletions(-)
diff --git gcc/ChangeLog.gomp gcc/ChangeLog.gomp
index bece7c1..06e8583 100644
--- gcc/ChangeLog.gomp
+++ gcc/ChangeLog.gomp
@@ -1,4 +1,15 @@
2014-12-10 Thomas Schwinge <tho...@codesourcery.com>
+ Bernd Schmidt <ber...@codesourcery.com>
+
+ * gimple.def (GIMPLE_OACC_KERNELS, GIMPLE_OACC_PARALLEL): Merge
+ into GIMPLE_OMP_TARGET. Update all users.
+
+2014-12-10 Thomas Schwinge <tho...@codesourcery.com>
+
+ * cgraphbuild.c (pass_build_cgraph_edges::execute): Remove
+ handling of GIMPLE_OACC_PARALLEL.
+ * gimple-pretty-print.c (dump_gimple_omp_target): Dump a bit more
+ data, pretty-printing.
* omp-low.c (build_omp_regions_1, make_gimple_omp_edges)
<GIMPLE_OMP_TARGET>: Handle
diff --git gcc/cgraphbuild.c gcc/cgraphbuild.c
index 9b078bc..c72ceab 100644
--- gcc/cgraphbuild.c
+++ gcc/cgraphbuild.c
@@ -368,21 +368,14 @@ pass_build_cgraph_edges::execute (function *fun)
bb->count, freq);
}
node->record_stmt_references (stmt);
- if (gimple_code (stmt) == GIMPLE_OACC_PARALLEL
- && gimple_oacc_parallel_child_fn (stmt))
- {
- tree fn = gimple_oacc_parallel_child_fn (stmt);
- node->create_reference (cgraph_node::get_create (fn),
- IPA_REF_ADDR, stmt);
- }
- else if (gimple_code (stmt) == GIMPLE_OMP_PARALLEL
- && gimple_omp_parallel_child_fn (stmt))
+ if (gimple_code (stmt) == GIMPLE_OMP_PARALLEL
+ && gimple_omp_parallel_child_fn (stmt))
{
tree fn = gimple_omp_parallel_child_fn (stmt);
node->create_reference (cgraph_node::get_create (fn),
IPA_REF_ADDR, stmt);
}
- else if (gimple_code (stmt) == GIMPLE_OMP_TASK)
+ if (gimple_code (stmt) == GIMPLE_OMP_TASK)
{
tree fn = gimple_omp_task_child_fn (stmt);
if (fn)
diff --git gcc/doc/gimple.texi gcc/doc/gimple.texi
index 4c59748..696c10e 100644
--- gcc/doc/gimple.texi
+++ gcc/doc/gimple.texi
@@ -439,8 +439,6 @@ 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
@item @code{GIMPLE_OMP_CONTINUE} @tab x @tab x
@@ -1008,8 +1006,6 @@ 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}::
* @code{GIMPLE_OMP_CONTINUE}::
@@ -1655,17 +1651,6 @@ Build a @code{GIMPLE_NOP} statement.
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}
-
-
@node @code{GIMPLE_OMP_ATOMIC_LOAD}
@subsection @code{GIMPLE_OMP_ATOMIC_LOAD}
@cindex @code{GIMPLE_OMP_ATOMIC_LOAD}
diff --git gcc/gimple-low.c gcc/gimple-low.c
index 60a7792e..3507d3c 100644
--- gcc/gimple-low.c
+++ gcc/gimple-low.c
@@ -368,8 +368,6 @@ 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:
case GIMPLE_OMP_TARGET:
diff --git gcc/gimple-pretty-print.c gcc/gimple-pretty-print.c
index 72dfac6..38d39f7 100644
--- gcc/gimple-pretty-print.c
+++ gcc/gimple-pretty-print.c
@@ -1338,15 +1338,21 @@ dump_gimple_omp_target (pretty_printer *buffer, gimple
gs, int spc, int flags)
case GF_OMP_TARGET_KIND_UPDATE:
kind = " update";
break;
+ case GF_OMP_TARGET_KIND_OACC_KERNELS:
+ kind = " oacc_kernels";
+ break;
+ case GF_OMP_TARGET_KIND_OACC_PARALLEL:
+ kind = " oacc_parallel";
+ break;
case GF_OMP_TARGET_KIND_OACC_DATA:
kind = " oacc_data";
break;
- case GF_OMP_TARGET_KIND_OACC_ENTER_EXIT_DATA:
- kind = " oacc_enter_exit_data";
- break;
case GF_OMP_TARGET_KIND_OACC_UPDATE:
kind = " oacc_update";
break;
+ case GF_OMP_TARGET_KIND_OACC_ENTER_EXIT_DATA:
+ kind = " oacc_enter_exit_data";
+ break;
default:
gcc_unreachable ();
}
@@ -1355,7 +1361,9 @@ dump_gimple_omp_target (pretty_printer *buffer, gimple
gs, int spc, int flags)
dump_gimple_fmt (buffer, spc, flags, "%G%s <%+BODY <%S>%nCLAUSES <", gs,
kind, gimple_omp_body (gs));
dump_omp_clauses (buffer, gimple_omp_target_clauses (gs), spc, flags);
- dump_gimple_fmt (buffer, spc, flags, " >");
+ dump_gimple_fmt (buffer, spc, flags, " >, %T, %T%n>",
+ gimple_omp_target_child_fn (gs),
+ gimple_omp_target_data_arg (gs));
}
else
{
@@ -1367,16 +1375,28 @@ dump_gimple_omp_target (pretty_printer *buffer, gimple
gs, int spc, int flags)
pp_string (buffer, " [child fn: ");
dump_generic_node (buffer, gimple_omp_target_child_fn (gs),
spc, flags, false);
- pp_right_bracket (buffer);
+ pp_string (buffer, " (");
+ if (gimple_omp_target_data_arg (gs))
+ dump_generic_node (buffer, gimple_omp_target_data_arg (gs),
+ spc, flags, false);
+ else
+ pp_string (buffer, "???");
+ pp_string (buffer, ")]");
}
- if (!gimple_seq_empty_p (gimple_omp_body (gs)))
+ gimple_seq body = gimple_omp_body (gs);
+ if (body && gimple_code (gimple_seq_first_stmt (body)) != GIMPLE_BIND)
{
newline_and_indent (buffer, spc + 2);
- pp_character (buffer, '{');
+ pp_left_brace (buffer);
pp_newline (buffer);
- dump_gimple_seq (buffer, gimple_omp_body (gs), spc + 4, flags);
+ dump_gimple_seq (buffer, body, spc + 4, flags);
newline_and_indent (buffer, spc + 2);
- pp_character (buffer, '}');
+ pp_right_brace (buffer);
+ }
+ else if (body)
+ {
+ pp_newline (buffer);
+ dump_gimple_seq (buffer, body, spc + 2, flags);
}
}
}
@@ -1878,81 +1898,6 @@ dump_gimple_phi (pretty_printer *buffer, gimple phi, int
spc, bool comment,
}
-/* 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_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_omp_clauses (gs), spc, flags);
- dump_gimple_fmt (buffer, spc, flags, " >, %T, %T%n>",
- gimple_omp_child_fn (gs), gimple_omp_data_arg (gs));
- }
- else
- {
- gimple_seq body;
- 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_omp_child_fn (gs),
- spc, flags, false);
- pp_string (buffer, " (");
- if (gimple_omp_data_arg (gs))
- dump_generic_node (buffer, gimple_omp_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). */
@@ -2237,11 +2182,6 @@ 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_offload (buffer, gs, spc, flags);
- break;
-
case GIMPLE_OMP_PARALLEL:
dump_gimple_omp_parallel (buffer, gs, spc, flags);
break;
diff --git gcc/gimple-walk.c gcc/gimple-walk.c
index cc74d34..bfa3532 100644
--- gcc/gimple-walk.c
+++ gcc/gimple-walk.c
@@ -304,36 +304,6 @@ 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);
- 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);
@@ -629,8 +599,6 @@ 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:
case GIMPLE_OMP_TASKGROUP:
diff --git gcc/gimple.c gcc/gimple.c
index e6de836..32615e8 100644
--- gcc/gimple.c
+++ gcc/gimple.c
@@ -811,40 +811,6 @@ 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.
- 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.
@@ -1077,7 +1043,7 @@ gimple_build_omp_single (gimple_seq body, tree clauses)
/* Build a GIMPLE_OMP_TARGET statement.
BODY is the sequence of statements that will be executed.
- KIND is the kind of target region.
+ KIND is the kind of the region.
CLAUSES are any of the construct's clauses. */
gimple
@@ -1719,10 +1685,6 @@ gimple_copy (gimple stmt)
gimple_try_set_cleanup (copy, new_seq);
break;
- case GIMPLE_OACC_KERNELS:
- case GIMPLE_OACC_PARALLEL:
- gcc_unreachable ();
-
case GIMPLE_OMP_FOR:
gcc_assert (!is_gimple_omp_oacc_specifically (stmt));
new_seq = gimple_seq_copy (gimple_omp_for_pre_body (stmt));
diff --git gcc/gimple.def gcc/gimple.def
index e2e912c..269c2d7 100644
--- gcc/gimple.def
+++ gcc/gimple.def
@@ -205,33 +205,8 @@ DEFGSCODE(GIMPLE_NOP, "gimple_nop", GSS_BASE)
/* IMPORTANT.
- 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_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 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)
+ Do not rearrange any of the GIMPLE_OMP_* codes. This ordering is
+ exposed by the range check in gimple_omp_subcode. */
/* 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
@@ -381,12 +356,12 @@ DEFGSCODE(GIMPLE_OMP_SECTIONS_SWITCH,
"gimple_omp_sections_switch", GSS_BASE)
DEFGSCODE(GIMPLE_OMP_SINGLE, "gimple_omp_single", GSS_OMP_SINGLE_LAYOUT)
/* GIMPLE_OMP_TARGET <BODY, CLAUSES, CHILD_FN> represents
- #pragma acc data
+ #pragma acc {kernels,parallel,data}
#pragma omp target {,data,update}
- BODY is the sequence of statements inside the target construct
+ BODY is the sequence of statements inside the construct
(NULL for target update).
CLAUSES is an OMP_CLAUSE chain holding the associated clauses.
- CHILD_FN is set when outlining the body of the target region.
+ CHILD_FN is set when outlining the body of the offloaded 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
diff --git gcc/gimple.h gcc/gimple.h
index a91bd4e..60c2469 100644
--- gcc/gimple.h
+++ gcc/gimple.h
@@ -108,9 +108,11 @@ enum gf_mask {
GF_OMP_TARGET_KIND_REGION = 0,
GF_OMP_TARGET_KIND_DATA = 1,
GF_OMP_TARGET_KIND_UPDATE = 2,
- GF_OMP_TARGET_KIND_OACC_DATA = 3,
- GF_OMP_TARGET_KIND_OACC_UPDATE = 4,
- GF_OMP_TARGET_KIND_OACC_ENTER_EXIT_DATA = 5,
+ GF_OMP_TARGET_KIND_OACC_PARALLEL = 3,
+ GF_OMP_TARGET_KIND_OACC_KERNELS = 4,
+ GF_OMP_TARGET_KIND_OACC_DATA = 5,
+ GF_OMP_TARGET_KIND_OACC_UPDATE = 6,
+ GF_OMP_TARGET_KIND_OACC_ENTER_EXIT_DATA = 7,
/* True on an GIMPLE_OMP_RETURN statement if the return does not require
a thread synchronization via some sort of barrier. The exact barrier
@@ -560,8 +562,8 @@ struct GTY((tag("GSS_OMP_FOR")))
};
-/* GIMPLE_OACC_KERNELS, GIMPLE_OACC_PARALLEL, GIMPLE_OMP_PARALLEL,
- GIMPLE_OMP_TARGET, GIMPLE_OMP_TASK */
+/* 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
{
@@ -580,22 +582,6 @@ 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
-{
- /* No extra fields; adds invariant:
- stmt->code == GIMPLE_OACC_PARALLEL. */
-};
-
/* GIMPLE_OMP_PARALLEL or GIMPLE_TASK */
struct GTY((tag("GSS_OMP_PARALLEL_LAYOUT")))
gimple_statement_omp_taskreg : public gimple_statement_omp_parallel_layout
@@ -913,22 +899,6 @@ 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;
-}
-
-template <>
-template <>
-inline bool
is_a_helper <gimple_statement_omp_taskreg *>::test (gimple gs)
{
return gs->code == GIMPLE_OMP_PARALLEL || gs->code == GIMPLE_OMP_TASK;
@@ -1121,22 +1091,6 @@ 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;
-}
-
-template <>
-template <>
-inline bool
is_a_helper <const gimple_statement_omp_taskreg *>::test (const_gimple gs)
{
return gs->code == GIMPLE_OMP_PARALLEL || gs->code == GIMPLE_OMP_TASK;
@@ -1260,8 +1214,6 @@ 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);
gimple gimple_build_omp_parallel (gimple_seq, tree, tree, tree);
@@ -1500,8 +1452,6 @@ 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:
case GIMPLE_OMP_TASKGROUP:
@@ -4362,197 +4312,6 @@ 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
-gimple_oacc_parallel_clauses (const_gimple gs)
-{
- const gimple_statement_oacc_parallel *oacc_parallel_stmt =
- as_a <const gimple_statement_oacc_parallel *> (gs);
- return oacc_parallel_stmt->clauses;
-}
-
-/* Return a pointer to the clauses associated with OACC_PARALLEL statement
- GS. */
-
-static inline tree *
-gimple_oacc_parallel_clauses_ptr (gimple gs)
-{
- gimple_statement_oacc_parallel *oacc_parallel_stmt =
- as_a <gimple_statement_oacc_parallel *> (gs);
- return &oacc_parallel_stmt->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_statement_oacc_parallel *oacc_parallel_stmt =
- as_a <gimple_statement_oacc_parallel *> (gs);
- oacc_parallel_stmt->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)
-{
- const gimple_statement_oacc_parallel *oacc_parallel_stmt =
- as_a <const gimple_statement_oacc_parallel *> (gs);
- return oacc_parallel_stmt->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_statement_oacc_parallel *oacc_parallel_stmt =
- as_a <gimple_statement_oacc_parallel *> (gs);
- return &oacc_parallel_stmt->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_statement_oacc_parallel *oacc_parallel_stmt =
- as_a <gimple_statement_oacc_parallel *> (gs);
- oacc_parallel_stmt->child_fn = child_fn;
-}
-
-/* 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)
-{
- const gimple_statement_oacc_parallel *oacc_parallel_stmt =
- as_a <const gimple_statement_oacc_parallel *> (gs);
- return oacc_parallel_stmt->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_statement_oacc_parallel *oacc_parallel_stmt =
- as_a <gimple_statement_oacc_parallel *> (gs);
- return &oacc_parallel_stmt->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_statement_oacc_parallel *oacc_parallel_stmt =
- as_a <gimple_statement_oacc_parallel *> (gs);
- oacc_parallel_stmt->data_arg = data_arg;
-}
-
-
/* Return the name associated with OMP_CRITICAL statement GS. */
static inline tree
@@ -5374,7 +5133,7 @@ gimple_omp_target_set_clauses (gimple gs, tree clauses)
}
-/* Return the kind of OMP target statemement. */
+/* Return the kind of the OMP_TARGET G. */
static inline int
gimple_omp_target_kind (const_gimple g)
@@ -5384,7 +5143,7 @@ gimple_omp_target_kind (const_gimple g)
}
-/* Set the OMP target kind. */
+/* Set the kind of the OMP_TARGET G. */
static inline void
gimple_omp_target_set_kind (gimple g, int kind)
@@ -5854,8 +5613,6 @@ 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_KERNELS: \
- case GIMPLE_OACC_PARALLEL: \
case GIMPLE_OMP_PARALLEL: \
case GIMPLE_OMP_TASK: \
case GIMPLE_OMP_FOR: \
@@ -5898,9 +5655,6 @@ 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_FOR:
switch (gimple_omp_for_kind (stmt))
{
@@ -5908,10 +5662,12 @@ is_gimple_omp_oacc_specifically (const_gimple stmt)
return true;
default:
return false;
- }
+ }
case GIMPLE_OMP_TARGET:
switch (gimple_omp_target_kind (stmt))
{
+ case GF_OMP_TARGET_KIND_OACC_PARALLEL:
+ case GF_OMP_TARGET_KIND_OACC_KERNELS:
case GF_OMP_TARGET_KIND_OACC_DATA:
case GF_OMP_TARGET_KIND_OACC_UPDATE:
case GF_OMP_TARGET_KIND_OACC_ENTER_EXIT_DATA:
@@ -5933,13 +5689,12 @@ is_gimple_omp_offloaded (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:
switch (gimple_omp_target_kind (stmt))
{
case GF_OMP_TARGET_KIND_REGION:
+ case GF_OMP_TARGET_KIND_OACC_PARALLEL:
+ case GF_OMP_TARGET_KIND_OACC_KERNELS:
return true;
default:
return false;
diff --git gcc/gimplify.c gcc/gimplify.c
index ad48d51..eb9d930 100644
--- gcc/gimplify.c
+++ gcc/gimplify.c
@@ -7315,10 +7315,12 @@ gimplify_omp_workshare (tree *expr_p, gimple_seq *pre_p)
OACC_DATA_CLAUSES (expr));
break;
case OACC_KERNELS:
- stmt = gimple_build_oacc_kernels (body, OACC_KERNELS_CLAUSES (expr));
+ stmt = gimple_build_omp_target (body, GF_OMP_TARGET_KIND_OACC_KERNELS,
+ OMP_CLAUSES (expr));
break;
case OACC_PARALLEL:
- stmt = gimple_build_oacc_parallel (body, OACC_PARALLEL_CLAUSES (expr));
+ stmt = gimple_build_omp_target (body, GF_OMP_TARGET_KIND_OACC_PARALLEL,
+ OMP_CLAUSES (expr));
break;
case OMP_SECTIONS:
stmt = gimple_build_omp_sections (body, OMP_CLAUSES (expr));
diff --git gcc/omp-low.c gcc/omp-low.c
index 6fed38f..39e2f22 100644
--- gcc/omp-low.c
+++ gcc/omp-low.c
@@ -272,10 +272,12 @@ oacc_max_threads (omp_context *ctx)
Scan for the innermost vector_length clause. */
for (omp_context *oc = ctx; oc; oc = oc->outer)
{
- if (gimple_code (oc->stmt) != GIMPLE_OACC_PARALLEL)
+ if (gimple_code (oc->stmt) != GIMPLE_OMP_TARGET
+ || (gimple_omp_target_kind (oc->stmt)
+ != GF_OMP_TARGET_KIND_OACC_PARALLEL))
continue;
- clauses = gimple_oacc_parallel_clauses (oc->stmt);
+ clauses = gimple_omp_target_clauses (oc->stmt);
vector_length = find_omp_clause (clauses, OMP_CLAUSE_VECTOR_LENGTH);
if (vector_length)
@@ -2643,28 +2645,7 @@ scan_omp_target (gimple stmt, omp_context *outer_ctx)
{
omp_context *ctx;
tree name;
- bool offloaded;
- void (*gimple_omp_set_child_fn) (gimple, tree);
- tree (*gimple_omp_clauses) (const_gimple);
-
- offloaded = is_gimple_omp_offloaded (stmt);
- 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;
- case GIMPLE_OMP_TARGET:
- gimple_omp_set_child_fn = gimple_omp_target_set_child_fn;
- gimple_omp_clauses = gimple_omp_target_clauses;
- break;
- default:
- gcc_unreachable ();
- }
+ bool offloaded = is_gimple_omp_offloaded (stmt);
if (is_gimple_omp_oacc_specifically (stmt))
{
@@ -2689,10 +2670,10 @@ scan_omp_target (gimple stmt, omp_context *outer_ctx)
0, 0);
create_omp_child_function (ctx, false);
- gimple_omp_set_child_fn (stmt, ctx->cb.dst_fn);
+ gimple_omp_target_set_child_fn (stmt, ctx->cb.dst_fn);
}
- scan_sharing_clauses (gimple_omp_clauses (stmt), ctx);
+ scan_sharing_clauses (gimple_omp_target_clauses (stmt), ctx);
scan_omp (gimple_omp_body_ptr (stmt), ctx);
if (TYPE_FIELDS (ctx->record_type) == NULL)
@@ -2944,8 +2925,6 @@ check_omp_nesting_restrictions (gimple stmt, omp_context
*ctx)
"of work-sharing, critical, ordered, master or explicit "
"task region");
return false;
- case GIMPLE_OACC_KERNELS:
- case GIMPLE_OACC_PARALLEL:
case GIMPLE_OMP_PARALLEL:
return true;
default:
@@ -3203,8 +3182,6 @@ scan_omp_1_stmt (gimple_stmt_iterator *gsi, bool
*handled_ops_p,
scan_omp (gimple_omp_body_ptr (stmt), ctx);
break;
- case GIMPLE_OACC_KERNELS:
- case GIMPLE_OACC_PARALLEL:
case GIMPLE_OMP_TARGET:
scan_omp_target (stmt, ctx);
break;
@@ -8780,42 +8757,24 @@ expand_omp_target (struct omp_region *region)
gimple entry_stmt, stmt;
edge e;
bool offloaded, data_region;
- tree (*gimple_omp_child_fn) (const_gimple);
- tree (*gimple_omp_data_arg) (const_gimple);
entry_stmt = last_stmt (region->entry);
new_bb = region->entry;
offloaded = is_gimple_omp_offloaded (entry_stmt);
- data_region = false;
- switch (region->type)
+ switch (gimple_omp_target_kind (entry_stmt))
{
- case GIMPLE_OACC_KERNELS:
- gimple_omp_child_fn = gimple_oacc_kernels_child_fn;
- gimple_omp_data_arg = gimple_oacc_kernels_data_arg;
+ case GF_OMP_TARGET_KIND_REGION:
+ case GF_OMP_TARGET_KIND_UPDATE:
+ case GF_OMP_TARGET_KIND_OACC_PARALLEL:
+ case GF_OMP_TARGET_KIND_OACC_KERNELS:
+ case GF_OMP_TARGET_KIND_OACC_UPDATE:
+ case GF_OMP_TARGET_KIND_OACC_ENTER_EXIT_DATA:
+ data_region = false;
break;
- case GIMPLE_OACC_PARALLEL:
- gimple_omp_child_fn = gimple_oacc_parallel_child_fn;
- gimple_omp_data_arg = gimple_oacc_parallel_data_arg;
- break;
- case GIMPLE_OMP_TARGET:
- switch (gimple_omp_target_kind (entry_stmt))
- {
- case GF_OMP_TARGET_KIND_DATA:
- case GF_OMP_TARGET_KIND_OACC_DATA:
- data_region = true;
- break;
- case GF_OMP_TARGET_KIND_REGION:
- case GF_OMP_TARGET_KIND_UPDATE:
- case GF_OMP_TARGET_KIND_OACC_UPDATE:
- case GF_OMP_TARGET_KIND_OACC_ENTER_EXIT_DATA:
- break;
- default:
- gcc_unreachable ();
- }
-
- gimple_omp_child_fn = gimple_omp_target_child_fn;
- gimple_omp_data_arg = gimple_omp_target_data_arg;
+ case GF_OMP_TARGET_KIND_DATA:
+ case GF_OMP_TARGET_KIND_OACC_DATA:
+ data_region = true;
break;
default:
gcc_unreachable ();
@@ -8825,7 +8784,7 @@ expand_omp_target (struct omp_region *region)
child_cfun = NULL;
if (offloaded)
{
- child_fn = gimple_omp_child_fn (entry_stmt);
+ child_fn = gimple_omp_target_child_fn (entry_stmt);
child_cfun = DECL_STRUCT_FUNCTION (child_fn);
}
@@ -8854,13 +8813,14 @@ expand_omp_target (struct omp_region *region)
a function call that has been inlined, the original PARM_DECL
.OMP_DATA_I may have been converted into a different local
variable. In which case, we need to keep the assignment. */
- if (gimple_omp_data_arg (entry_stmt))
+ tree data_arg = gimple_omp_target_data_arg (entry_stmt);
+ if (data_arg)
{
basic_block entry_succ_bb = single_succ (entry_bb);
gimple_stmt_iterator gsi;
tree arg;
gimple tgtcopy_stmt = NULL;
- tree sender = TREE_VEC_ELT (gimple_omp_data_arg (entry_stmt), 0);
+ tree sender = TREE_VEC_ELT (data_arg, 0);
for (gsi = gsi_start_bb (entry_succ_bb); ; gsi_next (&gsi))
{
@@ -8994,49 +8954,38 @@ expand_omp_target (struct omp_region *region)
tree t1, t2, t3, t4, device, cond, c, clauses;
enum built_in_function start_ix;
location_t clause_loc;
- tree (*gimple_omp_clauses) (const_gimple);
- switch (region->type)
+ switch (gimple_omp_target_kind (entry_stmt))
{
- case GIMPLE_OACC_KERNELS:
- gimple_omp_clauses = gimple_oacc_kernels_clauses;
- start_ix = BUILT_IN_GOACC_KERNELS;
+ case GF_OMP_TARGET_KIND_REGION:
+ start_ix = BUILT_IN_GOMP_TARGET;
break;
- case GIMPLE_OACC_PARALLEL:
- gimple_omp_clauses = gimple_oacc_parallel_clauses;
+ case GF_OMP_TARGET_KIND_DATA:
+ start_ix = BUILT_IN_GOMP_TARGET_DATA;
+ break;
+ case GF_OMP_TARGET_KIND_UPDATE:
+ start_ix = BUILT_IN_GOMP_TARGET_UPDATE;
+ break;
+ case GF_OMP_TARGET_KIND_OACC_PARALLEL:
start_ix = BUILT_IN_GOACC_PARALLEL;
break;
- case GIMPLE_OMP_TARGET:
- gimple_omp_clauses = gimple_omp_target_clauses;
- switch (gimple_omp_target_kind (entry_stmt))
- {
- case GF_OMP_TARGET_KIND_REGION:
- start_ix = BUILT_IN_GOMP_TARGET;
- break;
- case GF_OMP_TARGET_KIND_DATA:
- start_ix = BUILT_IN_GOMP_TARGET_DATA;
- break;
- case GF_OMP_TARGET_KIND_UPDATE:
- start_ix = BUILT_IN_GOMP_TARGET_UPDATE;
- break;
- case GF_OMP_TARGET_KIND_OACC_DATA:
- start_ix = BUILT_IN_GOACC_DATA_START;
- break;
- case GF_OMP_TARGET_KIND_OACC_ENTER_EXIT_DATA:
- start_ix = BUILT_IN_GOACC_ENTER_EXIT_DATA;
- break;
- case GF_OMP_TARGET_KIND_OACC_UPDATE:
- start_ix = BUILT_IN_GOACC_UPDATE;
- break;
- default:
- gcc_unreachable ();
- }
+ case GF_OMP_TARGET_KIND_OACC_KERNELS:
+ start_ix = BUILT_IN_GOACC_KERNELS;
+ break;
+ case GF_OMP_TARGET_KIND_OACC_DATA:
+ start_ix = BUILT_IN_GOACC_DATA_START;
+ break;
+ case GF_OMP_TARGET_KIND_OACC_UPDATE:
+ start_ix = BUILT_IN_GOACC_UPDATE;
+ break;
+ case GF_OMP_TARGET_KIND_OACC_ENTER_EXIT_DATA:
+ start_ix = BUILT_IN_GOACC_ENTER_EXIT_DATA;
break;
default:
gcc_unreachable ();
}
- clauses = gimple_omp_clauses (entry_stmt);
+ clauses = gimple_omp_target_clauses (entry_stmt);
/* By default, the value of DEVICE is -1 (let runtime library choose)
and there is no conditional. */
@@ -9119,7 +9068,7 @@ expand_omp_target (struct omp_region *region)
}
gsi = gsi_last_bb (new_bb);
- t = gimple_omp_data_arg (entry_stmt);
+ t = gimple_omp_target_data_arg (entry_stmt);
if (t == NULL)
{
t1 = size_zero_node;
@@ -9331,8 +9280,6 @@ expand_omp (struct omp_region *region)
expand_omp_atomic (region);
break;
- case GIMPLE_OACC_KERNELS:
- case GIMPLE_OACC_PARALLEL:
case GIMPLE_OMP_TARGET:
expand_omp_target (region);
break;
@@ -9679,22 +9626,18 @@ oacc_initialize_reduction_data (tree clauses, tree
nthreads,
tree c, t, oc;
gimple stmt;
omp_context *octx;
- tree (*gimple_omp_clauses) (const_gimple);
- void (*gimple_omp_set_clauses) (gimple, tree);
/* Find the innermost PARALLEL openmp context. FIXME: OpenACC kernels
may require extra care unless they are converted to openmp for loops. */
-
- if (gimple_code (ctx->stmt) == GIMPLE_OACC_PARALLEL)
+ if (gimple_code (ctx->stmt) == GIMPLE_OMP_TARGET
+ && (gimple_omp_target_kind (ctx->stmt)
+ == GF_OMP_TARGET_KIND_OACC_PARALLEL))
octx = ctx;
else
octx = ctx->outer;
- gimple_omp_clauses = gimple_oacc_parallel_clauses;
- gimple_omp_set_clauses = gimple_oacc_parallel_set_clauses;
-
/* Extract the clauses. */
- oc = gimple_omp_clauses (octx->stmt);
+ oc = gimple_omp_target_clauses (octx->stmt);
/* Find the last outer clause. */
for (; oc && OMP_CLAUSE_CHAIN (oc); oc = OMP_CLAUSE_CHAIN (oc))
@@ -9744,7 +9687,7 @@ oacc_initialize_reduction_data (tree clauses, tree
nthreads,
if (oc)
OMP_CLAUSE_CHAIN (oc) = t;
else
- gimple_omp_set_clauses (octx->stmt, t);
+ gimple_omp_target_set_clauses (octx->stmt, t);
OMP_CLAUSE_SIZE (t) = size;
oc = t;
}
@@ -11195,45 +11138,27 @@ lower_omp_target (gimple_stmt_iterator *gsi_p,
omp_context *ctx)
location_t loc = gimple_location (stmt);
bool offloaded, data_region;
unsigned int map_cnt = 0;
- tree (*gimple_omp_clauses) (const_gimple);
- void (*gimple_omp_set_data_arg) (gimple, tree);
offloaded = is_gimple_omp_offloaded (stmt);
- data_region = false;
- switch (gimple_code (stmt))
+ switch (gimple_omp_target_kind (stmt))
{
- case GIMPLE_OACC_KERNELS:
- gimple_omp_clauses = gimple_oacc_kernels_clauses;
- gimple_omp_set_data_arg = gimple_oacc_kernels_set_data_arg;
+ case GF_OMP_TARGET_KIND_REGION:
+ case GF_OMP_TARGET_KIND_UPDATE:
+ case GF_OMP_TARGET_KIND_OACC_PARALLEL:
+ case GF_OMP_TARGET_KIND_OACC_KERNELS:
+ case GF_OMP_TARGET_KIND_OACC_UPDATE:
+ case GF_OMP_TARGET_KIND_OACC_ENTER_EXIT_DATA:
+ data_region = false;
break;
- case GIMPLE_OACC_PARALLEL:
- gimple_omp_clauses = gimple_oacc_parallel_clauses;
- gimple_omp_set_data_arg = gimple_oacc_parallel_set_data_arg;
- break;
- case GIMPLE_OMP_TARGET:
- switch (gimple_omp_target_kind (stmt))
- {
- case GF_OMP_TARGET_KIND_DATA:
- case GF_OMP_TARGET_KIND_OACC_DATA:
- data_region = true;
- break;
- case GF_OMP_TARGET_KIND_REGION:
- case GF_OMP_TARGET_KIND_UPDATE:
- case GF_OMP_TARGET_KIND_OACC_UPDATE:
- case GF_OMP_TARGET_KIND_OACC_ENTER_EXIT_DATA:
- break;
- default:
- gcc_unreachable ();
- }
-
- gimple_omp_clauses = gimple_omp_target_clauses;
- gimple_omp_set_data_arg = gimple_omp_target_set_data_arg;
+ case GF_OMP_TARGET_KIND_DATA:
+ case GF_OMP_TARGET_KIND_OACC_DATA:
+ data_region = true;
break;
default:
gcc_unreachable ();
}
- clauses = gimple_omp_clauses (stmt);
+ clauses = gimple_omp_target_clauses (stmt);
tgt_bind = NULL;
tgt_body = NULL;
@@ -11250,15 +11175,9 @@ lower_omp_target (gimple_stmt_iterator *gsi_p,
omp_context *ctx)
irlist = NULL;
orlist = NULL;
- switch (gimple_code (stmt))
- {
- case GIMPLE_OACC_KERNELS:
- case GIMPLE_OACC_PARALLEL:
- oacc_process_reduction_data (&tgt_body, &irlist, &orlist, ctx);
- break;
- default:
- break;
- }
+ if (offloaded
+ && is_gimple_omp_oacc_specifically (stmt))
+ oacc_process_reduction_data (&tgt_body, &irlist, &orlist, ctx);
for (c = clauses; c ; c = OMP_CLAUSE_CHAIN (c))
switch (OMP_CLAUSE_CODE (c))
@@ -11390,7 +11309,7 @@ lower_omp_target (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_omp_set_data_arg (stmt, t);
+ gimple_omp_target_set_data_arg (stmt, t);
vec<constructor_elt, va_gc> *vsize;
vec<constructor_elt, va_gc> *vkind;
@@ -11457,8 +11376,8 @@ lower_omp_target (gimple_stmt_iterator *gsi_p,
omp_context *ctx)
|| TREE_CODE (TREE_TYPE (ovar)) != ARRAY_TYPE);
if (maybe_lookup_oacc_reduction (var, ctx))
{
- gcc_assert (gimple_code (stmt) == GIMPLE_OACC_KERNELS
- || gimple_code (stmt) == GIMPLE_OACC_PARALLEL);
+ gcc_assert (offloaded
+ && is_gimple_omp_oacc_specifically (stmt));
gimplify_assign (x, var, &ilist);
}
else if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
@@ -11802,8 +11721,6 @@ lower_omp_1 (gimple_stmt_iterator *gsi_p, omp_context
*ctx)
lower_omp_regimplify_p, ctx ? NULL : &wi, NULL))
gimple_regimplify_operands (stmt, gsi_p);
break;
- case GIMPLE_OACC_KERNELS:
- case GIMPLE_OACC_PARALLEL:
case GIMPLE_OMP_TARGET:
ctx = maybe_lookup_ctx (stmt);
gcc_assert (ctx);
@@ -12095,8 +12012,6 @@ 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:
case GIMPLE_OMP_SECTIONS:
@@ -12155,8 +12070,6 @@ 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:
case GIMPLE_OMP_SECTIONS:
@@ -12252,8 +12165,6 @@ 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:
case GIMPLE_OMP_FOR:
diff --git gcc/testsuite/gfortran.dg/goacc/private-1.f95
gcc/testsuite/gfortran.dg/goacc/private-1.f95
index 54c027d..23ce95a 100644
--- gcc/testsuite/gfortran.dg/goacc/private-1.f95
+++ gcc/testsuite/gfortran.dg/goacc/private-1.f95
@@ -31,7 +31,7 @@ program test
end do
!$acc end parallel
end program test
-! { dg-final { scan-tree-dump-times "pragma acc parallel" 3 "omplower" } }
+! { dg-final { scan-tree-dump-times "pragma omp target oacc_parallel" 3
"omplower" } }
! { dg-final { scan-tree-dump-times "private\\(i\\)" 3 "omplower" } }
! { dg-final { scan-tree-dump-times "private\\(j\\)" 2 "omplower" } }
! { dg-final { scan-tree-dump-times "private\\(k\\)" 1 "omplower" } }
diff --git gcc/tree-inline.c gcc/tree-inline.c
index 54b3514..89cb1eb 100644
--- gcc/tree-inline.c
+++ gcc/tree-inline.c
@@ -1395,10 +1395,6 @@ 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 ();
-
case GIMPLE_OMP_PARALLEL:
s1 = remap_gimple_seq (gimple_omp_body (stmt), id);
copy = gimple_build_omp_parallel
@@ -4112,8 +4108,6 @@ 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:
case GIMPLE_OMP_CRITICAL:
diff --git gcc/tree-nested.c gcc/tree-nested.c
index b5d6543..4da6297 100644
--- gcc/tree-nested.c
+++ gcc/tree-nested.c
@@ -1325,10 +1325,6 @@ convert_nonlocal_reference_stmt (gimple_stmt_iterator
*gsi, bool *handled_ops_p,
}
break;
- case GIMPLE_OACC_KERNELS:
- case GIMPLE_OACC_PARALLEL:
- gcc_unreachable ();
-
case GIMPLE_OMP_PARALLEL:
case GIMPLE_OMP_TASK:
save_suppress = info->suppress_expansion;
@@ -1898,10 +1894,6 @@ 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 ();
-
case GIMPLE_OMP_PARALLEL:
case GIMPLE_OMP_TASK:
save_suppress = info->suppress_expansion;
@@ -2289,10 +2281,6 @@ convert_tramp_reference_stmt (gimple_stmt_iterator *gsi,
bool *handled_ops_p,
break;
}
- case GIMPLE_OACC_KERNELS:
- case GIMPLE_OACC_PARALLEL:
- gcc_unreachable ();
-
case GIMPLE_OMP_TARGET:
gcc_assert (!is_gimple_omp_oacc_specifically (stmt));
if (gimple_omp_target_kind (stmt) != GF_OMP_TARGET_KIND_REGION)
@@ -2360,10 +2348,6 @@ convert_gimple_call (gimple_stmt_iterator *gsi, bool
*handled_ops_p,
}
break;
- case GIMPLE_OACC_KERNELS:
- case GIMPLE_OACC_PARALLEL:
- gcc_unreachable ();
-
case GIMPLE_OMP_PARALLEL:
case GIMPLE_OMP_TASK:
save_static_chain_added = info->static_chain_added;
Grüße,
Thomas
signature.asc
Description: PGP signature
