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 )
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;
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)
{
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);
| (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) )
| (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) )
| (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
}
}
- 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));
continue;
case OMP_CLAUSE_IF:
+ case OMP_CLAUSE_SELF:
case OMP_CLAUSE_NUM_THREADS:
case OMP_CLAUSE_NUM_TEAMS:
case OMP_CLAUSE_THREAD_LIMIT:
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;
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)
{
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);
| (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) )
| (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) )
| (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
}
}
- 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);
/* 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:
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:
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;
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
};
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],
| 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 \
}
}
+ 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;
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])
{
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;
/* 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;
}
}
/* Fall through. */
+ case OMP_CLAUSE_SELF:
case OMP_CLAUSE_FINAL:
OMP_CLAUSE_OPERAND (c, 0)
= gimple_boolify (OMP_CLAUSE_OPERAND (c, 0));
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:
}
}
+ 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);
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:
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:
--- /dev/null
+/* { 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" } */
+ ;
+}
--- /dev/null
+/* { 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;
+}
--- /dev/null
+! { 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
/* OpenACC clause: nohost. */
OMP_CLAUSE_NOHOST,
+
+ /* OpenACC clause: self. */
+ OMP_CLAUSE_SELF,
};
#undef DEFTREESTRUCT
/* 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:
/* 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:
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 ();
}
0, /* OMP_CLAUSE_IF_PRESENT */
0, /* OMP_CLAUSE_FINALIZE */
0, /* OMP_CLAUSE_NOHOST */
+ 1, /* OMP_CLAUSE_SELF */
};
const char * const omp_clause_code_name[] =
"if_present",
"finalize",
"nohost",
+ "self",
};
/* Unless specific to OpenACC, we tend to internally maintain OpenMP-centric
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)
/* 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. */
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);
--- /dev/null
+#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;
+}
projects / thirdparty / gcc.git / commitdiff
