diff --git a/gcc/c/c-parser.cc b/gcc/c/c-parser.cc
index 0d468b86bd8c7754c9f9310eca92daea5e090176..a82f5afeff7c7024185124d0274bb0a9e86959d2 100644
--- a/gcc/c/c-parser.cc
+++ b/gcc/c/c-parser.cc
@@ -15923,6 +15923,41 @@ c_parser_oacc_clause_wait (c_parser *parser, tree list)
return list;
}
+/* OpenACC 2.7:
+ self [( expression )] */
+
+static tree
+c_parser_oacc_compute_clause_self (c_parser *parser, tree list)
+{
+ tree t;
+ location_t location = c_parser_peek_token (parser)->location;
+ if (c_parser_peek_token (parser)->type == CPP_OPEN_PAREN)
+ {
+ matching_parens parens;
+ parens.consume_open (parser);
+
+ location_t loc = c_parser_peek_token (parser)->location;
+ c_expr expr = c_parser_expr_no_commas (parser, NULL);
+ expr = convert_lvalue_to_rvalue (loc, expr, true, true);
+ t = c_objc_common_truthvalue_conversion (loc, expr.value);
+ t = c_fully_fold (t, false, NULL);
+ parens.skip_until_found_close (parser);
+ }
+ else
+ t = truthvalue_true_node;
+
+ for (tree c = list; c; c = OMP_CLAUSE_CHAIN (c))
+ if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_SELF)
+ {
+ error_at (location, "too many %<self%> clauses");
+ return list;
+ }
+
+ tree c = build_omp_clause (location, OMP_CLAUSE_SELF);
+ OMP_CLAUSE_SELF_EXPR (c) = t;
+ OMP_CLAUSE_CHAIN (c) = list;
+ return c;
+}
/* OpenMP 5.0:
order ( concurrent )
@@ -18048,7 +18083,8 @@ c_parser_omp_clause_detach (c_parser *parser, tree list)
static tree
c_parser_oacc_all_clauses (c_parser *parser, omp_clause_mask mask,
- const char *where, bool finish_p = true)
+ const char *where, bool finish_p = true,
+ bool compute_p = false)
{
tree clauses = NULL;
bool first = true;
@@ -18064,7 +18100,18 @@ c_parser_oacc_all_clauses (c_parser *parser, omp_clause_mask mask,
c_parser_consume_token (parser);
here = c_parser_peek_token (parser)->location;
- c_kind = c_parser_omp_clause_name (parser);
+
+ /* For OpenACC compute directives */
+ if (compute_p
+ && c_parser_next_token_is (parser, CPP_NAME)
+ && !strcmp (IDENTIFIER_POINTER (c_parser_peek_token (parser)->value),
+ "self"))
+ {
+ c_kind = PRAGMA_OACC_CLAUSE_SELF;
+ c_parser_consume_token (parser);
+ }
+ else
+ c_kind = c_parser_omp_clause_name (parser);
switch (c_kind)
{
@@ -18196,6 +18243,10 @@ c_parser_oacc_all_clauses (c_parser *parser, omp_clause_mask mask,
false, clauses);
c_name = "reduction";
break;
+ case PRAGMA_OACC_CLAUSE_SELF:
+ clauses = c_parser_oacc_compute_clause_self (parser, clauses);
+ c_name = "self";
+ break;
case PRAGMA_OACC_CLAUSE_SEQ:
clauses = c_parser_oacc_simple_clause (here, OMP_CLAUSE_SEQ,
clauses);
@@ -19032,6 +19083,7 @@ c_parser_oacc_loop (location_t loc, c_parser *parser, char *p_name,
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_NUM_GANGS) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_NUM_WORKERS) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT) \
+ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_SELF) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_VECTOR_LENGTH) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_WAIT) )
@@ -19052,6 +19104,7 @@ c_parser_oacc_loop (location_t loc, c_parser *parser, char *p_name,
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_NUM_WORKERS) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_REDUCTION) \
+ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_SELF) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_VECTOR_LENGTH) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_WAIT) )
@@ -19070,6 +19123,7 @@ c_parser_oacc_loop (location_t loc, c_parser *parser, char *p_name,
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_FIRSTPRIVATE) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_REDUCTION) \
+ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_SELF) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_WAIT) )
static tree
@@ -19112,7 +19166,7 @@ c_parser_oacc_compute (location_t loc, c_parser *parser,
}
}
- tree clauses = c_parser_oacc_all_clauses (parser, mask, p_name);
+ tree clauses = c_parser_oacc_all_clauses (parser, mask, p_name, true, true);
tree block = c_begin_omp_parallel ();
add_stmt (c_parser_omp_structured_block (parser, if_p));
diff --git a/gcc/c/c-typeck.cc b/gcc/c/c-typeck.cc
index 112d28fd656d05bfb8a0af437ac7d6d2fd231a7f..1eee653b3b8b7b49b29ef8d7f74718cebb697ddb 100644
--- a/gcc/c/c-typeck.cc
+++ b/gcc/c/c-typeck.cc
@@ -15845,6 +15845,7 @@ c_finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
continue;
case OMP_CLAUSE_IF:
+ case OMP_CLAUSE_SELF:
case OMP_CLAUSE_NUM_THREADS:
case OMP_CLAUSE_NUM_TEAMS:
case OMP_CLAUSE_THREAD_LIMIT:
diff --git a/gcc/cp/parser.cc b/gcc/cp/parser.cc
index 5483121f51c16537e2982fb87dc903da718f037a..c5a9928ad27eb50f596c5a76f4890c8d20a7238c 100644
--- a/gcc/cp/parser.cc
+++ b/gcc/cp/parser.cc
@@ -41192,13 +41192,51 @@ cp_parser_oacc_clause_async (cp_parser *parser, tree list)
return list;
}
+/* OpenACC 2.7:
+ self [( expression )] */
+
+static tree
+cp_parser_oacc_compute_clause_self (cp_parser *parser, tree list)
+{
+ tree t;
+ location_t location = cp_lexer_peek_token (parser->lexer)->location;
+ if (cp_lexer_peek_token (parser->lexer)->type == CPP_OPEN_PAREN)
+ {
+ matching_parens parens;
+ parens.consume_open (parser);
+ t = cp_parser_assignment_expression (parser);
+ if (t == error_mark_node
+ || !parens.require_close (parser))
+ {
+ cp_parser_skip_to_closing_parenthesis (parser, /*recovering=*/true,
+ /*or_comma=*/false,
+ /*consume_paren=*/true);
+ return list;
+ }
+ }
+ else
+ t = truthvalue_true_node;
+
+ for (tree c = list; c; c = OMP_CLAUSE_CHAIN (c))
+ if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_SELF)
+ {
+ error_at (location, "too many %<self%> clauses");
+ return list;
+ }
+
+ tree c = build_omp_clause (location, OMP_CLAUSE_SELF);
+ OMP_CLAUSE_SELF_EXPR (c) = t;
+ OMP_CLAUSE_CHAIN (c) = list;
+ return c;
+}
+
/* Parse all OpenACC clauses. The set clauses allowed by the directive
is a bitmask in MASK. Return the list of clauses found. */
static tree
cp_parser_oacc_all_clauses (cp_parser *parser, omp_clause_mask mask,
const char *where, cp_token *pragma_tok,
- bool finish_p = true)
+ bool finish_p = true, bool compute_p = false)
{
tree clauses = NULL;
bool first = true;
@@ -41218,7 +41256,19 @@ cp_parser_oacc_all_clauses (cp_parser *parser, omp_clause_mask mask,
cp_lexer_consume_token (parser->lexer);
here = cp_lexer_peek_token (parser->lexer)->location;
- c_kind = cp_parser_omp_clause_name (parser);
+
+ /* For OpenACC compute directives */
+ if (compute_p
+ && cp_lexer_next_token_is (parser->lexer, CPP_NAME)
+ && !strcmp (IDENTIFIER_POINTER
+ (cp_lexer_peek_token (parser->lexer)->u.value),
+ "self"))
+ {
+ c_kind = PRAGMA_OACC_CLAUSE_SELF;
+ cp_lexer_consume_token (parser->lexer);
+ }
+ else
+ c_kind = cp_parser_omp_clause_name (parser);
switch (c_kind)
{
@@ -41352,6 +41402,10 @@ cp_parser_oacc_all_clauses (cp_parser *parser, omp_clause_mask mask,
false, clauses);
c_name = "reduction";
break;
+ case PRAGMA_OACC_CLAUSE_SELF:
+ clauses = cp_parser_oacc_compute_clause_self (parser, clauses);
+ c_name = "self";
+ break;
case PRAGMA_OACC_CLAUSE_SEQ:
clauses = cp_parser_oacc_simple_clause (here, OMP_CLAUSE_SEQ,
clauses);
@@ -46866,6 +46920,7 @@ cp_parser_oacc_loop (cp_parser *parser, cp_token *pragma_tok, char *p_name,
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_NUM_GANGS) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_NUM_WORKERS) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT) \
+ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_SELF) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_VECTOR_LENGTH) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_WAIT) )
@@ -46886,6 +46941,7 @@ cp_parser_oacc_loop (cp_parser *parser, cp_token *pragma_tok, char *p_name,
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRIVATE) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_REDUCTION) \
+ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_SELF) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_VECTOR_LENGTH) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_WAIT) )
@@ -46904,6 +46960,7 @@ cp_parser_oacc_loop (cp_parser *parser, cp_token *pragma_tok, char *p_name,
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_FIRSTPRIVATE) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_REDUCTION) \
+ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_SELF) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_WAIT) )
static tree
@@ -46949,7 +47006,8 @@ cp_parser_oacc_compute (cp_parser *parser, cp_token *pragma_tok,
}
}
- tree clauses = cp_parser_oacc_all_clauses (parser, mask, p_name, pragma_tok);
+ tree clauses = cp_parser_oacc_all_clauses (parser, mask, p_name, pragma_tok,
+ true, true);
tree block = begin_omp_parallel ();
unsigned int save = cp_parser_begin_omp_structured_block (parser);
diff --git a/gcc/cp/pt.cc b/gcc/cp/pt.cc
index 1c1c93131180c41bfb97b87f15c8413bb271787a..86c95b278ba397a02aabfe182e3a2f92f4e44df0 100644
--- a/gcc/cp/pt.cc
+++ b/gcc/cp/pt.cc
@@ -17418,6 +17418,7 @@ tsubst_omp_clauses (tree clauses, enum c_omp_region_type ort,
/* FALLTHRU */
case OMP_CLAUSE_TILE:
case OMP_CLAUSE_IF:
+ case OMP_CLAUSE_SELF:
case OMP_CLAUSE_NUM_THREADS:
case OMP_CLAUSE_SCHEDULE:
case OMP_CLAUSE_COLLAPSE:
diff --git a/gcc/cp/semantics.cc b/gcc/cp/semantics.cc
index dc3c11461fbc2f0c1e6d01b451ebc76b2546c6f7..72ec72de690b9a6170bd1a44a7417e009413bee2 100644
--- a/gcc/cp/semantics.cc
+++ b/gcc/cp/semantics.cc
@@ -7377,13 +7377,14 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
goto handle_field_decl;
case OMP_CLAUSE_IF:
- t = OMP_CLAUSE_IF_EXPR (c);
+ case OMP_CLAUSE_SELF:
+ t = OMP_CLAUSE_OPERAND (c, 0);
t = maybe_convert_cond (t);
if (t == error_mark_node)
remove = true;
else if (!processing_template_decl)
t = fold_build_cleanup_point_expr (TREE_TYPE (t), t);
- OMP_CLAUSE_IF_EXPR (c) = t;
+ OMP_CLAUSE_OPERAND (c, 0) = t;
break;
case OMP_CLAUSE_FINAL:
diff --git a/gcc/fortran/gfortran.h b/gcc/fortran/gfortran.h
index 9c1a39a19de4086cc795a9d06392e513245e586f..ceb2873c03f6e4952b4c5a69011db64046536499 100644
--- a/gcc/fortran/gfortran.h
+++ b/gcc/fortran/gfortran.h
@@ -1546,6 +1546,7 @@ typedef struct gfc_omp_clauses
gfc_omp_namelist *lists[OMP_LIST_NUM];
struct gfc_expr *if_expr;
struct gfc_expr *if_exprs[OMP_IF_LAST];
+ struct gfc_expr *self_expr;
struct gfc_expr *final_expr;
struct gfc_expr *num_threads;
struct gfc_expr *chunk_size;
diff --git a/gcc/fortran/openmp.cc b/gcc/fortran/openmp.cc
index caa5c9e15a30c7cbebbb3ba67e4bcbd4f47c2307..083c15e5599da49570d33a161cc9d02172278dfd 100644
--- a/gcc/fortran/openmp.cc
+++ b/gcc/fortran/openmp.cc
@@ -1094,6 +1094,7 @@ enum omp_mask2
OMP_CLAUSE_DOACROSS, /* OpenMP 5.2 */
OMP_CLAUSE_ASSUMPTIONS, /* OpenMP 5.1. */
OMP_CLAUSE_USES_ALLOCATORS, /* OpenMP 5.0 */
+ OMP_CLAUSE_SELF, /* OpenACC 2.7 */
/* This must come last. */
OMP_MASK2_LAST
};
@@ -3519,6 +3520,27 @@ gfc_match_omp_clauses (gfc_omp_clauses **cp, const omp_mask mask,
else
gfc_current_locus = old_loc;
}
+ if ((mask & OMP_CLAUSE_SELF)
+ && (m = gfc_match_dupl_check (!c->self_expr, "self"))
+ != MATCH_NO)
+ {
+ gcc_assert (!(mask & OMP_CLAUSE_HOST_SELF));
+ if (m == MATCH_ERROR)
+ goto error;
+ m = gfc_match (" ( %e )", &c->self_expr);
+ if (m == MATCH_ERROR)
+ {
+ gfc_current_locus = old_loc;
+ break;
+ }
+ else if (m == MATCH_NO)
+ {
+ c->self_expr = gfc_get_logical_expr (gfc_default_logical_kind,
+ NULL, true);
+ needs_space = true;
+ }
+ continue;
+ }
if ((mask & OMP_CLAUSE_HOST_SELF)
&& gfc_match ("self ( ") == MATCH_YES
&& gfc_match_omp_map_clause (&c->lists[OMP_LIST_MAP],
@@ -3791,19 +3813,22 @@ error:
| OMP_CLAUSE_COPY | OMP_CLAUSE_COPYIN | OMP_CLAUSE_COPYOUT \
| OMP_CLAUSE_CREATE | OMP_CLAUSE_NO_CREATE | OMP_CLAUSE_PRESENT \
| OMP_CLAUSE_DEVICEPTR | OMP_CLAUSE_PRIVATE | OMP_CLAUSE_FIRSTPRIVATE \
- | OMP_CLAUSE_DEFAULT | OMP_CLAUSE_WAIT | OMP_CLAUSE_ATTACH)
+ | OMP_CLAUSE_DEFAULT | OMP_CLAUSE_WAIT | OMP_CLAUSE_ATTACH \
+ | OMP_CLAUSE_SELF)
#define OACC_KERNELS_CLAUSES \
(omp_mask (OMP_CLAUSE_IF) | OMP_CLAUSE_ASYNC | OMP_CLAUSE_NUM_GANGS \
| OMP_CLAUSE_NUM_WORKERS | OMP_CLAUSE_VECTOR_LENGTH | OMP_CLAUSE_DEVICEPTR \
| OMP_CLAUSE_COPY | OMP_CLAUSE_COPYIN | OMP_CLAUSE_COPYOUT \
| OMP_CLAUSE_CREATE | OMP_CLAUSE_NO_CREATE | OMP_CLAUSE_PRESENT \
- | OMP_CLAUSE_DEFAULT | OMP_CLAUSE_WAIT | OMP_CLAUSE_ATTACH)
+ | OMP_CLAUSE_DEFAULT | OMP_CLAUSE_WAIT | OMP_CLAUSE_ATTACH \
+ | OMP_CLAUSE_SELF)
#define OACC_SERIAL_CLAUSES \
(omp_mask (OMP_CLAUSE_IF) | OMP_CLAUSE_ASYNC | OMP_CLAUSE_REDUCTION \
| OMP_CLAUSE_COPY | OMP_CLAUSE_COPYIN | OMP_CLAUSE_COPYOUT \
| OMP_CLAUSE_CREATE | OMP_CLAUSE_NO_CREATE | OMP_CLAUSE_PRESENT \
| OMP_CLAUSE_DEVICEPTR | OMP_CLAUSE_PRIVATE | OMP_CLAUSE_FIRSTPRIVATE \
- | OMP_CLAUSE_DEFAULT | OMP_CLAUSE_WAIT | OMP_CLAUSE_ATTACH)
+ | OMP_CLAUSE_DEFAULT | OMP_CLAUSE_WAIT | OMP_CLAUSE_ATTACH \
+ | OMP_CLAUSE_SELF)
#define OACC_DATA_CLAUSES \
(omp_mask (OMP_CLAUSE_IF) | OMP_CLAUSE_DEVICEPTR | OMP_CLAUSE_COPY \
| OMP_CLAUSE_COPYIN | OMP_CLAUSE_COPYOUT | OMP_CLAUSE_CREATE \
@@ -7540,6 +7565,15 @@ resolve_omp_clauses (gfc_code *code, gfc_omp_clauses *omp_clauses,
}
}
+ if (omp_clauses->self_expr)
+ {
+ gfc_expr *expr = omp_clauses->self_expr;
+ if (!gfc_resolve_expr (expr)
+ || expr->ts.type != BT_LOGICAL || expr->rank != 0)
+ gfc_error ("SELF clause at %L requires a scalar LOGICAL expression",
+ &expr->where);
+ }
+
if (omp_clauses->final_expr)
{
gfc_expr *expr = omp_clauses->final_expr;
diff --git a/gcc/fortran/trans-openmp.cc b/gcc/fortran/trans-openmp.cc
index 7930f2fd5d12e9e8368ef45f385d7635fb9fb53a..82bbc41b388683140475fa7a2379b979145ed547 100644
--- a/gcc/fortran/trans-openmp.cc
+++ b/gcc/fortran/trans-openmp.cc
@@ -3966,6 +3966,7 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses,
OMP_CLAUSE_IF_EXPR (c) = if_var;
omp_clauses = gfc_trans_add_clause (c, omp_clauses);
}
+
for (ifc = 0; ifc < OMP_IF_LAST; ifc++)
if (clauses->if_exprs[ifc])
{
@@ -4017,6 +4018,21 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses,
omp_clauses = gfc_trans_add_clause (c, omp_clauses);
}
+ if (clauses->self_expr)
+ {
+ tree self_var;
+
+ gfc_init_se (&se, NULL);
+ gfc_conv_expr (&se, clauses->self_expr);
+ gfc_add_block_to_block (block, &se.pre);
+ self_var = gfc_evaluate_now (se.expr, block);
+ gfc_add_block_to_block (block, &se.post);
+
+ c = build_omp_clause (gfc_get_location (&where), OMP_CLAUSE_SELF);
+ OMP_CLAUSE_SELF_EXPR (c) = self_var;
+ omp_clauses = gfc_trans_add_clause (c, omp_clauses);
+ }
+
if (clauses->final_expr)
{
tree final_var;
@@ -6615,6 +6631,8 @@ gfc_split_omp_clauses (gfc_code *code,
/* And this is copied to all. */
clausesa[GFC_OMP_SPLIT_TARGET].if_expr
= code->ext.omp_clauses->if_expr;
+ clausesa[GFC_OMP_SPLIT_TARGET].self_expr
+ = code->ext.omp_clauses->self_expr;
clausesa[GFC_OMP_SPLIT_TARGET].nowait
= code->ext.omp_clauses->nowait;
}
diff --git a/gcc/gimplify.cc b/gcc/gimplify.cc
index 22ff1075abbf1c0b02e1a19cd75162c7b494d89a..77f07af6ec5c170a24ae54df8229b5b93512d240 100644
--- a/gcc/gimplify.cc
+++ b/gcc/gimplify.cc
@@ -12121,6 +12121,7 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
}
/* Fall through. */
+ case OMP_CLAUSE_SELF:
case OMP_CLAUSE_FINAL:
OMP_CLAUSE_OPERAND (c, 0)
= gimple_boolify (OMP_CLAUSE_OPERAND (c, 0));
@@ -13342,6 +13343,7 @@ gimplify_adjust_omp_clauses (gimple_seq *pre_p, gimple_seq body, tree *list_p,
case OMP_CLAUSE_COPYIN:
case OMP_CLAUSE_COPYPRIVATE:
case OMP_CLAUSE_IF:
+ case OMP_CLAUSE_SELF:
case OMP_CLAUSE_NUM_THREADS:
case OMP_CLAUSE_NUM_TEAMS:
case OMP_CLAUSE_THREAD_LIMIT:
diff --git a/gcc/omp-expand.cc b/gcc/omp-expand.cc
index 6172839c719375239b79676091b25335a6d443f7..8576b9381025fb30414593d745cf22f8f1cf41ab 100644
--- a/gcc/omp-expand.cc
+++ b/gcc/omp-expand.cc
@@ -10332,6 +10332,47 @@ expand_omp_target (struct omp_region *region)
}
}
+ if ((c = omp_find_clause (clauses, OMP_CLAUSE_SELF)) != NULL_TREE)
+ {
+ gcc_assert (is_gimple_omp_oacc (entry_stmt) && offloaded);
+
+ edge e = split_block_after_labels (new_bb);
+ basic_block cond_bb = e->src;
+ new_bb = e->dest;
+ remove_edge (e);
+
+ basic_block then_bb = create_empty_bb (cond_bb);
+ basic_block else_bb = create_empty_bb (then_bb);
+ set_immediate_dominator (CDI_DOMINATORS, then_bb, cond_bb);
+ set_immediate_dominator (CDI_DOMINATORS, else_bb, cond_bb);
+
+ tree self_cond = gimple_boolify (OMP_CLAUSE_SELF_EXPR (c));
+ stmt = gimple_build_cond_empty (self_cond);
+ gsi = gsi_last_bb (cond_bb);
+ gsi_insert_after (&gsi, stmt, GSI_CONTINUE_LINKING);
+
+ tree tmp_var = create_tmp_var (TREE_TYPE (goacc_flags));
+ stmt = gimple_build_assign (tmp_var, BIT_IOR_EXPR, goacc_flags,
+ build_int_cst (integer_type_node,
+ GOACC_FLAG_LOCAL_DEVICE));
+ gsi = gsi_start_bb (then_bb);
+ gsi_insert_after (&gsi, stmt, GSI_CONTINUE_LINKING);
+
+ gsi = gsi_start_bb (else_bb);
+ stmt = gimple_build_assign (tmp_var, goacc_flags);
+ gsi_insert_after (&gsi, stmt, GSI_CONTINUE_LINKING);
+
+ make_edge (cond_bb, then_bb, EDGE_TRUE_VALUE);
+ make_edge (cond_bb, else_bb, EDGE_FALSE_VALUE);
+ add_bb_to_loop (then_bb, cond_bb->loop_father);
+ add_bb_to_loop (else_bb, cond_bb->loop_father);
+ make_edge (then_bb, new_bb, EDGE_FALLTHRU);
+ make_edge (else_bb, new_bb, EDGE_FALLTHRU);
+
+ goacc_flags = tmp_var;
+ gsi = gsi_last_nondebug_bb (new_bb);
+ }
+
if (need_device_adjustment)
{
tree uns = fold_convert (unsigned_type_node, device);
diff --git a/gcc/omp-low.cc b/gcc/omp-low.cc
index 91ef74f1f6a44bc689caaac4809409b98d394762..161bcfeec057a1dd26099ac457e565113500f79a 100644
--- a/gcc/omp-low.cc
+++ b/gcc/omp-low.cc
@@ -1493,6 +1493,7 @@ scan_sharing_clauses (tree clauses, omp_context *ctx)
case OMP_CLAUSE_FINAL:
case OMP_CLAUSE_IF:
+ case OMP_CLAUSE_SELF:
case OMP_CLAUSE_NUM_THREADS:
case OMP_CLAUSE_NUM_TEAMS:
case OMP_CLAUSE_THREAD_LIMIT:
@@ -1920,6 +1921,7 @@ scan_sharing_clauses (tree clauses, omp_context *ctx)
case OMP_CLAUSE_COPYIN:
case OMP_CLAUSE_DEFAULT:
case OMP_CLAUSE_IF:
+ case OMP_CLAUSE_SELF:
case OMP_CLAUSE_NUM_THREADS:
case OMP_CLAUSE_NUM_TEAMS:
case OMP_CLAUSE_THREAD_LIMIT:
diff --git a/gcc/testsuite/c-c++-common/goacc/self-clause-1.c b/gcc/testsuite/c-c++-common/goacc/self-clause-1.c
new file mode 100644
index 0000000000000000000000000000000000000000..ed5d072e81f17d051c2589d4842ba9d3f0912e59
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/goacc/self-clause-1.c
@@ -0,0 +1,22 @@
+/* { dg-skip-if "not yet" { c++ } } */
+
+void
+f (int b)
+{
+ struct { int i; } *p;
+
+#pragma acc parallel self self(b) /* { dg-error "too many 'self' clauses" } */
+ ;
+#pragma acc parallel self(*p) /* { dg-error "used struct type value where scalar is required" } */
+ ;
+
+#pragma acc kernels self self(b) /* { dg-error "too many 'self' clauses" } */
+ ;
+#pragma acc kernels self(*p) /* { dg-error "used struct type value where scalar is required" } */
+ ;
+
+#pragma acc serial self self(b) /* { dg-error "too many 'self' clauses" } */
+ ;
+#pragma acc serial self(*p) /* { dg-error "used struct type value where scalar is required" } */
+ ;
+}
diff --git a/gcc/testsuite/c-c++-common/goacc/self-clause-2.c b/gcc/testsuite/c-c++-common/goacc/self-clause-2.c
new file mode 100644
index 0000000000000000000000000000000000000000..d932ac9a4a608ae807faeb7b389989dcc258a581
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/goacc/self-clause-2.c
@@ -0,0 +1,17 @@
+/* { dg-additional-options "-fdump-tree-gimple" } */
+
+void
+f (short c)
+{
+#pragma acc parallel self(c) copy(c)
+ /* { dg-final { scan-tree-dump-times {(?n)#pragma omp target oacc_parallel map\(tofrom:c \[len: [0-9]+\]\) self\(_[0-9]+\)$} 1 "gimple" } } */
+ ++c;
+
+#pragma acc kernels self(c) copy(c)
+ /* { dg-final { scan-tree-dump-times {(?n)#pragma omp target oacc_kernels map\(tofrom:c \[len: [0-9]+\]\) self\(_[0-9]+\)$} 1 "gimple" } } */
+ ++c;
+
+#pragma acc serial self(c) copy(c)
+ /* { dg-final { scan-tree-dump-times {(?n)#pragma omp target oacc_serial map\(tofrom:c \[len: [0-9]+\]\) self\(_[0-9]+\)$} 1 "gimple" } } */
+ ++c;
+}
diff --git a/gcc/testsuite/gfortran.dg/goacc/self.f95 b/gcc/testsuite/gfortran.dg/goacc/self.f95
new file mode 100644
index 0000000000000000000000000000000000000000..4817f16be5610cf72a821d7e69578e75252d29fd
--- /dev/null
+++ b/gcc/testsuite/gfortran.dg/goacc/self.f95
@@ -0,0 +1,53 @@
+! { dg-do compile }
+
+program test
+ implicit none
+
+ logical :: x
+ integer :: i
+
+ !$acc parallel self () ! { dg-error "Invalid character" }
+ !$acc parallel self (i) ! { dg-error "scalar LOGICAL expression" }
+ !$acc end parallel
+ !$acc parallel self (1) ! { dg-error "scalar LOGICAL expression" }
+ !$acc end parallel
+
+ !$acc kernels self () ! { dg-error "Invalid character" }
+ !$acc kernels self (i) ! { dg-error "scalar LOGICAL expression" }
+ !$acc end kernels
+ !$acc kernels self (1) ! { dg-error "scalar LOGICAL expression" }
+ !$acc end kernels
+
+ !$acc serial self () ! { dg-error "Invalid character" }
+ !$acc serial self (i) ! { dg-error "scalar LOGICAL expression" }
+ !$acc end serial
+ !$acc serial self (1) ! { dg-error "scalar LOGICAL expression" }
+ !$acc end serial
+
+ ! at most one self clause may appear
+ !$acc parallel self (.false.) self (.false.) { dg-error "Duplicated 'self' clause" }
+ !$acc kernels self (.false.) self (.false.) { dg-error "Duplicated 'self' clause" }
+ !$acc serial self (.false.) self (.false.) { dg-error "Duplicated 'self' clause" }
+
+ !$acc parallel self (x)
+ !$acc end parallel
+ !$acc parallel self (.true.)
+ !$acc end parallel
+ !$acc parallel self (i.gt.1)
+ !$acc end parallel
+
+ !$acc kernels self (x)
+ !$acc end kernels
+ !$acc kernels self (.true.)
+ !$acc end kernels
+ !$acc kernels self (i.gt.1)
+ !$acc end kernels
+
+ !$acc serial self (x)
+ !$acc end serial
+ !$acc serial self (.true.)
+ !$acc end serial
+ !$acc serial self (i.gt.1)
+ !$acc end serial
+
+end program test
diff --git a/gcc/tree-core.h b/gcc/tree-core.h
index 2c89b655691b1ea01a9bf1ee177c336262f3d52b..cfe37c1d6273999ca4967de4af16eee435595286 100644
--- a/gcc/tree-core.h
+++ b/gcc/tree-core.h
@@ -527,6 +527,9 @@ enum omp_clause_code {
/* OpenACC clause: nohost. */
OMP_CLAUSE_NOHOST,
+
+ /* OpenACC clause: self. */
+ OMP_CLAUSE_SELF,
};
#undef DEFTREESTRUCT
diff --git a/gcc/tree-nested.cc b/gcc/tree-nested.cc
index d2fe3fca8af9277369bd11aae678de7750acbe67..51a7a032bc37936829ba3b63e99b54cf0e9c8b98 100644
--- a/gcc/tree-nested.cc
+++ b/gcc/tree-nested.cc
@@ -1374,6 +1374,7 @@ convert_nonlocal_omp_clauses (tree *pclauses, struct walk_stmt_info *wi)
/* FALLTHRU */
case OMP_CLAUSE_FINAL:
case OMP_CLAUSE_IF:
+ case OMP_CLAUSE_SELF:
case OMP_CLAUSE_NUM_THREADS:
case OMP_CLAUSE_DEPEND:
case OMP_CLAUSE_DOACROSS:
@@ -2165,6 +2166,7 @@ convert_local_omp_clauses (tree *pclauses, struct walk_stmt_info *wi)
/* FALLTHRU */
case OMP_CLAUSE_FINAL:
case OMP_CLAUSE_IF:
+ case OMP_CLAUSE_SELF:
case OMP_CLAUSE_NUM_THREADS:
case OMP_CLAUSE_DEPEND:
case OMP_CLAUSE_DOACROSS:
diff --git a/gcc/tree-pretty-print.cc b/gcc/tree-pretty-print.cc
index 58705c57ad5332d51eb4e5c67f53e5637532b19b..39ec1df93947c580209c7f4653f494783a7a388c 100644
--- a/gcc/tree-pretty-print.cc
+++ b/gcc/tree-pretty-print.cc
@@ -1453,7 +1453,12 @@ dump_omp_clause (pretty_printer *pp, tree clause, int spc, dump_flags_t flags)
false);
pp_right_paren (pp);
break;
-
+ case OMP_CLAUSE_SELF:
+ pp_string (pp, "self(");
+ dump_generic_node (pp, OMP_CLAUSE_SELF_EXPR (clause),
+ spc, flags, false);
+ pp_right_paren (pp);
+ break;
default:
gcc_unreachable ();
}
diff --git a/gcc/tree.cc b/gcc/tree.cc
index f9fa7b78ffffe1064d42ac37547a304f9cb8da7c..c38b09c431b8723f379f3ef1968ef528653dc8de 100644
--- a/gcc/tree.cc
+++ b/gcc/tree.cc
@@ -326,6 +326,7 @@ unsigned const char omp_clause_num_ops[] =
0, /* OMP_CLAUSE_IF_PRESENT */
0, /* OMP_CLAUSE_FINALIZE */
0, /* OMP_CLAUSE_NOHOST */
+ 1, /* OMP_CLAUSE_SELF */
};
const char * const omp_clause_code_name[] =
@@ -417,6 +418,7 @@ const char * const omp_clause_code_name[] =
"if_present",
"finalize",
"nohost",
+ "self",
};
/* Unless specific to OpenACC, we tend to internally maintain OpenMP-centric
diff --git a/gcc/tree.h b/gcc/tree.h
index 781297c7e770fef5d7d5be4c453e23af980e8f64..aaf744c060e578931b84b9c104c2fe23a7e7f4ed 100644
--- a/gcc/tree.h
+++ b/gcc/tree.h
@@ -1734,6 +1734,8 @@ class auto_suppress_location_wrappers
OMP_CLAUSE_OPERAND (OMP_CLAUSE_SUBCODE_CHECK (NODE, OMP_CLAUSE_HINT), 0)
#define OMP_CLAUSE_FILTER_EXPR(NODE) \
OMP_CLAUSE_OPERAND (OMP_CLAUSE_SUBCODE_CHECK (NODE, OMP_CLAUSE_FILTER), 0)
+#define OMP_CLAUSE_SELF_EXPR(NODE) \
+ OMP_CLAUSE_OPERAND (OMP_CLAUSE_SUBCODE_CHECK (NODE, OMP_CLAUSE_SELF), 0)
#define OMP_CLAUSE_GRAINSIZE_EXPR(NODE) \
OMP_CLAUSE_OPERAND (OMP_CLAUSE_SUBCODE_CHECK (NODE, OMP_CLAUSE_GRAINSIZE),0)
diff --git a/include/gomp-constants.h b/include/gomp-constants.h
index 8d4e8e813031884c9cdf285050db55a24db0e9dc..89b966e63c6e9b3ee106f02fb030e58e6eb82ad0 100644
--- a/include/gomp-constants.h
+++ b/include/gomp-constants.h
@@ -304,6 +304,8 @@ enum gomp_map_kind
/* Force host fallback execution. */
#define GOACC_FLAG_HOST_FALLBACK (1 << 0)
+/* Execute on local device (i.e. host multicore CPU). */
+#define GOACC_FLAG_LOCAL_DEVICE (1 << 1)
/* For legacy reasons, in the ABI, the GOACC_FLAGs are encoded as an inverted
bitmask. */
diff --git a/libgomp/oacc-parallel.c b/libgomp/oacc-parallel.c
index 363e6656982584bf269921be2c0f957182aa8ac9..cf37a1bdd7dbe0c0580397887774621c7d4012cf 100644
--- a/libgomp/oacc-parallel.c
+++ b/libgomp/oacc-parallel.c
@@ -193,6 +193,17 @@ GOACC_parallel_keyed (int flags_m, void (*fn) (void *),
goacc_restore_bind ();
goto out_prof;
}
+ else if (flags & GOACC_FLAG_LOCAL_DEVICE)
+ {
+ /* TODO: a proper pthreads based "multi-core CPU" local device
+ implementation. Currently, this is still the same as host-fallback. */
+ prof_info.device_type = acc_device_host;
+ api_info.device_type = prof_info.device_type;
+ goacc_save_and_set_bind (acc_device_host);
+ fn (hostaddrs);
+ goacc_restore_bind ();
+ goto out_prof;
+ }
else if (acc_device_type (acc_dev->type) == acc_device_host)
{
fn (hostaddrs);
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/self-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/self-1.c
new file mode 100644
index 0000000000000000000000000000000000000000..752e16e85455575898b58ad31cda9f25fe077980
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/self-1.c
@@ -0,0 +1,962 @@
+#include <openacc.h>
+#include <stdlib.h>
+#include <stdbool.h>
+
+#define N 32
+
+int
+main(int argc, char **argv)
+{
+ float *a, *b, *d_a, *d_b, exp, exp2;
+ int i;
+ const int one = 1;
+ const int zero = 0;
+ int n;
+
+ a = (float *) malloc (N * sizeof (float));
+ b = (float *) malloc (N * sizeof (float));
+ d_a = (float *) acc_malloc (N * sizeof (float));
+ d_b = (float *) acc_malloc (N * sizeof (float));
+
+ for (i = 0; i < N; i++)
+ a[i] = 4.0;
+
+#pragma acc parallel copyin(a[0:N]) copyout(b[0:N]) self(0)
+ {
+ int ii;
+
+ for (ii = 0; ii < N; ii++)
+ {
+ if (acc_on_device (acc_device_host))
+ b[ii] = a[ii] + 1;
+ else
+ b[ii] = a[ii];
+ }
+ }
+
+#if ACC_MEM_SHARED
+ exp = 5.0;
+#else
+ exp = 4.0;
+#endif
+
+ for (i = 0; i < N; i++)
+ {
+ if (b[i] != exp)
+ abort();
+ }
+
+ for (i = 0; i < N; i++)
+ a[i] = 16.0;
+
+#pragma acc parallel self(1)
+ {
+ int ii;
+
+ for (ii = 0; ii < N; ii++)
+ {
+ if (acc_on_device (acc_device_host))
+ b[ii] = a[ii] + 1;
+ else
+ b[ii] = a[ii];
+ }
+ }
+
+ for (i = 0; i < N; i++)
+ {
+ if (b[i] != 17.0)
+ abort();
+ }
+
+ for (i = 0; i < N; i++)
+ a[i] = 8.0;
+
+#pragma acc parallel copyin(a[0:N]) copyout(b[0:N]) self(!one)
+ {
+ int ii;
+
+ for (ii = 0; ii < N; ii++)
+ {
+ if (acc_on_device (acc_device_host))
+ b[ii] = a[ii] + 1;
+ else
+ b[ii] = a[ii];
+ }
+ }
+
+#if ACC_MEM_SHARED
+ exp = 9.0;
+#else
+ exp = 8.0;
+#endif
+
+ for (i = 0; i < N; i++)
+ {
+ if (b[i] != exp)
+ abort();
+ }
+
+ for (i = 0; i < N; i++)
+ a[i] = 22.0;
+
+#pragma acc parallel self(!zero)
+ {
+ int ii;
+
+ for (ii = 0; ii < N; ii++)
+ {
+ if (acc_on_device (acc_device_host))
+ b[ii] = a[ii] + 1;
+ else
+ b[ii] = a[ii];
+ }
+ }
+
+ for (i = 0; i < N; i++)
+ {
+ if (b[i] != 23.0)
+ abort();
+ }
+
+ for (i = 0; i < N; i++)
+ a[i] = 16.0;
+
+#pragma acc parallel copyin(a[0:N]) copyout(b[0:N]) self(false)
+ {
+ int ii;
+
+ for (ii = 0; ii < N; ii++)
+ {
+ if (acc_on_device (acc_device_host))
+ b[ii] = a[ii] + 1;
+ else
+ b[ii] = a[ii];
+ }
+ }
+
+#if ACC_MEM_SHARED
+ exp = 17.0;
+#else
+ exp = 16.0;
+#endif
+
+ for (i = 0; i < N; i++)
+ {
+ if (b[i] != exp)
+ abort();
+ }
+
+ for (i = 0; i < N; i++)
+ a[i] = 76.0;
+
+#pragma acc parallel self(true)
+ {
+ int ii;
+
+ for (ii = 0; ii < N; ii++)
+ {
+ if (acc_on_device (acc_device_host))
+ b[ii] = a[ii] + 1;
+ else
+ b[ii] = a[ii];
+ }
+ }
+
+ for (i = 0; i < N; i++)
+ {
+ if (b[i] != 77.0)
+ abort();
+ }
+
+ for (i = 0; i < N; i++)
+ a[i] = 22.0;
+
+ n = 1;
+
+#pragma acc parallel copyin(a[0:N]) copyout(b[0:N]) self(!n)
+ {
+ int ii;
+
+ for (ii = 0; ii < N; ii++)
+ {
+ if (acc_on_device (acc_device_host))
+ b[ii] = a[ii] + 1;
+ else
+ b[ii] = a[ii];
+ }
+ }
+
+#if ACC_MEM_SHARED
+ exp = 23.0;
+#else
+ exp = 22.0;
+#endif
+
+ for (i = 0; i < N; i++)
+ {
+ if (b[i] != exp)
+ abort();
+ }
+
+ for (i = 0; i < N; i++)
+ a[i] = 18.0;
+
+ n = 0;
+
+#pragma acc parallel self(!n)
+ {
+ int ii;
+
+ for (ii = 0; ii < N; ii++)
+ {
+ if (acc_on_device (acc_device_host))
+ b[ii] = a[ii] + 1;
+ else
+ b[ii] = a[ii];
+ }
+ }
+
+ for (i = 0; i < N; i++)
+ {
+ if (b[i] != 19.0)
+ abort();
+ }
+
+ for (i = 0; i < N; i++)
+ a[i] = 49.0;
+
+ n = 1;
+
+#pragma acc parallel copyin(a[0:N]) copyout(b[0:N]) self(!(n + n))
+ {
+ int ii;
+
+ for (ii = 0; ii < N; ii++)
+ {
+ if (acc_on_device (acc_device_host))
+ b[ii] = a[ii] + 1;
+ else
+ b[ii] = a[ii];
+ }
+ }
+
+#if ACC_MEM_SHARED
+ exp = 50.0;
+#else
+ exp = 49.0;
+#endif
+
+ for (i = 0; i < N; i++)
+ {
+ if (b[i] != exp)
+ abort();
+ }
+
+ for (i = 0; i < N; i++)
+ a[i] = 38.0;
+
+ n = 0;
+
+#pragma acc parallel self(!(n + n))
+ {
+ int ii;
+
+ for (ii = 0; ii < N; ii++)
+ {
+ if (acc_on_device (acc_device_host))
+ b[ii] = a[ii] + 1;
+ else
+ b[ii] = a[ii];
+ }
+ }
+
+ for (i = 0; i < N; i++)
+ {
+ if (b[i] != 39.0)
+ abort();
+ }
+
+ for (i = 0; i < N; i++)
+ a[i] = 91.0;
+
+#pragma acc parallel copyin(a[0:N]) copyout(b[0:N]) self(!(-2))
+ {
+ int ii;
+
+ for (ii = 0; ii < N; ii++)
+ {
+ if (acc_on_device (acc_device_host))
+ b[ii] = a[ii] + 1;
+ else
+ b[ii] = a[ii];
+ }
+ }
+
+#if ACC_MEM_SHARED
+ exp = 92.0;
+#else
+ exp = 91.0;
+#endif
+
+ for (i = 0; i < N; i++)
+ {
+ if (b[i] != exp)
+ abort();
+ }
+
+ for (i = 0; i < N; i++)
+ a[i] = 43.0;
+
+#pragma acc parallel copyin(a[0:N]) copyout(b[0:N]) self(one != 1)
+ {
+ int ii;
+
+ for (ii = 0; ii < N; ii++)
+ {
+ if (acc_on_device (acc_device_host))
+ b[ii] = a[ii] + 1;
+ else
+ b[ii] = a[ii];
+ }
+ }
+
+#if ACC_MEM_SHARED
+ exp = 44.0;
+#else
+ exp = 43.0;
+#endif
+
+ for (i = 0; i < N; i++)
+ {
+ if (b[i] != exp)
+ abort();
+ }
+
+ for (i = 0; i < N; i++)
+ a[i] = 87.0;
+
+#pragma acc parallel self(one != 0)
+ {
+ int ii;
+
+ for (ii = 0; ii < N; ii++)
+ {
+ if (acc_on_device (acc_device_host))
+ b[ii] = a[ii] + 1;
+ else
+ b[ii] = a[ii];
+ }
+ }
+
+ for (i = 0; i < N; i++)
+ {
+ if (b[i] != 88.0)
+ abort();
+ }
+
+ for (i = 0; i < N; i++)
+ {
+ a[i] = 3.0;
+ b[i] = 9.0;
+ }
+
+#if ACC_MEM_SHARED
+ exp = 0.0;
+ exp2 = 0.0;
+#else
+ acc_map_data (a, d_a, N * sizeof (float));
+ acc_map_data (b, d_b, N * sizeof (float));
+ exp = 3.0;
+ exp2 = 9.0;
+#endif
+
+#pragma acc update device(a[0:N], b[0:N]) if(1)
+
+ for (i = 0; i < N; i++)
+ {
+ a[i] = 0.0;
+ b[i] = 0.0;
+ }
+
+#pragma acc update host(a[0:N], b[0:N]) if(1)
+
+ for (i = 0; i < N; i++)
+ {
+ if (a[i] != exp)
+ abort();
+
+ if (b[i] != exp2)
+ abort();
+ }
+
+ for (i = 0; i < N; i++)
+ {
+ a[i] = 6.0;
+ b[i] = 12.0;
+ }
+
+#pragma acc update device(a[0:N], b[0:N]) if(0)
+
+ for (i = 0; i < N; i++)
+ {
+ a[i] = 0.0;
+ b[i] = 0.0;
+ }
+
+#pragma acc update host(a[0:N], b[0:N]) if(1)
+
+ for (i = 0; i < N; i++)
+ {
+ if (a[i] != exp)
+ abort();
+
+ if (b[i] != exp2)
+ abort();
+ }
+
+ for (i = 0; i < N; i++)
+ {
+ a[i] = 26.0;
+ b[i] = 21.0;
+ }
+
+#pragma acc update device(a[0:N], b[0:N]) if(1)
+
+ for (i = 0; i < N; i++)
+ {
+ a[i] = 0.0;
+ b[i] = 0.0;
+ }
+
+#pragma acc update host(a[0:N], b[0:N]) if(0)
+
+ for (i = 0; i < N; i++)
+ {
+ if (a[i] != 0.0)
+ abort();
+
+ if (b[i] != 0.0)
+ abort();
+ }
+
+#if !ACC_MEM_SHARED
+ acc_unmap_data (a);
+ acc_unmap_data (b);
+#endif
+
+ acc_free (d_a);
+ acc_free (d_b);
+
+ for (i = 0; i < N; i++)
+ {
+ a[i] = 4.0;
+ b[i] = 0.0;
+ }
+
+#pragma acc data copyin(a[0:N]) copyout(b[0:N]) if(1)
+{
+#pragma acc parallel present(a[0:N])
+ {
+ int ii;
+
+ for (ii = 0; ii < N; ii++)
+ {
+ b[ii] = a[ii];
+ }
+ }
+}
+
+ for (i = 0; i < N; i++)
+ {
+ if (b[i] != 4.0)
+ abort();
+ }
+
+ for (i = 0; i < N; i++)
+ {
+ a[i] = 8.0;
+ b[i] = 1.0;
+ }
+
+#pragma acc data copyin(a[0:N]) copyout(b[0:N]) if(0)
+{
+#if !ACC_MEM_SHARED
+ if (acc_is_present (a, N * sizeof (float)))
+ abort ();
+#endif
+
+#if !ACC_MEM_SHARED
+ if (acc_is_present (b, N * sizeof (float)))
+ abort ();
+#endif
+}
+
+ for (i = 0; i < N; i++)
+ {
+ a[i] = 18.0;
+ b[i] = 21.0;
+ }
+
+#pragma acc data copyin(a[0:N]) if(1)
+{
+#if !ACC_MEM_SHARED
+ if (!acc_is_present (a, N * sizeof (float)))
+ abort ();
+#endif
+
+#pragma acc data copyout(b[0:N]) if(0)
+ {
+#if !ACC_MEM_SHARED
+ if (acc_is_present (b, N * sizeof (float)))
+ abort ();
+#endif
+
+#pragma acc data copyout(b[0:N]) if(1)
+ {
+#pragma acc parallel present(a[0:N]) present(b[0:N])
+ {
+ int ii;
+
+ for (ii = 0; ii < N; ii++)
+ {
+ b[ii] = a[ii];
+ }
+ }
+ }
+
+#if !ACC_MEM_SHARED
+ if (acc_is_present (b, N * sizeof (float)))
+ abort ();
+#endif
+ }
+}
+
+ for (i = 0; i < N; i++)
+ {
+ if (b[i] != 18.0)
+ abort ();
+ }
+
+#pragma acc enter data copyin (b[0:N]) if (0)
+
+#if !ACC_MEM_SHARED
+ if (acc_is_present (b, N * sizeof (float)))
+ abort ();
+#endif
+
+#pragma acc exit data delete (b[0:N]) if (0)
+
+#pragma acc enter data copyin (b[0:N]) if (1)
+
+#if !ACC_MEM_SHARED
+ if (!acc_is_present (b, N * sizeof (float)))
+ abort ();
+#endif
+
+#pragma acc exit data delete (b[0:N]) if (1)
+
+#if !ACC_MEM_SHARED
+ if (acc_is_present (b, N * sizeof (float)))
+ abort ();
+#endif
+
+#pragma acc enter data copyin (b[0:N]) if (zero)
+
+#if !ACC_MEM_SHARED
+ if (acc_is_present (b, N * sizeof (float)))
+ abort ();
+#endif
+
+#pragma acc exit data delete (b[0:N]) if (zero)
+
+#pragma acc enter data copyin (b[0:N]) if (one)
+
+#if !ACC_MEM_SHARED
+ if (!acc_is_present (b, N * sizeof (float)))
+ abort ();
+#endif
+
+#pragma acc exit data delete (b[0:N]) if (one)
+
+#if !ACC_MEM_SHARED
+ if (acc_is_present (b, N * sizeof (float)))
+ abort ();
+#endif
+
+#pragma acc enter data copyin (b[0:N]) if (one == 0)
+
+#if !ACC_MEM_SHARED
+ if (acc_is_present (b, N * sizeof (float)))
+ abort ();
+#endif
+
+#pragma acc exit data delete (b[0:N]) if (one == 0)
+
+#pragma acc enter data copyin (b[0:N]) if (one == 1)
+
+#if !ACC_MEM_SHARED
+ if (!acc_is_present (b, N * sizeof (float)))
+ abort ();
+#endif
+
+#pragma acc exit data delete (b[0:N]) if (one == 1)
+
+#if !ACC_MEM_SHARED
+ if (acc_is_present (b, N * sizeof (float)))
+ abort ();
+#endif
+
+ for (i = 0; i < N; i++)
+ a[i] = 4.0;
+
+#pragma acc kernels copyin(a[0:N]) copyout(b[0:N]) self(0)
+ {
+ int ii;
+
+ for (ii = 0; ii < N; ii++)
+ {
+ if (acc_on_device (acc_device_host))
+ b[ii] = a[ii] + 1;
+ else
+ b[ii] = a[ii];
+ }
+ }
+
+#if ACC_MEM_SHARED
+ exp = 5.0;
+#else
+ exp = 4.0;
+#endif
+
+ for (i = 0; i < N; i++)
+ {
+ if (b[i] != exp)
+ abort();
+ }
+
+ for (i = 0; i < N; i++)
+ a[i] = 16.0;
+
+#pragma acc kernels self(1)
+ {
+ int ii;
+
+ for (ii = 0; ii < N; ii++)
+ {
+ if (acc_on_device (acc_device_host))
+ b[ii] = a[ii] + 1;
+ else
+ b[ii] = a[ii];
+ }
+ }
+
+ for (i = 0; i < N; i++)
+ {
+ if (b[i] != 17.0)
+ abort();
+ }
+
+ for (i = 0; i < N; i++)
+ a[i] = 8.0;
+
+#pragma acc kernels copyin(a[0:N]) copyout(b[0:N]) self(!one)
+ {
+ int ii;
+
+ for (ii = 0; ii < N; ii++)
+ {
+ if (acc_on_device (acc_device_host))
+ b[ii] = a[ii] + 1;
+ else
+ b[ii] = a[ii];
+ }
+ }
+
+#if ACC_MEM_SHARED
+ exp = 9.0;
+#else
+ exp = 8.0;
+#endif
+
+ for (i = 0; i < N; i++)
+ {
+ if (b[i] != exp)
+ abort();
+ }
+
+ for (i = 0; i < N; i++)
+ a[i] = 22.0;
+
+#pragma acc kernels self(!zero)
+ {
+ int ii;
+
+ for (ii = 0; ii < N; ii++)
+ {
+ if (acc_on_device (acc_device_host))
+ b[ii] = a[ii] + 1;
+ else
+ b[ii] = a[ii];
+ }
+ }
+
+ for (i = 0; i < N; i++)
+ {
+ if (b[i] != 23.0)
+ abort();
+ }
+
+ for (i = 0; i < N; i++)
+ a[i] = 16.0;
+
+#pragma acc kernels copyin(a[0:N]) copyout(b[0:N]) self(false)
+ {
+ int ii;
+
+ for (ii = 0; ii < N; ii++)
+ {
+ if (acc_on_device (acc_device_host))
+ b[ii] = a[ii] + 1;
+ else
+ b[ii] = a[ii];
+ }
+ }
+
+#if ACC_MEM_SHARED
+ exp = 17.0;
+#else
+ exp = 16.0;
+#endif
+
+ for (i = 0; i < N; i++)
+ {
+ if (b[i] != exp)
+ abort();
+ }
+
+ for (i = 0; i < N; i++)
+ a[i] = 76.0;
+
+#pragma acc kernels self(true)
+ {
+ int ii;
+
+ for (ii = 0; ii < N; ii++)
+ {
+ if (acc_on_device (acc_device_host))
+ b[ii] = a[ii] + 1;
+ else
+ b[ii] = a[ii];
+ }
+ }
+
+ for (i = 0; i < N; i++)
+ {
+ if (b[i] != 77.0)
+ abort();
+ }
+
+ for (i = 0; i < N; i++)
+ a[i] = 22.0;
+
+ n = 1;
+
+#pragma acc kernels copyin(a[0:N]) copyout(b[0:N]) self(!n)
+ {
+ int ii;
+
+ for (ii = 0; ii < N; ii++)
+ {
+ if (acc_on_device (acc_device_host))
+ b[ii] = a[ii] + 1;
+ else
+ b[ii] = a[ii];
+ }
+ }
+
+#if ACC_MEM_SHARED
+ exp = 23.0;
+#else
+ exp = 22.0;
+#endif
+
+ for (i = 0; i < N; i++)
+ {
+ if (b[i] != exp)
+ abort();
+ }
+
+ for (i = 0; i < N; i++)
+ a[i] = 18.0;
+
+ n = 0;
+
+#pragma acc kernels self(!n)
+ {
+ int ii;
+
+ for (ii = 0; ii < N; ii++)
+ {
+ if (acc_on_device (acc_device_host))
+ b[ii] = a[ii] + 1;
+ else
+ b[ii] = a[ii];
+ }
+ }
+
+ for (i = 0; i < N; i++)
+ {
+ if (b[i] != 19.0)
+ abort();
+ }
+
+ for (i = 0; i < N; i++)
+ a[i] = 49.0;
+
+ n = 1;
+
+#pragma acc kernels copyin(a[0:N]) copyout(b[0:N]) self((n + n) == 0)
+ {
+ int ii;
+
+ for (ii = 0; ii < N; ii++)
+ {
+ if (acc_on_device (acc_device_host))
+ b[ii] = a[ii] + 1;
+ else
+ b[ii] = a[ii];
+ }
+ }
+
+#if ACC_MEM_SHARED
+ exp = 50.0;
+#else
+ exp = 49.0;
+#endif
+
+ for (i = 0; i < N; i++)
+ {
+ if (b[i] != exp)
+ abort();
+ }
+
+ for (i = 0; i < N; i++)
+ a[i] = 38.0;
+
+ n = 0;
+
+#pragma acc kernels self(!(n + n))
+ {
+ int ii;
+
+ for (ii = 0; ii < N; ii++)
+ {
+ if (acc_on_device (acc_device_host))
+ b[ii] = a[ii] + 1;
+ else
+ b[ii] = a[ii];
+ }
+ }
+
+ for (i = 0; i < N; i++)
+ {
+ if (b[i] != 39.0)
+ abort();
+ }
+
+ for (i = 0; i < N; i++)
+ a[i] = 91.0;
+
+#pragma acc kernels copyin(a[0:N]) copyout(b[0:N]) self(!(-2))
+ {
+ int ii;
+
+ for (ii = 0; ii < N; ii++)
+ {
+ if (acc_on_device (acc_device_host))
+ b[ii] = a[ii] + 1;
+ else
+ b[ii] = a[ii];
+ }
+ }
+
+#if ACC_MEM_SHARED
+ exp = 92.0;
+#else
+ exp = 91.0;
+#endif
+
+ for (i = 0; i < N; i++)
+ {
+ if (b[i] != exp)
+ abort();
+ }
+
+ for (i = 0; i < N; i++)
+ a[i] = 43.0;
+
+#pragma acc kernels copyin(a[0:N]) copyout(b[0:N]) self(one != 1)
+ {
+ int ii;
+
+ for (ii = 0; ii < N; ii++)
+ {
+ if (acc_on_device (acc_device_host))
+ b[ii] = a[ii] + 1;
+ else
+ b[ii] = a[ii];
+ }
+ }
+
+#if ACC_MEM_SHARED
+ exp = 44.0;
+#else
+ exp = 43.0;
+#endif
+
+ for (i = 0; i < N; i++)
+ {
+ if (b[i] != exp)
+ abort();
+ }
+
+ for (i = 0; i < N; i++)
+ a[i] = 87.0;
+
+#pragma acc kernels self(one != 0)
+ {
+ int ii;
+
+ for (ii = 0; ii < N; ii++)
+ {
+ if (acc_on_device (acc_device_host))
+ b[ii] = a[ii] + 1;
+ else
+ b[ii] = a[ii];
+ }
+ }
+
+ for (i = 0; i < N; i++)
+ {
+ if (b[i] != 88.0)
+ abort();
+ }
+
+ for (i = 0; i < N; i++)
+ {
+ a[i] = 3.0;
+ b[i] = 9.0;
+ }
+
+#if ACC_MEM_SHARED
+ exp = 0.0;
+ exp2 = 0.0;
+#else
+ acc_map_data (a, d_a, N * sizeof (float));
+ acc_map_data (b, d_b, N * sizeof (float));
+ exp = 3.0;
+ exp2 = 9.0;
+#endif
+
+ return 0;
+}