OpenACC 2.7: Implement self clause for compute constructs

This patch implements the 'self' clause for compute constructs: parallel,
kernels, and serial. This clause conditionally uses the local device
(the host mult-core CPU) as the executing device of the compute region.

The actual implementation of the "local device" device type inside libgomp
(presumably using pthreads) is still not yet completed, so the libgomp
side is still implemented the exact same as host-fallback mode. (so as of now,
it essentially behaves like the 'if' clause with the condition inverted)

gcc/c/ChangeLog:

	* c-parser.cc (c_parser_oacc_compute_clause_self): New function.
	(c_parser_oacc_all_clauses): Add new 'bool compute_p = false'
	parameter, add parsing of self clause when compute_p is true.
	(OACC_KERNELS_CLAUSE_MASK): Add PRAGMA_OACC_CLAUSE_SELF.
	(OACC_PARALLEL_CLAUSE_MASK): Likewise,
	(OACC_SERIAL_CLAUSE_MASK): Likewise.
	(c_parser_oacc_compute): Adjust call to c_parser_oacc_all_clauses to
	set compute_p argument to true.
	* c-typeck.cc (c_finish_omp_clauses): Add OMP_CLAUSE_SELF case.

gcc/cp/ChangeLog:

	* parser.cc (cp_parser_oacc_compute_clause_self): New function.
	(cp_parser_oacc_all_clauses): Add new 'bool compute_p = false'
	parameter, add parsing of self clause when compute_p is true.
	(OACC_KERNELS_CLAUSE_MASK): Add PRAGMA_OACC_CLAUSE_SELF.
	(OACC_PARALLEL_CLAUSE_MASK): Likewise,
	(OACC_SERIAL_CLAUSE_MASK): Likewise.
	(cp_parser_oacc_compute): Adjust call to c_parser_oacc_all_clauses to
	set compute_p argument to true.
	* pt.cc (tsubst_omp_clauses): Add OMP_CLAUSE_SELF case.
	* semantics.cc (c_finish_omp_clauses): Add OMP_CLAUSE_SELF case, merged
	with OMP_CLAUSE_IF case.

gcc/fortran/ChangeLog:

	* gfortran.h (typedef struct gfc_omp_clauses): Add self_expr field.
	* openmp.cc (enum omp_mask2): Add OMP_CLAUSE_SELF.
	(gfc_match_omp_clauses): Add handling for OMP_CLAUSE_SELF.
	(OACC_PARALLEL_CLAUSES): Add OMP_CLAUSE_SELF.
	(OACC_KERNELS_CLAUSES): Likewise.
	(OACC_SERIAL_CLAUSES): Likewise.
	(resolve_omp_clauses): Add handling for omp_clauses->self_expr.
	* trans-openmp.cc (gfc_trans_omp_clauses): Add handling of
	clauses->self_expr and building of OMP_CLAUSE_SELF tree clause.
	(gfc_split_omp_clauses): Add handling of self_expr field copy.

gcc/ChangeLog:

	* gimplify.cc (gimplify_scan_omp_clauses): Add OMP_CLAUSE_SELF case.
	(gimplify_adjust_omp_clauses): Likewise.
	* omp-expand.cc (expand_omp_target): Add OMP_CLAUSE_SELF expansion code,
	* omp-low.cc (scan_sharing_clauses): Add OMP_CLAUSE_SELF case.
	* tree-core.h (enum omp_clause_code): Add OMP_CLAUSE_SELF enum.
	* tree-nested.cc (convert_nonlocal_omp_clauses): Add OMP_CLAUSE_SELF
	case.
	(convert_local_omp_clauses): Likewise.
	* tree-pretty-print.cc (dump_omp_clause): Add OMP_CLAUSE_SELF case.
	* tree.cc (omp_clause_num_ops): Add OMP_CLAUSE_SELF entry.
	(omp_clause_code_name): Likewise.
	* tree.h (OMP_CLAUSE_SELF_EXPR): New macro.

gcc/testsuite/ChangeLog:

	* c-c++-common/goacc/self-clause-1.c: New test.
	* c-c++-common/goacc/self-clause-2.c: New test.
	* gfortran.dg/goacc/self.f95: New test.

include/ChangeLog:

	* gomp-constants.h (GOACC_FLAG_LOCAL_DEVICE): New flag bit value.

libgomp/ChangeLog:

	* oacc-parallel.c (GOACC_parallel_keyed): Add code to handle
	GOACC_FLAG_LOCAL_DEVICE case.
	* testsuite/libgomp.oacc-c-c++-common/self-1.c: New test.
This commit is contained in:
Chung-Lin Tang 2023-06-13 08:44:31 -07:00 committed by Thomas Schwinge
parent fa68e04e76
commit 3a3596389c
22 changed files with 1306 additions and 12 deletions

View file

@ -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));

View file

@ -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:

View file

@ -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);

View file

@ -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:

View file

@ -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:

View file

@ -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;

View file

@ -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;

View file

@ -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;
}

View file

@ -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:

View file

@ -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);

View file

@ -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:

View file

@ -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" } */
;
}

View file

@ -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;
}

View file

@ -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

View file

@ -527,6 +527,9 @@ enum omp_clause_code {
/* OpenACC clause: nohost. */
OMP_CLAUSE_NOHOST,
/* OpenACC clause: self. */
OMP_CLAUSE_SELF,
};
#undef DEFTREESTRUCT

View file

@ -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:

View file

@ -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 ();
}

View file

@ -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

View file

@ -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)

View file

@ -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. */

View file

@ -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);

View file

@ -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;
}