Hi!
This patch contains C++ parser changes etc. to handle
#pragma omp {teams,target {,data,update},distribute} parsing
all the way through till omp lowering (it bombs badly in omp expansion,
but already omp lowering will need to be tought out).
Things not handled yet are #pragma omp declare target/#pragma omp end
declare target (I assume we can handle it by automatically adding
"omp target" attribute to vars/functions in there) and there is no support
for array sections yet (also relevant for OMP_CLAUSE_DEPEND).
Say:
void baz (float *, float *, int);
float
foo (int x)
{
float b[1024], c[1024], s = 0;
int i;
baz (b, c, x);
#pragma omp target map(to: b, c)
#pragma omp parallel for reduction(+:s)
for (i = 0; i < x; i++)
s += b[i] * c[i];
return s;
}
float
bar (int x, int y, int z)
{
float b[1024], c[1024], s = 0;
int i, j;
baz (b, c, x);
#pragma omp target data map(to: b)
{
#pragma omp target map(to: c)
#pragma omp teams num_teams(y) num_threads(z) reduction(+:s)
#pragma omp distribute dist_schedule(static, 4) collapse(1)
for (j=0; j < x; j += y)
#pragma omp parallel for reduction(+:s)
for (i = j; i < j + y; i++)
s += b[i] * c[i];
#pragma omp target update from(b, c)
}
return s;
}
now parses with g++ -S -fopenmp -fdump-tree-all and shows dump up to
*.cfg (crash during ompexp).
Comments on this before I commit it to gomp-4_0-branch?
2013-05-24 Jakub Jelinek <ja...@redhat.com>
* tree.def (OMP_TEAMS, OMP_TARGET_DATA, OMP_TARGET,
OMP_TARGET_UPDATE): New tree codes.
* tree-cfg.c (make_edges): Handle GIMPLE_OMP_TARGET
and GIMPLE_OMP_TEAMS.
* omp-low.c (scan_sharing_clauses): Handle OMP_CLAUSE_DIST_SCHEDULE.
* gimple-low.c (lower_stmt): Handle GIMPLE_OMP_TARGET
and GIMPLE_OMP_TEAMS.
* tree.h (OMP_TEAMS_BODY, OMP_TEAMS_CLAUSES, OMP_TARGET_DATA_BODY,
OMP_TARGET_DATA_CLAUSES, OMP_TARGET_BODY, OMP_TARGET_CLAUSES,
OMP_TARGET_UPDATE_CLAUSES): Define.
* tree-nested.c (convert_nonlocal_reference_stmt,
convert_local_reference_stmt, convert_gimple_call): Handle
GIMPLE_OMP_TARGET and GIMPLE_OMP_TEAMS.
* tree-inline.c (estimate_num_insns): Likewise.
(remap_gimple_stmt): Likewise. Adjust gimple_build_omp_for
caller.
* gimple.def: Adjust comments describing OMP_CLAUSEs.
(GIMPLE_OMP_TARGET, GIMPLE_OMP_TEAMS): New GIMPLE stmts.
* tree-parloops.c (create_parallel_loop): Adjust gimple_build_omp_for
caller.
* tree-pretty-print.c (dump_generic_node): Handle OMP_TEAMS,
OMP_TARGET, OMP_TARGET_DATA and OMP_TARGET_UPDATE.
* gimple.h (GF_OMP_TARGET_KIND_MASK, GF_OMP_TARGET_KIND_REGION,
GF_OMP_TARGET_KIND_DATA, GF_OMP_TARGET_KIND_UPDATE): New.
(gimple_build_omp_for): Add kind argument to prototype.
(gimple_build_omp_target, gimple_build_omp_teams): New prototypes.
(gimple_has_substatements): Handle GIMPLE_OMP_TARGET and
GIMPLE_OMP_TEAMS.
(gimple_omp_subcode): Change GIMPLE_OMP_SINGLE to GIMPLE_OMP_TEAMS.
(gimple_omp_target_clauses, gimple_omp_target_clauses_ptr,
gimple_omp_target_set_clauses, gimple_omp_target_kind,
gimple_omp_target_set_kind, gimple_omp_teams_clauses,
gimple_omp_teams_clauses_ptr, gimple_omp_teams_set_clauses): New
inline functions.
(gimple_return_set_retval): Handle GIMPLE_OMP_TARGET and
GIMPLE_OMP_TEAMS.
* gimple.c (gimple_build_omp_for): Add kind argument, call
gimple_omp_for_set_kind.
(gimple_build_omp_target, gimple_build_omp_teams): New functions.
(walk_gimple_op, walk_gimple_stmt, gimple_copy): Handle
GIMPLE_OMP_TARGET and GIMPLE_OMP_TEAMS.
* gimple-pretty-print.c (dump_gimple_omp_target,
dump_gimple_omp_teams): New functions.
(pp_gimple_stmt_1): Handle GIMPLE_OMP_TARGET and GIMPLE_OMP_TEAMS.
* gimplify.c (enum gimplify_omp_var_data): Add GOVD_MAP.
(enum omp_region_type): Add ORT_TEAMS, ORT_TARGET and ORT_TARGET_DATA.
(omp_add_variable): Add temporary assertions.
(omp_notice_threadprivate_variable): Complain if threadprivate vars
appear in target region.
(omp_notice_variable): ORT_TARGET, ORT_TARGET_DATA and ORT_TEAMS
handling.
(omp_check_private): Ignore ORT_TARGET and ORT_TARGET_DATA regions.
(gimplify_scan_omp_clauses): Handle OMP_CLAUSE_MAP, OMP_CLAUSE_TO,
OMP_CLAUSE_FROM, OMP_CLAUSE_NUM_TEAMS, OMP_CLAUSE_DIST_SCHEDULE
and OMP_CLAUSE_DEVICE.
(gimplify_adjust_omp_clauses): Likewise.
(gimplify_adjust_omp_clauses_1): Handle GOVD_MAP. Fix up
check for privatization by also testing for GOVD_LINEAR.
(gimplify_omp_for): Adjust gimple_build_omp_for caller.
Clear *expr_p.
(gimplify_omp_workshare): Handle also OMP_TARGET, OMP_TARGET_DATA
and OMP_TEAMS. Clear *expr_p.
(gimplify_omp_target_update): New function.
(gimplify_expr): Handle OMP_TARGET, OMP_TARGET_DATA, OMP_TARGET_UPDATE
and OMP_TEAMS.
cp/
* parser.c (cp_parser_omp_clause_cancelkind): Remove diagnostics.
(cp_parser_omp_all_clauses): Require that OMP_CLAUSE_{TO,FROM}
and OMP_CLAUSE_{PARALLEL,FOR,SECTIONS,TASKGROUP} must be first in
the list of clauses.
(OMP_TEAMS_CLAUSE_MASK, OMP_TARGET_CLAUSE_MASK,
OMP_TARGET_DATA_CLAUSE_MASK, OMP_TARGET_UPDATE_CLAUSE_MASK,
OMP_DISTRIBUTE_CLAUSE_MASK): Define.
(cp_parser_omp_teams, cp_parser_omp_target, cp_parser_omp_target_data,
cp_parser_omp_target_update, cp_parser_omp_distribute): New functions.
(cp_parser_omp_construct): Handle PRAGMA_OMP_DISTRIBUTE and
PRAGMA_OMP_TEAMS.
(cp_parser_pragma): Handle PRAGMA_OMP_DISTRIBUTE, PRAGMA_OMP_TEAMS
and PRAGMA_OMP_TARGET.
* pt.c (tsubst_expr): Handle OMP_TEAMS, OMP_TARGET, OMP_TARGET_DATA
and OMP_TARGET_UPDATE.
--- gcc/tree.def.jj 2013-04-30 18:03:33.000000000 +0200
+++ gcc/tree.def 2013-05-22 15:22:45.143759788 +0200
@@ -1042,6 +1042,21 @@ DEFTREECODE (OMP_FOR_SIMD, "omp_for_simd
Operands like for OMP_FOR. */
DEFTREECODE (OMP_DISTRIBUTE, "omp_distribute", tcc_statement, 6)
+/* OpenMP - #pragma omp teams [clause1 ... clauseN]
+ Operand 0: OMP_TEAMS_BODY: Teams body.
+ Operand 1: OMP_TEAMS_CLAUSES: List of clauses. */
+DEFTREECODE (OMP_TEAMS, "omp_teams", tcc_statement, 2)
+
+/* OpenMP - #pragma omp target data [clause1 ... clauseN]
+ Operand 0: OMP_TARGET_DATA_BODY: Target data construct body.
+ Operand 1: OMP_TARGET_DATA_CLAUSES: List of clauses. */
+DEFTREECODE (OMP_TARGET_DATA, "omp_target_data", tcc_statement, 2)
+
+/* OpenMP - #pragma omp target [clause1 ... clauseN]
+ Operand 0: OMP_TARGET_BODY: Target construct body.
+ Operand 1: OMP_TARGET_CLAUSES: List of clauses. */
+DEFTREECODE (OMP_TARGET, "omp_target", tcc_statement, 2)
+
/* OpenMP - #pragma omp sections [clause1 ... clauseN]
Operand 0: OMP_SECTIONS_BODY: Sections body.
Operand 1: OMP_SECTIONS_CLAUSES: List of clauses. */
@@ -1069,6 +1084,10 @@ DEFTREECODE (OMP_ORDERED, "omp_ordered",
Operand 1: OMP_CRITICAL_NAME: Identifier for critical section. */
DEFTREECODE (OMP_CRITICAL, "omp_critical", tcc_statement, 2)
+/* OpenMP - #pragma omp target update [clause1 ... clauseN]
+ Operand 0: OMP_TARGET_UPDATE_CLAUSES: List of clauses. */
+DEFTREECODE (OMP_TARGET_UPDATE, "omp_target_update", tcc_statement, 1)
+
/* OMP_ATOMIC through OMP_ATOMIC_CAPTURE_NEW must be consecutive,
or OMP_ATOMIC_SEQ_CST needs adjusting. */
--- gcc/tree-cfg.c.jj 2013-05-20 13:21:43.000000000 +0200
+++ gcc/tree-cfg.c 2013-05-24 13:36:50.784061934 +0200
@@ -592,6 +592,8 @@ make_edges (void)
case GIMPLE_OMP_TASK:
case GIMPLE_OMP_FOR:
case GIMPLE_OMP_SINGLE:
+ case GIMPLE_OMP_TARGET:
+ case GIMPLE_OMP_TEAMS:
case GIMPLE_OMP_MASTER:
case GIMPLE_OMP_ORDERED:
case GIMPLE_OMP_CRITICAL:
--- gcc/omp-low.c.jj 2013-05-20 15:07:59.000000000 +0200
+++ gcc/omp-low.c 2013-05-24 13:47:32.465024619 +0200
@@ -1483,6 +1483,7 @@ scan_sharing_clauses (tree clauses, omp_
case OMP_CLAUSE_IF:
case OMP_CLAUSE_NUM_THREADS:
case OMP_CLAUSE_SCHEDULE:
+ case OMP_CLAUSE_DIST_SCHEDULE:
if (ctx->outer)
scan_omp_op (&OMP_CLAUSE_OPERAND (c, 0), ctx->outer);
break;
@@ -1548,6 +1549,7 @@ scan_sharing_clauses (tree clauses, omp_
case OMP_CLAUSE_IF:
case OMP_CLAUSE_NUM_THREADS:
case OMP_CLAUSE_SCHEDULE:
+ case OMP_CLAUSE_DIST_SCHEDULE:
case OMP_CLAUSE_NOWAIT:
case OMP_CLAUSE_ORDERED:
case OMP_CLAUSE_COLLAPSE:
--- gcc/gimple-low.c.jj 2013-03-20 10:07:24.000000000 +0100
+++ gcc/gimple-low.c 2013-05-24 13:40:19.382080371 +0200
@@ -444,6 +444,8 @@ lower_stmt (gimple_stmt_iterator *gsi, s
case GIMPLE_OMP_PARALLEL:
case GIMPLE_OMP_TASK:
+ case GIMPLE_OMP_TARGET:
+ case GIMPLE_OMP_TEAMS:
data->cannot_fallthru = false;
lower_omp_directive (gsi, data);
data->cannot_fallthru = false;
--- gcc/tree.h.jj 2013-05-20 13:18:33.000000000 +0200
+++ gcc/tree.h 2013-05-22 15:07:13.696048592 +0200
@@ -1863,6 +1863,20 @@ extern void protected_set_expr_location
#define OMP_CRITICAL_BODY(NODE) TREE_OPERAND (OMP_CRITICAL_CHECK (NODE), 0)
#define OMP_CRITICAL_NAME(NODE) TREE_OPERAND (OMP_CRITICAL_CHECK (NODE), 1)
+#define OMP_TEAMS_BODY(NODE) TREE_OPERAND (OMP_TEAMS_CHECK (NODE), 0)
+#define OMP_TEAMS_CLAUSES(NODE) TREE_OPERAND (OMP_TEAMS_CHECK
(NODE), 1)
+
+#define OMP_TARGET_DATA_BODY(NODE) \
+ TREE_OPERAND (OMP_TARGET_DATA_CHECK (NODE), 0)
+#define OMP_TARGET_DATA_CLAUSES(NODE)\
+ TREE_OPERAND (OMP_TARGET_DATA_CHECK (NODE), 1)
+
+#define OMP_TARGET_BODY(NODE) TREE_OPERAND (OMP_TARGET_CHECK (NODE), 0)
+#define OMP_TARGET_CLAUSES(NODE) TREE_OPERAND (OMP_TARGET_CHECK (NODE), 1)
+
+#define OMP_TARGET_UPDATE_CLAUSES(NODE)\
+ TREE_OPERAND (OMP_TARGET_UPDATE_CHECK (NODE), 0)
+
#define OMP_CLAUSE_CHAIN(NODE) TREE_CHAIN (OMP_CLAUSE_CHECK (NODE))
#define OMP_CLAUSE_DECL(NODE) \
OMP_CLAUSE_OPERAND (OMP_CLAUSE_RANGE_CHECK (OMP_CLAUSE_CHECK (NODE), \
--- gcc/tree-nested.c.jj 2013-03-20 10:08:27.000000000 +0100
+++ gcc/tree-nested.c 2013-05-23 13:23:48.785572843 +0200
@@ -1291,6 +1291,22 @@ convert_nonlocal_reference_stmt (gimple_
info->suppress_expansion = save_suppress;
break;
+ case GIMPLE_OMP_TARGET:
+ save_suppress = info->suppress_expansion;
+ convert_nonlocal_omp_clauses (gimple_omp_target_clauses_ptr (stmt), wi);
+ walk_body (convert_nonlocal_reference_stmt,
convert_nonlocal_reference_op,
+ info, gimple_omp_body_ptr (stmt));
+ info->suppress_expansion = save_suppress;
+ break;
+
+ case GIMPLE_OMP_TEAMS:
+ save_suppress = info->suppress_expansion;
+ convert_nonlocal_omp_clauses (gimple_omp_teams_clauses_ptr (stmt), wi);
+ walk_body (convert_nonlocal_reference_stmt,
convert_nonlocal_reference_op,
+ info, gimple_omp_body_ptr (stmt));
+ info->suppress_expansion = save_suppress;
+ break;
+
case GIMPLE_OMP_SECTION:
case GIMPLE_OMP_MASTER:
case GIMPLE_OMP_ORDERED:
@@ -1714,6 +1730,22 @@ convert_local_reference_stmt (gimple_stm
info->suppress_expansion = save_suppress;
break;
+ case GIMPLE_OMP_TARGET:
+ save_suppress = info->suppress_expansion;
+ convert_local_omp_clauses (gimple_omp_target_clauses_ptr (stmt), wi);
+ walk_body (convert_local_reference_stmt, convert_local_reference_op,
+ info, gimple_omp_body_ptr (stmt));
+ info->suppress_expansion = save_suppress;
+ break;
+
+ case GIMPLE_OMP_TEAMS:
+ save_suppress = info->suppress_expansion;
+ convert_local_omp_clauses (gimple_omp_teams_clauses_ptr (stmt), wi);
+ walk_body (convert_local_reference_stmt, convert_local_reference_op,
+ info, gimple_omp_body_ptr (stmt));
+ info->suppress_expansion = save_suppress;
+ break;
+
case GIMPLE_OMP_SECTION:
case GIMPLE_OMP_MASTER:
case GIMPLE_OMP_ORDERED:
@@ -2071,6 +2103,8 @@ convert_gimple_call (gimple_stmt_iterato
case GIMPLE_OMP_SECTIONS:
case GIMPLE_OMP_SECTION:
case GIMPLE_OMP_SINGLE:
+ case GIMPLE_OMP_TARGET:
+ case GIMPLE_OMP_TEAMS:
case GIMPLE_OMP_MASTER:
case GIMPLE_OMP_ORDERED:
case GIMPLE_OMP_CRITICAL:
--- gcc/tree-inline.c.jj 2013-05-13 16:49:40.000000000 +0200
+++ gcc/tree-inline.c 2013-05-23 13:53:52.258166517 +0200
@@ -1298,7 +1298,8 @@ remap_gimple_stmt (gimple stmt, copy_bod
case GIMPLE_OMP_FOR:
s1 = remap_gimple_seq (gimple_omp_body (stmt), id);
s2 = remap_gimple_seq (gimple_omp_for_pre_body (stmt), id);
- copy = gimple_build_omp_for (s1, gimple_omp_for_clauses (stmt),
+ copy = gimple_build_omp_for (s1, gimple_omp_for_kind (stmt),
+ gimple_omp_for_clauses (stmt),
gimple_omp_for_collapse (stmt), s2);
{
size_t i;
@@ -1345,6 +1346,19 @@ remap_gimple_stmt (gimple stmt, copy_bod
(s1, gimple_omp_single_clauses (stmt));
break;
+ case GIMPLE_OMP_TARGET:
+ s1 = remap_gimple_seq (gimple_omp_body (stmt), id);
+ copy = gimple_build_omp_target
+ (s1, gimple_omp_target_kind (stmt),
+ gimple_omp_target_clauses (stmt));
+ break;
+
+ case GIMPLE_OMP_TEAMS:
+ s1 = remap_gimple_seq (gimple_omp_body (stmt), id);
+ copy = gimple_build_omp_teams
+ (s1, gimple_omp_teams_clauses (stmt));
+ break;
+
case GIMPLE_OMP_CRITICAL:
s1 = remap_gimple_seq (gimple_omp_body (stmt), id);
copy
@@ -3716,6 +3730,8 @@ estimate_num_insns (gimple stmt, eni_wei
case GIMPLE_OMP_SECTION:
case GIMPLE_OMP_SECTIONS:
case GIMPLE_OMP_SINGLE:
+ case GIMPLE_OMP_TARGET:
+ case GIMPLE_OMP_TEAMS:
return (weights->omp_cost
+ estimate_num_insns_seq (gimple_omp_body (stmt), weights));
--- gcc/gimple.def.jj 2013-03-20 10:05:01.000000000 +0100
+++ gcc/gimple.def 2013-05-23 12:45:29.902586194 +0200
@@ -287,7 +287,7 @@ DEFGSCODE(GIMPLE_OMP_ORDERED, "gimple_om
BODY is a the sequence of statements to be executed by all threads.
- CLAUSES is a TREE_LIST node with all the clauses.
+ CLAUSES is an OMP_CLAUSE chain with all the clauses.
CHILD_FN is set when outlining the body of the parallel region.
All the statements in BODY are moved into this newly created
@@ -306,7 +306,7 @@ DEFGSCODE(GIMPLE_OMP_PARALLEL, "gimple_o
BODY is a the sequence of statements to be executed by all threads.
- CLAUSES is a TREE_LIST node with all the clauses.
+ CLAUSES is an OMP_CLAUSE chain with all the clauses.
CHILD_FN is set when outlining the body of the explicit task region.
All the statements in BODY are moved into this newly created
@@ -334,7 +334,7 @@ DEFGSCODE(GIMPLE_OMP_SECTION, "gimple_om
/* OMP_SECTIONS <BODY, CLAUSES, CONTROL> represents #pragma omp sections.
BODY is the sequence of statements in the sections body.
- CLAUSES is a TREE_LIST node holding the list of associated clauses.
+ CLAUSES is an OMP_CLAUSE chain holding the list of associated clauses.
CONTROL is a VAR_DECL used for deciding which of the sections
to execute. */
DEFGSCODE(GIMPLE_OMP_SECTIONS, "gimple_omp_sections", GSS_OMP_SECTIONS)
@@ -346,9 +346,21 @@ DEFGSCODE(GIMPLE_OMP_SECTIONS_SWITCH, "g
/* GIMPLE_OMP_SINGLE <BODY, CLAUSES> represents #pragma omp single
BODY is the sequence of statements inside the single section.
- CLAUSES is a TREE_LIST node holding the associated clauses. */
+ CLAUSES is an OMP_CLAUSE chain holding the associated clauses. */
DEFGSCODE(GIMPLE_OMP_SINGLE, "gimple_omp_single", GSS_OMP_SINGLE)
+/* GIMPLE_OMP_TARGET <BODY, CLAUSES> represents
+ #pragma omp target {,data,update}
+ BODY is the sequence of statements inside the target construct
+ (NULL for target update).
+ CLAUSES is an OMP_CLAUSE chain holding the associated clauses. */
+DEFGSCODE(GIMPLE_OMP_TARGET, "gimple_omp_target", GSS_OMP_SINGLE)
+
+/* GIMPLE_OMP_TEAMS <BODY, CLAUSES> represents #pragma omp teams
+ BODY is the sequence of statements inside the single section.
+ CLAUSES is an OMP_CLAUSE chain holding the associated clauses. */
+DEFGSCODE(GIMPLE_OMP_TEAMS, "gimple_omp_teams", GSS_OMP_SINGLE)
+
/* GIMPLE_PREDICT <PREDICT, OUTCOME> specifies a hint for branch prediction.
PREDICT is one of the predictors from predict.def.
--- gcc/cp/parser.c.jj 2013-05-20 13:21:27.000000000 +0200
+++ gcc/cp/parser.c 2013-05-22 19:12:25.523182096 +0200
@@ -26823,22 +26823,7 @@ cp_parser_omp_clause_cancelkind (cp_pars
enum omp_clause_code code,
tree list, location_t location)
{
- tree c;
-
- for (c = list; c; c = OMP_CLAUSE_CHAIN (c))
- switch (OMP_CLAUSE_CODE (c))
- {
- case OMP_CLAUSE_PARALLEL:
- case OMP_CLAUSE_FOR:
- case OMP_CLAUSE_SECTIONS:
- case OMP_CLAUSE_TASKGROUP:
- error_at (location, "only one of %<parallel%>, %<for%>, %<sections%> "
- "and %<taskgroup%> clauses can be specified");
- break;
- default:
- break;
- }
- c = build_omp_clause (location, code);
+ tree c = build_omp_clause (location, code);
OMP_CLAUSE_CHAIN (c) = list;
return c;
}
@@ -27260,7 +27245,6 @@ cp_parser_omp_all_clauses (cp_parser *pa
token = cp_lexer_peek_token (parser->lexer);
c_kind = cp_parser_omp_clause_name (parser);
- first = false;
switch (c_kind)
{
@@ -27359,31 +27343,48 @@ cp_parser_omp_all_clauses (cp_parser *pa
clauses = cp_parser_omp_clause_cancelkind (parser,
OMP_CLAUSE_PARALLEL,
clauses, token->location);
c_name = "parallel";
+ if (!first)
+ {
+ clause_not_first:
+ error_at (token->location, "%qs must be the first clause of %qs",
+ c_name, where);
+ clauses = prev;
+ }
break;
case PRAGMA_OMP_CLAUSE_FOR:
clauses = cp_parser_omp_clause_cancelkind (parser, OMP_CLAUSE_FOR,
clauses, token->location);
c_name = "for";
+ if (!first)
+ goto clause_not_first;
break;
case PRAGMA_OMP_CLAUSE_SECTIONS:
clauses = cp_parser_omp_clause_cancelkind (parser,
OMP_CLAUSE_SECTIONS,
clauses, token->location);
c_name = "sections";
+ if (!first)
+ goto clause_not_first;
break;
case PRAGMA_OMP_CLAUSE_TASKGROUP:
clauses = cp_parser_omp_clause_cancelkind (parser,
OMP_CLAUSE_TASKGROUP,
clauses, token->location);
c_name = "taskgroup";
+ if (!first)
+ goto clause_not_first;
break;
case PRAGMA_OMP_CLAUSE_TO:
clauses = cp_parser_omp_var_list (parser, OMP_CLAUSE_TO,
clauses);
c_name = "to";
+ if (!first)
+ goto clause_not_first;
break;
case PRAGMA_OMP_CLAUSE_FROM:
clauses = cp_parser_omp_var_list (parser, OMP_CLAUSE_FROM,
clauses);
c_name = "from";
+ if (!first)
+ goto clause_not_first;
break;
case PRAGMA_OMP_CLAUSE_UNIFORM:
clauses = cp_parser_omp_var_list (parser, OMP_CLAUSE_UNIFORM,
@@ -27441,6 +27442,8 @@ cp_parser_omp_all_clauses (cp_parser *pa
goto saw_error;
}
+ first = false;
+
if (((mask >> c_kind) & 1) == 0)
{
/* Remove the invalid clause(s) from the list to avoid
@@ -29013,6 +29016,180 @@ cp_parser_omp_cancellation_point (cp_par
}
/* OpenMP 4.0:
+ # pragma omp teams teams-clause[optseq] new-line
+ structured-block */
+
+#define OMP_TEAMS_CLAUSE_MASK \
+ ( (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_PRIVATE) \
+ | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_FIRSTPRIVATE) \
+ | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_SHARED) \
+ | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_REDUCTION) \
+ | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_NUM_TEAMS) \
+ | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_NUM_THREADS) \
+ | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_DEFAULT))
+
+static tree
+cp_parser_omp_teams (cp_parser *parser, cp_token *pragma_tok)
+{
+ tree stmt = make_node (OMP_TEAMS);
+ TREE_TYPE (stmt) = void_type_node;
+
+ OMP_TEAMS_CLAUSES (stmt)
+ = cp_parser_omp_all_clauses (parser, OMP_TEAMS_CLAUSE_MASK,
+ "#pragma omp teams", pragma_tok);
+ OMP_TEAMS_BODY (stmt) = cp_parser_omp_structured_block (parser);
+
+ return add_stmt (stmt);
+}
+
+/* OpenMP 4.0:
+ # pragma omp target data target-data-clause[optseq] new-line
+ structured-block */
+
+#define OMP_TARGET_DATA_CLAUSE_MASK \
+ ( (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_DEVICE) \
+ | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_MAP) \
+ | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_IF))
+
+static tree
+cp_parser_omp_target_data (cp_parser *parser, cp_token *pragma_tok)
+{
+ tree stmt = make_node (OMP_TARGET_DATA);
+ TREE_TYPE (stmt) = void_type_node;
+
+ OMP_TARGET_DATA_CLAUSES (stmt)
+ = cp_parser_omp_all_clauses (parser, OMP_TARGET_DATA_CLAUSE_MASK,
+ "#pragma omp target data", pragma_tok);
+ OMP_TARGET_DATA_BODY (stmt) = cp_parser_omp_structured_block (parser);
+
+ SET_EXPR_LOCATION (stmt, pragma_tok->location);
+ return add_stmt (stmt);
+}
+
+/* OpenMP 4.0:
+ # pragma omp target update target-update-clause[optseq] new-line */
+
+#define OMP_TARGET_UPDATE_CLAUSE_MASK \
+ ( (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_FROM) \
+ | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_TO) \
+ | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_DEVICE) \
+ | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_IF))
+
+static bool
+cp_parser_omp_target_update (cp_parser *parser, cp_token *pragma_tok,
+ enum pragma_context context)
+{
+ if (context == pragma_stmt)
+ {
+ error_at (pragma_tok->location,
+ "%<#pragma omp target update%> may only be "
+ "used in compound statements");
+ cp_parser_skip_to_pragma_eol (parser, pragma_tok);
+ return false;
+ }
+
+ tree clauses
+ = cp_parser_omp_all_clauses (parser, OMP_TARGET_UPDATE_CLAUSE_MASK,
+ "#pragma omp target update", pragma_tok);
+ if (find_omp_clause (clauses, OMP_CLAUSE_TO) == NULL_TREE
+ && find_omp_clause (clauses, OMP_CLAUSE_FROM) == NULL_TREE)
+ {
+ error_at (pragma_tok->location,
+ "%<#pragma omp target update must contain either "
+ "%<from%> or %<to%> clauses");
+ return false;
+ }
+
+ tree stmt = make_node (OMP_TARGET_UPDATE);
+ TREE_TYPE (stmt) = void_type_node;
+ OMP_TARGET_UPDATE_CLAUSES (stmt) = clauses;
+ SET_EXPR_LOCATION (stmt, pragma_tok->location);
+ add_stmt (stmt);
+ return false;
+}
+
+/* OpenMP 4.0:
+ # pragma omp target target-clause[optseq] new-line
+ structured-block */
+
+#define OMP_TARGET_CLAUSE_MASK \
+ ( (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_DEVICE) \
+ | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_MAP) \
+ | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_IF))
+
+static bool
+cp_parser_omp_target (cp_parser *parser, cp_token *pragma_tok,
+ enum pragma_context context)
+{
+ if (context != pragma_stmt && context != pragma_compound)
+ {
+ cp_parser_error (parser, "expected declaration specifiers");
+ cp_parser_skip_to_pragma_eol (parser, pragma_tok);
+ return false;
+ }
+
+ if (cp_lexer_next_token_is (parser->lexer, CPP_NAME))
+ {
+ tree id = cp_lexer_peek_token (parser->lexer)->u.value;
+ const char *p = IDENTIFIER_POINTER (id);
+
+ if (strcmp (p, "data") == 0)
+ {
+ cp_lexer_consume_token (parser->lexer);
+ cp_parser_omp_target_data (parser, pragma_tok);
+ return true;
+ }
+ else if (strcmp (p, "update") == 0)
+ {
+ cp_lexer_consume_token (parser->lexer);
+ return cp_parser_omp_target_update (parser, pragma_tok, context);
+ }
+ }
+
+ tree stmt = make_node (OMP_TARGET);
+ TREE_TYPE (stmt) = void_type_node;
+
+ OMP_TARGET_CLAUSES (stmt)
+ = cp_parser_omp_all_clauses (parser, OMP_TARGET_CLAUSE_MASK,
+ "#pragma omp target", pragma_tok);
+ OMP_TARGET_BODY (stmt) = cp_parser_omp_structured_block (parser);
+
+ SET_EXPR_LOCATION (stmt, pragma_tok->location);
+ add_stmt (stmt);
+ return true;
+}
+
+/* OpenMP 4.0:
+ #pragma omp distribute distribute-clause[optseq] new-line
+ for-loop */
+
+#define OMP_DISTRIBUTE_CLAUSE_MASK \
+ ( (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_PRIVATE) \
+ | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_FIRSTPRIVATE) \
+ | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_DIST_SCHEDULE)\
+ | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_COLLAPSE))
+
+static tree
+cp_parser_omp_distribute (cp_parser *parser, cp_token *pragma_tok)
+{
+ tree clauses, sb, ret;
+ unsigned int save;
+
+ clauses = cp_parser_omp_all_clauses (parser, OMP_DISTRIBUTE_CLAUSE_MASK,
+ "#pragma omp distribute", pragma_tok);
+
+ sb = begin_omp_structured_block ();
+ save = cp_parser_begin_omp_structured_block (parser);
+
+ ret = cp_parser_omp_for_loop (parser, OMP_DISTRIBUTE, clauses, NULL);
+
+ cp_parser_end_omp_structured_block (parser, save);
+ add_stmt (finish_omp_structured_block (sb));
+
+ return ret;
+}
+
+/* OpenMP 4.0:
# pragma omp declare simd declare-simd-clauses[optseq] new-line */
#define OMP_DECLARE_SIMD_CLAUSE_MASK \
@@ -29112,6 +29289,9 @@ cp_parser_omp_construct (cp_parser *pars
case PRAGMA_OMP_CRITICAL:
stmt = cp_parser_omp_critical (parser, pragma_tok);
break;
+ case PRAGMA_OMP_DISTRIBUTE:
+ stmt = cp_parser_omp_distribute (parser, pragma_tok);
+ break;
case PRAGMA_OMP_FOR:
stmt = cp_parser_omp_for (parser, pragma_tok);
break;
@@ -29139,6 +29319,9 @@ cp_parser_omp_construct (cp_parser *pars
case PRAGMA_OMP_TASKGROUP:
cp_parser_omp_taskgroup (parser, pragma_tok);
return;
+ case PRAGMA_OMP_TEAMS:
+ stmt = cp_parser_omp_teams (parser, pragma_tok);
+ break;
default:
gcc_unreachable ();
}
@@ -29609,6 +29792,7 @@ cp_parser_pragma (cp_parser *parser, enu
case PRAGMA_OMP_ATOMIC:
case PRAGMA_OMP_CRITICAL:
+ case PRAGMA_OMP_DISTRIBUTE:
case PRAGMA_OMP_FOR:
case PRAGMA_OMP_MASTER:
case PRAGMA_OMP_ORDERED:
@@ -29618,11 +29802,15 @@ cp_parser_pragma (cp_parser *parser, enu
case PRAGMA_OMP_SINGLE:
case PRAGMA_OMP_TASK:
case PRAGMA_OMP_TASKGROUP:
+ case PRAGMA_OMP_TEAMS:
if (context != pragma_stmt && context != pragma_compound)
goto bad_stmt;
cp_parser_omp_construct (parser, pragma_tok);
return true;
+ case PRAGMA_OMP_TARGET:
+ return cp_parser_omp_target (parser, pragma_tok, context);
+
case PRAGMA_OMP_SECTION:
error_at (pragma_tok->location,
"%<#pragma omp section%> may only be used in "
--- gcc/cp/pt.c.jj 2013-05-20 13:21:25.000000000 +0200
+++ gcc/cp/pt.c 2013-05-22 18:44:31.959390265 +0200
@@ -13330,6 +13330,9 @@ tsubst_expr (tree t, tree args, tsubst_f
case OMP_SECTIONS:
case OMP_SINGLE:
+ case OMP_TEAMS:
+ case OMP_TARGET_DATA:
+ case OMP_TARGET:
tmp = tsubst_omp_clauses (OMP_CLAUSES (t), false,
args, complain, in_decl);
stmt = push_stmt_list ();
@@ -13341,6 +13344,14 @@ tsubst_expr (tree t, tree args, tsubst_f
OMP_CLAUSES (t) = tmp;
add_stmt (t);
break;
+
+ case OMP_TARGET_UPDATE:
+ tmp = tsubst_omp_clauses (OMP_TARGET_UPDATE_CLAUSES (t), false,
+ args, complain, in_decl);
+ t = copy_node (t);
+ OMP_CLAUSES (t) = tmp;
+ add_stmt (t);
+ break;
case OMP_SECTION:
case OMP_CRITICAL:
--- gcc/tree-parloops.c.jj 2013-05-13 16:46:37.000000000 +0200
+++ gcc/tree-parloops.c 2013-05-23 13:54:21.854439781 +0200
@@ -1686,7 +1686,7 @@ create_parallel_loop (struct loop *loop,
t = build_omp_clause (loc, OMP_CLAUSE_SCHEDULE);
OMP_CLAUSE_SCHEDULE_KIND (t) = OMP_CLAUSE_SCHEDULE_STATIC;
- for_stmt = gimple_build_omp_for (NULL, t, 1, NULL);
+ for_stmt = gimple_build_omp_for (NULL, GF_OMP_FOR_KIND_FOR, t, 1, NULL);
gimple_set_location (for_stmt, loc);
gimple_omp_for_set_index (for_stmt, 0, initvar);
gimple_omp_for_set_initial (for_stmt, 0, cvar_init);
--- gcc/tree-pretty-print.c.jj 2013-05-20 13:18:24.000000000 +0200
+++ gcc/tree-pretty-print.c 2013-05-22 19:00:26.349842964 +0200
@@ -2347,6 +2347,27 @@ dump_generic_node (pretty_printer *buffe
pp_string (buffer, "#pragma omp distribute");
goto dump_omp_loop;
+ case OMP_TEAMS:
+ pp_string (buffer, "#pragma omp teams");
+ dump_omp_clauses (buffer, OMP_TEAMS_CLAUSES (node), spc, flags);
+ goto dump_omp_body;
+
+ case OMP_TARGET_DATA:
+ pp_string (buffer, "#pragma omp target data");
+ dump_omp_clauses (buffer, OMP_TARGET_DATA_CLAUSES (node), spc, flags);
+ goto dump_omp_body;
+
+ case OMP_TARGET:
+ pp_string (buffer, "#pragma omp target");
+ dump_omp_clauses (buffer, OMP_TARGET_CLAUSES (node), spc, flags);
+ goto dump_omp_body;
+
+ case OMP_TARGET_UPDATE:
+ pp_string (buffer, "#pragma omp target update");
+ dump_omp_clauses (buffer, OMP_TARGET_UPDATE_CLAUSES (node), spc, flags);
+ is_expr = false;
+ break;
+
dump_omp_loop:
dump_omp_clauses (buffer, OMP_FOR_CLAUSES (node), spc, flags);
--- gcc/gimple.h.jj 2013-05-13 16:49:47.000000000 +0200
+++ gcc/gimple.h 2013-05-23 14:06:41.729631141 +0200
@@ -115,6 +115,10 @@ enum gf_mask {
GF_OMP_FOR_KIND_SIMD = 1 << 0,
GF_OMP_FOR_KIND_FOR_SIMD = 2 << 0,
GF_OMP_FOR_KIND_DISTRIBUTE = 3 << 0,
+ GF_OMP_TARGET_KIND_MASK = 3 << 0,
+ GF_OMP_TARGET_KIND_REGION = 0 << 0,
+ GF_OMP_TARGET_KIND_DATA = 1 << 0,
+ GF_OMP_TARGET_KIND_UPDATE = 2 << 0,
/* True on an GIMPLE_OMP_RETURN statement if the return does not require
a thread synchronization via some sort of barrier. The exact barrier
@@ -618,7 +622,7 @@ struct GTY(()) gimple_statement_omp_cont
tree control_use;
};
-/* GIMPLE_OMP_SINGLE */
+/* GIMPLE_OMP_SINGLE, GIMPLE_OMP_TARGET, GIMPLE_OMP_TEAMS */
struct GTY(()) gimple_statement_omp_single {
/* [ WORD 1-7 ] */
@@ -805,7 +809,7 @@ gimple gimple_build_switch_nlabels (unsi
gimple gimple_build_switch (tree, tree, vec<tree> );
gimple gimple_build_omp_parallel (gimple_seq, tree, tree, tree);
gimple gimple_build_omp_task (gimple_seq, tree, tree, tree, tree, tree, tree);
-gimple gimple_build_omp_for (gimple_seq, tree, size_t, gimple_seq);
+gimple gimple_build_omp_for (gimple_seq, int, tree, size_t, gimple_seq);
gimple gimple_build_omp_critical (gimple_seq, tree);
gimple gimple_build_omp_section (gimple_seq);
gimple gimple_build_omp_continue (tree, tree);
@@ -815,6 +819,8 @@ gimple gimple_build_omp_ordered (gimple_
gimple gimple_build_omp_sections (gimple_seq, tree);
gimple gimple_build_omp_sections_switch (void);
gimple gimple_build_omp_single (gimple_seq, tree);
+gimple gimple_build_omp_target (gimple_seq, int, tree);
+gimple gimple_build_omp_teams (gimple_seq, tree);
gimple gimple_build_cdt (tree, tree);
gimple gimple_build_omp_atomic_load (tree, tree);
gimple gimple_build_omp_atomic_store (tree);
@@ -1264,6 +1270,8 @@ gimple_has_substatements (gimple g)
case GIMPLE_OMP_TASK:
case GIMPLE_OMP_SECTIONS:
case GIMPLE_OMP_SINGLE:
+ case GIMPLE_OMP_TARGET:
+ case GIMPLE_OMP_TEAMS:
case GIMPLE_OMP_CRITICAL:
case GIMPLE_WITH_CLEANUP_EXPR:
case GIMPLE_TRANSACTION:
@@ -1691,7 +1699,7 @@ static inline unsigned
gimple_omp_subcode (const_gimple s)
{
gcc_gimple_checking_assert (gimple_code (s) >= GIMPLE_OMP_ATOMIC_LOAD
- && gimple_code (s) <= GIMPLE_OMP_SINGLE);
+ && gimple_code (s) <= GIMPLE_OMP_TEAMS);
return s->gsbase.subcode;
}
@@ -4604,6 +4612,87 @@ gimple_omp_single_set_clauses (gimple gs
}
+/* Return the clauses associated with OMP_TARGET GS. */
+
+static inline tree
+gimple_omp_target_clauses (const_gimple gs)
+{
+ GIMPLE_CHECK (gs, GIMPLE_OMP_TARGET);
+ return gs->gimple_omp_single.clauses;
+}
+
+
+/* Return a pointer to the clauses associated with OMP_TARGET GS. */
+
+static inline tree *
+gimple_omp_target_clauses_ptr (gimple gs)
+{
+ GIMPLE_CHECK (gs, GIMPLE_OMP_TARGET);
+ return &gs->gimple_omp_single.clauses;
+}
+
+
+/* Set CLAUSES to be the clauses associated with OMP_TARGET GS. */
+
+static inline void
+gimple_omp_target_set_clauses (gimple gs, tree clauses)
+{
+ GIMPLE_CHECK (gs, GIMPLE_OMP_TARGET);
+ gs->gimple_omp_single.clauses = clauses;
+}
+
+
+/* Return the kind of OMP target statemement. */
+
+static inline int
+gimple_omp_target_kind (const_gimple g)
+{
+ GIMPLE_CHECK (g, GIMPLE_OMP_TARGET);
+ return (gimple_omp_subcode (g) & GF_OMP_TARGET_KIND_MASK);
+}
+
+
+/* Set the OMP target kind. */
+
+static inline void
+gimple_omp_target_set_kind (gimple g, int kind)
+{
+ GIMPLE_CHECK (g, GIMPLE_OMP_TARGET);
+ g->gsbase.subcode = (g->gsbase.subcode & ~GF_OMP_TARGET_KIND_MASK)
+ | (kind & GF_OMP_TARGET_KIND_MASK);
+}
+
+
+/* Return the clauses associated with OMP_TEAMS GS. */
+
+static inline tree
+gimple_omp_teams_clauses (const_gimple gs)
+{
+ GIMPLE_CHECK (gs, GIMPLE_OMP_TEAMS);
+ return gs->gimple_omp_single.clauses;
+}
+
+
+/* Return a pointer to the clauses associated with OMP_TEAMS GS. */
+
+static inline tree *
+gimple_omp_teams_clauses_ptr (gimple gs)
+{
+ GIMPLE_CHECK (gs, GIMPLE_OMP_TEAMS);
+ return &gs->gimple_omp_single.clauses;
+}
+
+
+/* Set CLAUSES to be the clauses associated with OMP_TEAMS GS. */
+
+static inline void
+gimple_omp_teams_set_clauses (gimple gs, tree clauses)
+{
+ GIMPLE_CHECK (gs, GIMPLE_OMP_TEAMS);
+ gs->gimple_omp_single.clauses = clauses;
+}
+
+
/* Return the clauses associated with OMP_SECTIONS GS. */
static inline tree
@@ -4946,6 +5035,8 @@ gimple_return_set_retval (gimple gs, tre
case GIMPLE_OMP_SECTIONS: \
case GIMPLE_OMP_SECTIONS_SWITCH: \
case GIMPLE_OMP_SINGLE: \
+ case GIMPLE_OMP_TARGET: \
+ case GIMPLE_OMP_TEAMS: \
case GIMPLE_OMP_SECTION: \
case GIMPLE_OMP_MASTER: \
case GIMPLE_OMP_ORDERED: \
--- gcc/gimple.c.jj 2013-05-13 16:49:46.000000000 +0200
+++ gcc/gimple.c 2013-05-23 13:40:12.487789257 +0200
@@ -908,13 +908,14 @@ gimple_build_omp_critical (gimple_seq bo
PRE_BODY is the sequence of statements that are loop invariant. */
gimple
-gimple_build_omp_for (gimple_seq body, tree clauses, size_t collapse,
+gimple_build_omp_for (gimple_seq body, int kind, tree clauses, size_t collapse,
gimple_seq pre_body)
{
gimple p = gimple_alloc (GIMPLE_OMP_FOR, 0);
if (body)
gimple_omp_set_body (p, body);
gimple_omp_for_set_clauses (p, clauses);
+ gimple_omp_for_set_kind (p, kind);
p->gimple_omp_for.collapse = collapse;
p->gimple_omp_for.iter
= ggc_alloc_cleared_vec_gimple_omp_for_iter (collapse);
@@ -1094,6 +1095,41 @@ gimple_build_omp_single (gimple_seq body
}
+/* Build a GIMPLE_OMP_TARGET statement.
+
+ BODY is the sequence of statements that will be executed.
+ CLAUSES are any of the OMP target construct's clauses. */
+
+gimple
+gimple_build_omp_target (gimple_seq body, int kind, tree clauses)
+{
+ gimple p = gimple_alloc (GIMPLE_OMP_TARGET, 0);
+ if (body)
+ gimple_omp_set_body (p, body);
+ gimple_omp_target_set_clauses (p, clauses);
+ gimple_omp_target_set_kind (p, kind);
+
+ return p;
+}
+
+
+/* Build a GIMPLE_OMP_TEAMS statement.
+
+ BODY is the sequence of statements that will be executed.
+ CLAUSES are any of the OMP teams construct's clauses. */
+
+gimple
+gimple_build_omp_teams (gimple_seq body, tree clauses)
+{
+ gimple p = gimple_alloc (GIMPLE_OMP_TEAMS, 0);
+ if (body)
+ gimple_omp_set_body (p, body);
+ gimple_omp_teams_set_clauses (p, clauses);
+
+ return p;
+}
+
+
/* Build a GIMPLE_OMP_ATOMIC_LOAD statement. */
gimple
@@ -1610,6 +1646,20 @@ walk_gimple_op (gimple stmt, walk_tree_f
return ret;
break;
+ case GIMPLE_OMP_TARGET:
+ ret = walk_tree (gimple_omp_target_clauses_ptr (stmt), callback_op, wi,
+ pset);
+ if (ret)
+ return ret;
+ break;
+
+ case GIMPLE_OMP_TEAMS:
+ ret = walk_tree (gimple_omp_teams_clauses_ptr (stmt), callback_op, wi,
+ pset);
+ if (ret)
+ return ret;
+ break;
+
case GIMPLE_OMP_ATOMIC_LOAD:
ret = walk_tree (gimple_omp_atomic_load_lhs_ptr (stmt), callback_op, wi,
pset);
@@ -1786,6 +1836,8 @@ walk_gimple_stmt (gimple_stmt_iterator *
case GIMPLE_OMP_TASK:
case GIMPLE_OMP_SECTIONS:
case GIMPLE_OMP_SINGLE:
+ case GIMPLE_OMP_TARGET:
+ case GIMPLE_OMP_TEAMS:
ret = walk_gimple_seq_mod (gimple_omp_body_ptr (stmt), callback_stmt,
callback_op, wi);
if (ret)
@@ -2308,6 +2360,8 @@ gimple_copy (gimple stmt)
/* FALLTHRU */
case GIMPLE_OMP_SINGLE:
+ case GIMPLE_OMP_TARGET:
+ case GIMPLE_OMP_TEAMS:
case GIMPLE_OMP_SECTION:
case GIMPLE_OMP_MASTER:
case GIMPLE_OMP_ORDERED:
--- gcc/gimple-pretty-print.c.jj 2013-05-13 16:49:01.000000000 +0200
+++ gcc/gimple-pretty-print.c 2013-05-23 13:17:37.589649980 +0200
@@ -1264,6 +1264,78 @@ dump_gimple_omp_single (pretty_printer *
}
}
+/* Dump a GIMPLE_OMP_TARGET tuple on the pretty_printer BUFFER. */
+
+static void
+dump_gimple_omp_target (pretty_printer *buffer, gimple gs, int spc, int flags)
+{
+ const char *kind;
+ switch (gimple_omp_target_kind (gs))
+ {
+ case GF_OMP_TARGET_KIND_REGION:
+ kind = "";
+ break;
+ case GF_OMP_TARGET_KIND_DATA:
+ kind = " data";
+ break;
+ case GF_OMP_TARGET_KIND_UPDATE:
+ kind = " update";
+ break;
+ default:
+ gcc_unreachable ();
+ }
+ if (flags & TDF_RAW)
+ {
+ 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, " >");
+ }
+ else
+ {
+ pp_string (buffer, "#pragma omp target");
+ pp_string (buffer, kind);
+ dump_omp_clauses (buffer, gimple_omp_target_clauses (gs), spc, flags);
+ if (!gimple_seq_empty_p (gimple_omp_body (gs)))
+ {
+ newline_and_indent (buffer, spc + 2);
+ pp_character (buffer, '{');
+ pp_newline (buffer);
+ dump_gimple_seq (buffer, gimple_omp_body (gs), spc + 4, flags);
+ newline_and_indent (buffer, spc + 2);
+ pp_character (buffer, '}');
+ }
+ }
+}
+
+/* Dump a GIMPLE_OMP_TEAMS tuple on the pretty_printer BUFFER. */
+
+static void
+dump_gimple_omp_teams (pretty_printer *buffer, gimple gs, int spc, int flags)
+{
+ if (flags & TDF_RAW)
+ {
+ dump_gimple_fmt (buffer, spc, flags, "%G <%+BODY <%S>%nCLAUSES <", gs,
+ gimple_omp_body (gs));
+ dump_omp_clauses (buffer, gimple_omp_teams_clauses (gs), spc, flags);
+ dump_gimple_fmt (buffer, spc, flags, " >");
+ }
+ else
+ {
+ pp_string (buffer, "#pragma omp teams");
+ dump_omp_clauses (buffer, gimple_omp_teams_clauses (gs), spc, flags);
+ if (!gimple_seq_empty_p (gimple_omp_body (gs)))
+ {
+ newline_and_indent (buffer, spc + 2);
+ pp_character (buffer, '{');
+ pp_newline (buffer);
+ dump_gimple_seq (buffer, gimple_omp_body (gs), spc + 4, flags);
+ newline_and_indent (buffer, spc + 2);
+ pp_character (buffer, '}');
+ }
+ }
+}
+
/* Dump a GIMPLE_OMP_SECTIONS tuple on the pretty_printer BUFFER. */
static void
@@ -2038,6 +2110,14 @@ pp_gimple_stmt_1 (pretty_printer *buffer
dump_gimple_omp_single (buffer, gs, spc, flags);
break;
+ case GIMPLE_OMP_TARGET:
+ dump_gimple_omp_target (buffer, gs, spc, flags);
+ break;
+
+ case GIMPLE_OMP_TEAMS:
+ dump_gimple_omp_teams (buffer, gs, spc, flags);
+ break;
+
case GIMPLE_OMP_RETURN:
dump_gimple_omp_return (buffer, gs, spc, flags);
break;
--- gcc/gimplify.c.jj 2013-05-13 16:49:10.000000000 +0200
+++ gcc/gimplify.c 2013-05-24 14:04:01.534225316 +0200
@@ -57,10 +57,11 @@ enum gimplify_omp_var_data
GOVD_LASTPRIVATE = 32,
GOVD_REDUCTION = 64,
GOVD_LOCAL = 128,
- GOVD_DEBUG_PRIVATE = 256,
- GOVD_PRIVATE_OUTER_REF = 512,
- GOVD_LINEAR = 1024,
- GOVD_ALIGNED = 2048,
+ GOVD_MAP = 256,
+ GOVD_DEBUG_PRIVATE = 512,
+ GOVD_PRIVATE_OUTER_REF = 1024,
+ GOVD_LINEAR = 2048,
+ GOVD_ALIGNED = 4096,
GOVD_DATA_SHARE_CLASS = (GOVD_SHARED | GOVD_PRIVATE | GOVD_FIRSTPRIVATE
| GOVD_LASTPRIVATE | GOVD_REDUCTION | GOVD_LINEAR
| GOVD_LOCAL)
@@ -74,7 +75,10 @@ enum omp_region_type
ORT_PARALLEL = 2,
ORT_COMBINED_PARALLEL = 3,
ORT_TASK = 4,
- ORT_UNTIED_TASK = 5
+ ORT_UNTIED_TASK = 5,
+ ORT_TEAMS = 8,
+ ORT_TARGET_DATA = 16,
+ ORT_TARGET = 32
};
struct gimplify_omp_ctx
@@ -5829,6 +5833,9 @@ omp_add_variable (struct gimplify_omp_ct
the parameters of the type. */
if (DECL_SIZE (decl) && TREE_CODE (DECL_SIZE (decl)) != INTEGER_CST)
{
+ /* To be handled later. */
+ gcc_assert ((flags & GOVD_MAP) == 0);
+
/* Add the pointer replacement variable as PRIVATE if the variable
replacement is private, else FIRSTPRIVATE since we'll need the
address of the original variable either for SHARED, or for the
@@ -5870,6 +5877,9 @@ omp_add_variable (struct gimplify_omp_ct
}
else if (lang_hooks.decls.omp_privatize_by_reference (decl))
{
+ /* To be handled later. */
+ gcc_assert ((flags & GOVD_MAP) == 0);
+
gcc_assert ((flags & GOVD_LOCAL) == 0);
omp_firstprivatize_type_sizes (ctx, TREE_TYPE (decl));
@@ -5896,6 +5906,22 @@ omp_notice_threadprivate_variable (struc
tree decl2)
{
splay_tree_node n;
+ struct gimplify_omp_ctx *octx;
+
+ for (octx = ctx; octx; octx = octx->outer_context)
+ if (octx->region_type == ORT_TARGET)
+ {
+ n = splay_tree_lookup (ctx->variables, (splay_tree_key)decl);
+ if (n == NULL)
+ {
+ error ("threadprivate variable %qE used in target region",
+ DECL_NAME (decl));
+ error_at (octx->location, "enclosing target region");
+ splay_tree_insert (octx->variables, (splay_tree_key)decl, 0);
+ }
+ if (decl2)
+ splay_tree_insert (octx->variables, (splay_tree_key)decl2, 0);
+ }
if (ctx->region_type != ORT_UNTIED_TASK)
return false;
@@ -5944,13 +5970,24 @@ omp_notice_variable (struct gimplify_omp
}
n = splay_tree_lookup (ctx->variables, (splay_tree_key)decl);
+ if (ctx->region_type == ORT_TARGET)
+ {
+ if (n == NULL)
+ omp_add_variable (ctx, decl, GOVD_MAP | flags);
+ else
+ n->value |= flags;
+ ret = lang_hooks.decls.omp_disregard_value_expr (decl, true);
+ goto do_outer;
+ }
+
if (n == NULL)
{
enum omp_clause_default_kind default_kind, kind;
struct gimplify_omp_ctx *octx;
if (ctx->region_type == ORT_WORKSHARE
- || ctx->region_type == ORT_SIMD)
+ || ctx->region_type == ORT_SIMD
+ || ctx->region_type == ORT_TARGET_DATA)
goto do_outer;
/* ??? Some compiler-generated variables (like SAVE_EXPRs) could be
@@ -5964,12 +6001,24 @@ omp_notice_variable (struct gimplify_omp
switch (default_kind)
{
case OMP_CLAUSE_DEFAULT_NONE:
- error ("%qE not specified in enclosing parallel",
- DECL_NAME (lang_hooks.decls.omp_report_decl (decl)));
if ((ctx->region_type & ORT_TASK) != 0)
- error_at (ctx->location, "enclosing task");
+ {
+ error ("%qE not specified in enclosing task",
+ DECL_NAME (lang_hooks.decls.omp_report_decl (decl)));
+ error_at (ctx->location, "enclosing task");
+ }
+ else if (ctx->region_type == ORT_TEAMS)
+ {
+ error ("%qE not specified in enclosing teams construct",
+ DECL_NAME (lang_hooks.decls.omp_report_decl (decl)));
+ error_at (ctx->location, "enclosing teams construct");
+ }
else
- error_at (ctx->location, "enclosing parallel");
+ {
+ error ("%qE not specified in enclosing parallel",
+ DECL_NAME (lang_hooks.decls.omp_report_decl (decl)));
+ error_at (ctx->location, "enclosing parallel");
+ }
/* FALLTHRU */
case OMP_CLAUSE_DEFAULT_SHARED:
flags |= GOVD_SHARED;
@@ -5989,13 +6038,15 @@ omp_notice_variable (struct gimplify_omp
{
splay_tree_node n2;
+ if ((octx->region_type & (ORT_TARGET_DATA | ORT_TARGET)) != 0)
+ continue;
n2 = splay_tree_lookup (octx->variables, (splay_tree_key) decl);
if (n2 && (n2->value & GOVD_DATA_SHARE_CLASS) != GOVD_SHARED)
{
flags |= GOVD_FIRSTPRIVATE;
break;
}
- if ((octx->region_type & ORT_PARALLEL) != 0)
+ if ((octx->region_type & (ORT_PARALLEL | ORT_TEAMS)) != 0)
break;
}
if (flags & GOVD_FIRSTPRIVATE)
@@ -6137,6 +6188,9 @@ omp_check_private (struct gimplify_omp_c
/* References might be private, but might be shared too. */
|| lang_hooks.decls.omp_privatize_by_reference (decl));
+ if ((ctx->region_type & (ORT_TARGET | ORT_TARGET_DATA)) != 0)
+ continue;
+
n = splay_tree_lookup (ctx->variables, (splay_tree_key) decl);
if (n != NULL)
return (n->value & GOVD_SHARED) == 0;
@@ -6204,6 +6258,20 @@ gimplify_scan_omp_clauses (tree *list_p,
}
flags = GOVD_LINEAR | GOVD_EXPLICIT;
goto do_add;
+ case OMP_CLAUSE_MAP:
+ flags = GOVD_MAP | GOVD_EXPLICIT;
+ notice_outer = false;
+ goto do_add;
+
+ case OMP_CLAUSE_TO:
+ case OMP_CLAUSE_FROM:
+ decl = OMP_CLAUSE_DECL (c);
+ if (error_operand_p (decl))
+ {
+ remove = true;
+ break;
+ }
+ goto do_notice;
do_add:
decl = OMP_CLAUSE_DECL (c);
@@ -6292,6 +6360,9 @@ gimplify_scan_omp_clauses (tree *list_p,
case OMP_CLAUSE_SCHEDULE:
case OMP_CLAUSE_NUM_THREADS:
+ case OMP_CLAUSE_NUM_TEAMS:
+ case OMP_CLAUSE_DIST_SCHEDULE:
+ case OMP_CLAUSE_DEVICE:
if (gimplify_expr (&OMP_CLAUSE_OPERAND (c, 0), pre_p, NULL,
is_gimple_val, fb_rvalue) == GS_ERROR)
remove = true;
@@ -6357,12 +6428,40 @@ gimplify_adjust_omp_clauses_1 (splay_tre
gcc_assert ((flags & GOVD_DATA_SHARE_CLASS) == GOVD_PRIVATE);
private_debug = true;
}
+ else if (flags & GOVD_MAP)
+ private_debug = false;
else
private_debug
= lang_hooks.decls.omp_private_debug_clause (decl,
!!(flags & GOVD_SHARED));
if (private_debug)
code = OMP_CLAUSE_PRIVATE;
+ else if (flags & GOVD_MAP)
+ {
+ /* If decl is already in the enclosing device data environment,
+ the spec says that it should just be used and no init/assignment
+ should be done. If there was any privatization in between though,
+ it means that original decl might be in the enclosing device data
+ environment, but the privatized might not. */
+ struct gimplify_omp_ctx *ctx;
+ for (ctx = gimplify_omp_ctxp->outer_context;
+ ctx; ctx = ctx->outer_context)
+ {
+ n = splay_tree_lookup (ctx->variables, (splay_tree_key) decl);
+ if (n == NULL)
+ continue;
+ if (ctx->region_type == ORT_TARGET_DATA)
+ {
+ if ((n->value & GOVD_MAP) != 0)
+ return 0;
+ }
+ else if ((n->value & (GOVD_FIRSTPRIVATE | GOVD_LASTPRIVATE
+ | GOVD_PRIVATE | GOVD_REDUCTION
+ | GOVD_LINEAR)) != 0)
+ break;
+ }
+ code = OMP_CLAUSE_MAP;
+ }
else if (flags & GOVD_SHARED)
{
if (is_global_var (decl))
@@ -6373,7 +6472,8 @@ gimplify_adjust_omp_clauses_1 (splay_tre
splay_tree_node on
= splay_tree_lookup (ctx->variables, (splay_tree_key) decl);
if (on && (on->value & (GOVD_FIRSTPRIVATE | GOVD_LASTPRIVATE
- | GOVD_PRIVATE | GOVD_REDUCTION)) != 0)
+ | GOVD_PRIVATE | GOVD_REDUCTION
+ | GOVD_LINEAR)) != 0)
break;
ctx = ctx->outer_context;
}
@@ -6400,6 +6500,8 @@ gimplify_adjust_omp_clauses_1 (splay_tre
OMP_CLAUSE_PRIVATE_DEBUG (clause) = 1;
else if (code == OMP_CLAUSE_PRIVATE && (flags & GOVD_PRIVATE_OUTER_REF))
OMP_CLAUSE_PRIVATE_OUTER_REF (clause) = 1;
+ else if (code == OMP_CLAUSE_MAP)
+ OMP_CLAUSE_MAP_KIND (clause) = OMP_CLAUSE_MAP_TOFROM;
*list_p = clause;
lang_hooks.decls.omp_finish_clause (clause);
@@ -6517,11 +6619,47 @@ gimplify_adjust_omp_clauses (tree *list_
}
break;
+ case OMP_CLAUSE_MAP:
+ decl = OMP_CLAUSE_DECL (c);
+ n = splay_tree_lookup (ctx->variables, (splay_tree_key) decl);
+ remove = false;
+ if (ctx->region_type == ORT_TARGET && !(n->value & GOVD_SEEN))
+ remove = true;
+ else
+ {
+ /* If decl is already in the enclosing device data environment,
+ the spec says that it should just be used and no
init/assignment
+ should be done. If there was any privatization in between
though,
+ it means that original decl might be in the enclosing device
data
+ environment, but the privatized might not. */
+ struct gimplify_omp_ctx *octx;
+ for (octx = ctx->outer_context; octx; octx = octx->outer_context)
+ {
+ n = splay_tree_lookup (octx->variables,
+ (splay_tree_key) decl);
+ if (n == NULL)
+ continue;
+ if (octx->region_type == ORT_TARGET_DATA)
+ {
+ if ((n->value & GOVD_MAP) != 0)
+ remove = true;
+ }
+ else if ((n->value & (GOVD_FIRSTPRIVATE | GOVD_LASTPRIVATE
+ | GOVD_PRIVATE | GOVD_REDUCTION
+ | GOVD_LINEAR)) != 0)
+ break;
+ }
+ }
+ break;
+
case OMP_CLAUSE_REDUCTION:
case OMP_CLAUSE_COPYIN:
case OMP_CLAUSE_COPYPRIVATE:
case OMP_CLAUSE_IF:
case OMP_CLAUSE_NUM_THREADS:
+ case OMP_CLAUSE_NUM_TEAMS:
+ case OMP_CLAUSE_DIST_SCHEDULE:
+ case OMP_CLAUSE_DEVICE:
case OMP_CLAUSE_SCHEDULE:
case OMP_CLAUSE_NOWAIT:
case OMP_CLAUSE_ORDERED:
@@ -6532,6 +6670,8 @@ gimplify_adjust_omp_clauses (tree *list_
case OMP_CLAUSE_MERGEABLE:
case OMP_CLAUSE_PROC_BIND:
case OMP_CLAUSE_SAFELEN:
+ case OMP_CLAUSE_TO:
+ case OMP_CLAUSE_FROM:
break;
default:
@@ -6847,25 +6987,19 @@ gimplify_omp_for (tree *expr_p, gimple_s
gimplify_adjust_omp_clauses (&OMP_FOR_CLAUSES (for_stmt));
- gfor = gimple_build_omp_for (for_body, OMP_FOR_CLAUSES (for_stmt),
- TREE_VEC_LENGTH (OMP_FOR_INIT (for_stmt)),
- for_pre_body);
+ int kind;
switch (TREE_CODE (for_stmt))
{
- case OMP_FOR:
- break;
- case OMP_SIMD:
- gimple_omp_for_set_kind (gfor, GF_OMP_FOR_KIND_SIMD);
- break;
- case OMP_FOR_SIMD:
- gimple_omp_for_set_kind (gfor, GF_OMP_FOR_KIND_FOR_SIMD);
- break;
- case OMP_DISTRIBUTE:
- gimple_omp_for_set_kind (gfor, GF_OMP_FOR_KIND_DISTRIBUTE);
- break;
+ case OMP_FOR: kind = GF_OMP_FOR_KIND_FOR; break;
+ case OMP_SIMD: kind = GF_OMP_FOR_KIND_SIMD; break;
+ case OMP_FOR_SIMD: kind = GF_OMP_FOR_KIND_FOR_SIMD; break;
+ case OMP_DISTRIBUTE: kind = GF_OMP_FOR_KIND_DISTRIBUTE; break;
default:
gcc_unreachable ();
}
+ gfor = gimple_build_omp_for (for_body, kind, OMP_FOR_CLAUSES (for_stmt),
+ TREE_VEC_LENGTH (OMP_FOR_INIT (for_stmt)),
+ for_pre_body);
for (i = 0; i < TREE_VEC_LENGTH (OMP_FOR_INIT (for_stmt)); i++)
{
@@ -6880,11 +7014,15 @@ gimplify_omp_for (tree *expr_p, gimple_s
}
gimplify_seq_add_stmt (pre_p, gfor);
- return ret == GS_ALL_DONE ? GS_ALL_DONE : GS_ERROR;
+ if (ret != GS_ALL_DONE)
+ return GS_ERROR;
+ *expr_p = NULL_TREE;
+ return GS_ALL_DONE;
}
-/* Gimplify the gross structure of other OpenMP worksharing constructs.
- In particular, OMP_SECTIONS and OMP_SINGLE. */
+/* Gimplify the gross structure of other OpenMP constructs.
+ In particular, OMP_SECTIONS, OMP_SINGLE, OMP_TARGET, OMP_TARGET_DATA
+ and OMP_TEAMS. */
static void
gimplify_omp_workshare (tree *expr_p, gimple_seq *pre_p)
@@ -6892,19 +7030,72 @@ gimplify_omp_workshare (tree *expr_p, gi
tree expr = *expr_p;
gimple stmt;
gimple_seq body = NULL;
+ enum omp_region_type ort = ORT_WORKSHARE;
- gimplify_scan_omp_clauses (&OMP_CLAUSES (expr), pre_p, ORT_WORKSHARE);
+ switch (TREE_CODE (expr))
+ {
+ case OMP_SECTIONS:
+ case OMP_SINGLE:
+ break;
+ case OMP_TARGET:
+ ort = ORT_TARGET;
+ break;
+ case OMP_TARGET_DATA:
+ ort = ORT_TARGET_DATA;
+ break;
+ case OMP_TEAMS:
+ ort = ORT_TEAMS;
+ break;
+ default:
+ gcc_unreachable ();
+ }
+ gimplify_scan_omp_clauses (&OMP_CLAUSES (expr), pre_p, ort);
gimplify_and_add (OMP_BODY (expr), &body);
gimplify_adjust_omp_clauses (&OMP_CLAUSES (expr));
- if (TREE_CODE (expr) == OMP_SECTIONS)
- stmt = gimple_build_omp_sections (body, OMP_CLAUSES (expr));
- else if (TREE_CODE (expr) == OMP_SINGLE)
- stmt = gimple_build_omp_single (body, OMP_CLAUSES (expr));
- else
- gcc_unreachable ();
+ switch (TREE_CODE (expr))
+ {
+ case OMP_SECTIONS:
+ stmt = gimple_build_omp_sections (body, OMP_CLAUSES (expr));
+ break;
+ case OMP_SINGLE:
+ stmt = gimple_build_omp_single (body, OMP_CLAUSES (expr));
+ break;
+ case OMP_TARGET:
+ stmt = gimple_build_omp_target (body, GF_OMP_TARGET_KIND_REGION,
+ OMP_CLAUSES (expr));
+ break;
+ case OMP_TARGET_DATA:
+ stmt = gimple_build_omp_target (body, GF_OMP_TARGET_KIND_DATA,
+ OMP_CLAUSES (expr));
+ break;
+ case OMP_TEAMS:
+ stmt = gimple_build_omp_teams (body, OMP_CLAUSES (expr));
+ break;
+ default:
+ gcc_unreachable ();
+ }
+
+ gimplify_seq_add_stmt (pre_p, stmt);
+ *expr_p = NULL_TREE;
+}
+
+/* Gimplify the gross structure of OpenMP target update construct. */
+
+static void
+gimplify_omp_target_update (tree *expr_p, gimple_seq *pre_p)
+{
+ tree expr = *expr_p;
+ gimple stmt;
+
+ gimplify_scan_omp_clauses (&OMP_TARGET_UPDATE_CLAUSES (expr), pre_p,
+ ORT_WORKSHARE);
+ gimplify_adjust_omp_clauses (&OMP_TARGET_UPDATE_CLAUSES (expr));
+ stmt = gimple_build_omp_target (NULL, GF_OMP_TARGET_KIND_UPDATE,
+ OMP_TARGET_UPDATE_CLAUSES (expr));
gimplify_seq_add_stmt (pre_p, stmt);
+ *expr_p = NULL_TREE;
}
/* A subroutine of gimplify_omp_atomic. The front end is supposed to have
@@ -7811,10 +8002,18 @@ gimplify_expr (tree *expr_p, gimple_seq
case OMP_SECTIONS:
case OMP_SINGLE:
+ case OMP_TARGET:
+ case OMP_TARGET_DATA:
+ case OMP_TEAMS:
gimplify_omp_workshare (expr_p, pre_p);
ret = GS_ALL_DONE;
break;
+ case OMP_TARGET_UPDATE:
+ gimplify_omp_target_update (expr_p, pre_p);
+ ret = GS_ALL_DONE;
+ break;
+
case OMP_SECTION:
case OMP_MASTER:
case OMP_ORDERED:
Jakub
