2016-02-01 Cesar Philippidis <cesar@codesourcery.com>
gcc/c/
* c-typeck.c (c_finish_omp_clauses): Report OMP_CLAUSE_MAP errors
as data clause errors for OpenACC.
gcc/cp/
* cp-tree.h (finish_omp_clauses): Update prototype.
* parser.c (cp_parser_oacc_all_clauses): Update call to
finish_omp_clauses.
(cp_parser_omp_all_clauses): Likewise.
(cp_parser_omp_for_loop): Likewise.
(cp_omp_split_clauses): Likewise.
(cp_parser_oacc_cache): Likewise.
(cp_parser_oacc_loop): Likewise.
(cp_parser_omp_declare_target): Likewise.
(cp_parser_cilk_for): Likewise.
* pt.c (tsubst_attribute): Update call to tsubst_omp_clauses.
(tsubst_omp_clauses): New is_oacc argument. Use it when calling
finish_omp_clauses.
(tsubst_omp_for_iterator): Update call to finish_omp_clauses.
(tsubst_expr): Update calls to tsubst_omp_clauses.
* semantics.c (finish_omp_clauses): Add is_oacc argument. Report
OMP_CLAUSE_MAP errors as data clause errors for OpenACC. Check for
duplicate reference mappings. Exclude "omp declare simd"-isms when
processing OpenACC clauses.
(finish_omp_for): Update call to finish_omp_clauses.
gcc/
* gimplify.c (gimplify_scan_omp_clauses): Consider OACC_{DATA, PARALLEL,
KERNELS} when setting target_firstprivatize_array_bases. Consider
OACC_{DATA, ENTER_DATA, EXIT_DATA, UPDATE} when filtering out pointer
mappings. Also filter out GOMP_MAP_POINTER.
gcc/testsuite/
* c-c++-common/goacc/data-clause-duplicate-1.c: Adjust test.
* c-c++-common/goacc/declare-2.c: Likewise.
* c-c++-common/goacc/deviceptr-1.c: Likewise.
* c-c++-common/goacc/kernels-alias-ipa-pta-3.c: Likewise.
* c-c++-common/goacc/pcopy.c: Likewise.
* c-c++-common/goacc/pcopyin.c: Likewise.
* c-c++-common/goacc/pcopyout.c: Likewise.
* c-c++-common/goacc/pcreate.c: Likewise.
* c-c++-common/goacc/present-1.c: Likewise.
* g++.dg/goacc/data-1.C: New test.
* g++.dg/goacc/data-2.C: New test.
* g++.dg/goacc/reduction-1.C: New test.
* g++.dg/goacc/update.C: New test.
* g++.dg/gomp/template-data.C: New test.
libgomp/
* testsuite/libgomp.c++/non-scalar-data.C: New test.
* testsuite/libgomp.oacc-c++/data-references.C: New test.
* testsuite/libgomp.oacc-c++/data-templates.C: New test.
* testsuite/libgomp.oacc-c++/non-scalar-data-templates.C: New test.
* testsuite/libgomp.oacc-c++/update-reference.C: New test.
* testsuite/libgomp.oacc-c++/update-template.C: New test.
* testsuite/libgomp.oacc-c-c++-common/asyncwait-1.c: Adjust test.
* testsuite/libgomp.oacc-c-c++-common/data-2.c: Likewise.
* testsuite/libgomp.oacc-c-c++-common/data-3.c: Likewise.
* testsuite/libgomp.oacc-c-c++-common/if-1.c: Likewise.
* testsuite/libgomp.oacc-c-c++-common/nested-1.c: Likewise.
* testsuite/libgomp.oacc-c-c++-common/present-2.c: Likewise.
@@ -13157,8 +13157,10 @@ c_finish_omp_clauses (tree clauses, bool is_omp, bool declare_simd)
{
if (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_MAP)
error ("%qD appears more than once in motion clauses", t);
- else
+ else if (is_omp)
error ("%qD appears more than once in map clauses", t);
+ else
+ error ("%qD appears more than once in data clauses", t);
remove = true;
}
else if (bitmap_bit_p (&generic_head, DECL_UID (t))
@@ -6415,7 +6415,7 @@ extern tree omp_reduction_id (enum tree_code, tree, tree);
extern tree cp_remove_omp_priv_cleanup_stmt (tree *, int *, void *);
extern void cp_check_omp_declare_reduction (tree);
extern void finish_omp_declare_simd_methods (tree);
-extern tree finish_omp_clauses (tree, bool, bool = false);
+extern tree finish_omp_clauses (tree, bool, bool, bool = false);
extern tree push_omp_privatization_clauses (bool);
extern void pop_omp_privatization_clauses (tree);
extern void save_omp_privatization_clauses (vec<tree> &);
@@ -32168,7 +32168,7 @@ cp_parser_oacc_all_clauses (cp_parser *parser, omp_clause_mask mask,
cp_parser_skip_to_pragma_eol (parser, pragma_tok);
if (finish_p)
- return finish_omp_clauses (clauses, false);
+ return finish_omp_clauses (clauses, true, true);
return clauses;
}
@@ -32487,9 +32487,9 @@ cp_parser_omp_all_clauses (cp_parser *parser, omp_clause_mask mask,
if (finish_p)
{
if ((mask & (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_UNIFORM)) != 0)
- return finish_omp_clauses (clauses, false, true);
+ return finish_omp_clauses (clauses, false, false, true);
else
- return finish_omp_clauses (clauses, true);
+ return finish_omp_clauses (clauses, false, true);
}
return clauses;
}
@@ -33564,7 +33564,7 @@ cp_parser_omp_for_loop (cp_parser *parser, enum tree_code code, tree clauses,
else
c = build_omp_clause (loc, OMP_CLAUSE_LASTPRIVATE);
OMP_CLAUSE_DECL (c) = add_private_clause;
- c = finish_omp_clauses (c, true);
+ c = finish_omp_clauses (c, false, true);
if (c)
{
OMP_CLAUSE_CHAIN (c) = clauses;
@@ -33713,7 +33713,7 @@ cp_omp_split_clauses (location_t loc, enum tree_code code,
c_omp_split_clauses (loc, code, mask, clauses, cclauses);
for (i = 0; i < C_OMP_CLAUSE_SPLIT_COUNT; i++)
if (cclauses[i])
- cclauses[i] = finish_omp_clauses (cclauses[i], true);
+ cclauses[i] = finish_omp_clauses (cclauses[i], false, true);
}
/* OpenMP 4.0:
@@ -34985,7 +34985,7 @@ cp_parser_oacc_cache (cp_parser *parser, cp_token *pragma_tok)
tree stmt, clauses;
clauses = cp_parser_omp_var_list (parser, OMP_CLAUSE__CACHE_, NULL_TREE);
- clauses = finish_omp_clauses (clauses, false);
+ clauses = finish_omp_clauses (clauses, true, true);
cp_parser_require_pragma_eol (parser, cp_lexer_peek_token (parser->lexer));
@@ -35310,9 +35310,9 @@ cp_parser_oacc_loop (cp_parser *parser, cp_token *pragma_tok, char *p_name,
{
clauses = c_oacc_split_loop_clauses (clauses, cclauses);
if (*cclauses)
- finish_omp_clauses (*cclauses, false);
+ finish_omp_clauses (*cclauses, true, true);
if (clauses)
- finish_omp_clauses (clauses, false);
+ finish_omp_clauses (clauses, true, true);
}
tree block = begin_omp_structured_block ();
@@ -35678,7 +35678,7 @@ cp_parser_omp_declare_target (cp_parser *parser, cp_token *pragma_tok)
{
clauses = cp_parser_omp_var_list (parser, OMP_CLAUSE_TO_DECLARE,
clauses);
- clauses = finish_omp_clauses (clauses, true);
+ clauses = finish_omp_clauses (clauses, false, true);
cp_parser_require_pragma_eol (parser, pragma_tok);
}
else
@@ -37652,7 +37652,7 @@ cp_parser_cilk_for (cp_parser *parser, tree grain)
tree clauses = build_omp_clause (EXPR_LOCATION (grain), OMP_CLAUSE_SCHEDULE);
OMP_CLAUSE_SCHEDULE_KIND (clauses) = OMP_CLAUSE_SCHEDULE_CILKFOR;
OMP_CLAUSE_SCHEDULE_CHUNK_EXPR (clauses) = grain;
- clauses = finish_omp_clauses (clauses, false);
+ clauses = finish_omp_clauses (clauses, false, false);
tree ret = cp_parser_omp_for_loop (parser, CILK_FOR, clauses, NULL);
if (ret)
@@ -9541,7 +9541,8 @@ can_complete_type_without_circularity (tree type)
return 1;
}
-static tree tsubst_omp_clauses (tree, bool, bool, tree, tsubst_flags_t, tree);
+static tree tsubst_omp_clauses (tree, bool, bool, tree, tsubst_flags_t, tree,
+ bool);
/* Instantiate a single dependent attribute T (a TREE_LIST), and return either
T or a new TREE_LIST, possibly a chain in the case of a pack expansion. */
@@ -9561,9 +9562,9 @@ tsubst_attribute (tree t, tree *decl_p, tree args,
{
tree clauses = TREE_VALUE (val);
clauses = tsubst_omp_clauses (clauses, true, false, args,
- complain, in_decl);
+ complain, in_decl, false);
c_omp_declare_simd_clauses_to_decls (*decl_p, clauses);
- clauses = finish_omp_clauses (clauses, false, true);
+ clauses = finish_omp_clauses (clauses, false, false, true);
tree parms = DECL_ARGUMENTS (*decl_p);
clauses
= c_omp_declare_simd_clauses_to_numbers (parms, clauses);
@@ -14457,7 +14458,8 @@ tsubst_omp_clause_decl (tree decl, tree args, tsubst_flags_t complain,
static tree
tsubst_omp_clauses (tree clauses, bool declare_simd, bool allow_fields,
- tree args, tsubst_flags_t complain, tree in_decl)
+ tree args, tsubst_flags_t complain, tree in_decl,
+ bool is_oacc)
{
tree new_clauses = NULL_TREE, nc, oc;
tree linear_no_step = NULL_TREE;
@@ -14670,7 +14672,7 @@ tsubst_omp_clauses (tree clauses, bool declare_simd, bool allow_fields,
new_clauses = nreverse (new_clauses);
if (!declare_simd)
{
- new_clauses = finish_omp_clauses (new_clauses, allow_fields);
+ new_clauses = finish_omp_clauses (new_clauses, is_oacc, allow_fields);
if (linear_no_step)
for (nc = new_clauses; nc; nc = OMP_CLAUSE_CHAIN (nc))
if (nc == linear_no_step)
@@ -14891,7 +14893,7 @@ tsubst_omp_for_iterator (tree t, int i, tree declv, tree orig_declv,
{
tree c = build_omp_clause (input_location, OMP_CLAUSE_PRIVATE);
OMP_CLAUSE_DECL (c) = decl;
- c = finish_omp_clauses (c, true);
+ c = finish_omp_clauses (c, false, true);
if (c)
{
OMP_CLAUSE_CHAIN (c) = *clauses;
@@ -15374,8 +15376,8 @@ tsubst_expr (tree t, tree args, tsubst_flags_t complain, tree in_decl,
case OACC_KERNELS:
case OACC_PARALLEL:
- tmp = tsubst_omp_clauses (OMP_CLAUSES (t), false, false, args, complain,
- in_decl);
+ tmp = tsubst_omp_clauses (OMP_CLAUSES (t), false, true, args, complain,
+ in_decl, true);
stmt = begin_omp_parallel ();
RECUR (OMP_BODY (t));
finish_omp_construct (TREE_CODE (t), stmt, tmp);
@@ -15384,7 +15386,7 @@ tsubst_expr (tree t, tree args, tsubst_flags_t complain, tree in_decl,
case OMP_PARALLEL:
r = push_omp_privatization_clauses (OMP_PARALLEL_COMBINED (t));
tmp = tsubst_omp_clauses (OMP_PARALLEL_CLAUSES (t), false, true,
- args, complain, in_decl);
+ args, complain, in_decl, false);
if (OMP_PARALLEL_COMBINED (t))
omp_parallel_combined_clauses = &tmp;
stmt = begin_omp_parallel ();
@@ -15398,7 +15400,7 @@ tsubst_expr (tree t, tree args, tsubst_flags_t complain, tree in_decl,
case OMP_TASK:
r = push_omp_privatization_clauses (false);
tmp = tsubst_omp_clauses (OMP_TASK_CLAUSES (t), false, true,
- args, complain, in_decl);
+ args, complain, in_decl, false);
stmt = begin_omp_task ();
RECUR (OMP_TASK_BODY (t));
finish_omp_task (tmp, stmt);
@@ -15420,9 +15422,9 @@ tsubst_expr (tree t, tree args, tsubst_flags_t complain, tree in_decl,
int i;
r = push_omp_privatization_clauses (OMP_FOR_INIT (t) == NULL_TREE);
- clauses = tsubst_omp_clauses (OMP_FOR_CLAUSES (t), false,
- TREE_CODE (t) != OACC_LOOP,
- args, complain, in_decl);
+ clauses = tsubst_omp_clauses (OMP_FOR_CLAUSES (t), false, true,
+ args, complain, in_decl,
+ TREE_CODE (t) == OACC_LOOP);
if (OMP_FOR_INIT (t) != NULL_TREE)
{
declv = make_tree_vec (TREE_VEC_LENGTH (OMP_FOR_INIT (t)));
@@ -15479,7 +15481,7 @@ tsubst_expr (tree t, tree args, tsubst_flags_t complain, tree in_decl,
r = push_omp_privatization_clauses (TREE_CODE (t) == OMP_TEAMS
&& OMP_TEAMS_COMBINED (t));
tmp = tsubst_omp_clauses (OMP_CLAUSES (t), false, true,
- args, complain, in_decl);
+ args, complain, in_decl, false);
stmt = push_stmt_list ();
RECUR (OMP_BODY (t));
stmt = pop_stmt_list (stmt);
@@ -15494,9 +15496,8 @@ tsubst_expr (tree t, tree args, tsubst_flags_t complain, tree in_decl,
case OACC_DATA:
case OMP_TARGET_DATA:
case OMP_TARGET:
- tmp = tsubst_omp_clauses (OMP_CLAUSES (t), false,
- TREE_CODE (t) != OACC_DATA,
- args, complain, in_decl);
+ tmp = tsubst_omp_clauses (OMP_CLAUSES (t), false, true, args, complain,
+ in_decl, TREE_CODE (t) == OACC_DATA);
keep_next_level (true);
stmt = begin_omp_structured_block ();
@@ -15541,8 +15542,8 @@ tsubst_expr (tree t, tree args, tsubst_flags_t complain, tree in_decl,
case OACC_DECLARE:
t = copy_node (t);
- tmp = tsubst_omp_clauses (OACC_DECLARE_CLAUSES (t), false, false,
- args, complain, in_decl);
+ tmp = tsubst_omp_clauses (OACC_DECLARE_CLAUSES (t), false, true,
+ args, complain, in_decl, true);
OACC_DECLARE_CLAUSES (t) = tmp;
add_stmt (t);
break;
@@ -15551,7 +15552,7 @@ tsubst_expr (tree t, tree args, tsubst_flags_t complain, tree in_decl,
case OMP_TARGET_ENTER_DATA:
case OMP_TARGET_EXIT_DATA:
tmp = tsubst_omp_clauses (OMP_STANDALONE_CLAUSES (t), false, true,
- args, complain, in_decl);
+ args, complain, in_decl, false);
t = copy_node (t);
OMP_STANDALONE_CLAUSES (t) = tmp;
add_stmt (t);
@@ -15560,8 +15561,8 @@ tsubst_expr (tree t, tree args, tsubst_flags_t complain, tree in_decl,
case OACC_ENTER_DATA:
case OACC_EXIT_DATA:
case OACC_UPDATE:
- tmp = tsubst_omp_clauses (OMP_STANDALONE_CLAUSES (t), false, false,
- args, complain, in_decl);
+ tmp = tsubst_omp_clauses (OMP_STANDALONE_CLAUSES (t), false, true,
+ args, complain, in_decl, true);
t = copy_node (t);
OMP_STANDALONE_CLAUSES (t) = tmp;
add_stmt (t);
@@ -15569,7 +15570,7 @@ tsubst_expr (tree t, tree args, tsubst_flags_t complain, tree in_decl,
case OMP_ORDERED:
tmp = tsubst_omp_clauses (OMP_ORDERED_CLAUSES (t), false, true,
- args, complain, in_decl);
+ args, complain, in_decl, false);
stmt = push_stmt_list ();
RECUR (OMP_BODY (t));
stmt = pop_stmt_list (stmt);
@@ -5733,7 +5733,8 @@ cp_finish_omp_clause_depend_sink (tree sink_clause)
Remove any elements from the list that are invalid. */
tree
-finish_omp_clauses (tree clauses, bool allow_fields, bool declare_simd)
+finish_omp_clauses (tree clauses, bool is_oacc, bool allow_fields,
+ bool declare_simd)
{
bitmap_head generic_head, firstprivate_head, lastprivate_head;
bitmap_head aligned_head, map_head, map_field_head;
@@ -5973,7 +5974,10 @@ finish_omp_clauses (tree clauses, bool allow_fields, bool declare_simd)
else if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_PRIVATE
&& bitmap_bit_p (&map_head, DECL_UID (t)))
{
- error ("%qD appears both in data and map clauses", t);
+ if (is_oacc)
+ error ("%qD appears more than once in data clauses", t);
+ else
+ error ("%qD appears both in data and map clauses", t);
remove = true;
}
else
@@ -5999,7 +6003,8 @@ finish_omp_clauses (tree clauses, bool allow_fields, bool declare_simd)
omp_note_field_privatization (t, OMP_CLAUSE_DECL (c));
else
t = OMP_CLAUSE_DECL (c);
- if (t == current_class_ptr)
+ if (!is_oacc
+ && t == current_class_ptr)
{
error ("%<this%> allowed in OpenMP only in %<declare simd%>"
" clauses");
@@ -6025,7 +6030,10 @@ finish_omp_clauses (tree clauses, bool allow_fields, bool declare_simd)
}
else if (bitmap_bit_p (&map_head, DECL_UID (t)))
{
- error ("%qD appears both in data and map clauses", t);
+ if (is_oacc)
+ error ("%qD appears more than once in data clauses", t);
+ else
+ error ("%qD appears both in data and map clauses", t);
remove = true;
}
else
@@ -6535,6 +6543,9 @@ finish_omp_clauses (tree clauses, bool allow_fields, bool declare_simd)
if (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_MAP)
error ("%qD appears more than once in motion"
" clauses", t);
+ else if (is_oacc)
+ error ("%qD appears more than once in data"
+ " clauses", t);
else
error ("%qD appears more than once in map"
" clauses", t);
@@ -6546,6 +6557,27 @@ finish_omp_clauses (tree clauses, bool allow_fields, bool declare_simd)
bitmap_set_bit (&map_field_head, DECL_UID (t));
}
}
+ else if (TREE_CODE (t) == TREE_LIST)
+ {
+ while (TREE_CODE (t = TREE_CHAIN (t)) == TREE_LIST)
+ ;
+
+ if (DECL_P (t))
+ {
+ if (bitmap_bit_p (&map_head, DECL_UID (t)))
+ {
+ if (is_oacc)
+ error ("%qD appears more than once in data "
+ "clauses", t);
+ else
+ error ("%qD appears more than once in map "
+ "clauses", t);
+ remove = true;
+ }
+ else
+ bitmap_set_bit (&map_head, DECL_UID (t));
+ }
+ }
}
break;
}
@@ -6622,7 +6654,8 @@ finish_omp_clauses (tree clauses, bool allow_fields, bool declare_simd)
omp_clause_code_name[OMP_CLAUSE_CODE (c)]);
remove = true;
}
- else if (t == current_class_ptr)
+ else if (!is_oacc
+ && t == current_class_ptr)
{
error ("%<this%> allowed in OpenMP only in %<declare simd%>"
" clauses");
@@ -6663,7 +6696,10 @@ finish_omp_clauses (tree clauses, bool allow_fields, bool declare_simd)
}
else if (bitmap_bit_p (&map_head, DECL_UID (t)))
{
- error ("%qD appears both in data and map clauses", t);
+ if (is_oacc)
+ error ("%qD appears more than once in data clauses", t);
+ else
+ error ("%qD appears both in data and map clauses", t);
remove = true;
}
else
@@ -6673,6 +6709,8 @@ finish_omp_clauses (tree clauses, bool allow_fields, bool declare_simd)
{
if (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_MAP)
error ("%qD appears more than once in motion clauses", t);
+ if (is_oacc)
+ error ("%qD appears more than once in data clauses", t);
else
error ("%qD appears more than once in map clauses", t);
remove = true;
@@ -6680,7 +6718,10 @@ finish_omp_clauses (tree clauses, bool allow_fields, bool declare_simd)
else if (bitmap_bit_p (&generic_head, DECL_UID (t))
|| bitmap_bit_p (&firstprivate_head, DECL_UID (t)))
{
- error ("%qD appears both in data and map clauses", t);
+ if (is_oacc)
+ error ("%qD appears more than once in data clauses", t);
+ else
+ error ("%qD appears both in data and map clauses", t);
remove = true;
}
else
@@ -8257,7 +8298,7 @@ finish_omp_for (location_t locus, enum tree_code code, tree declv,
OMP_CLAUSE_OPERAND (c, 0)
= cilk_for_number_of_iterations (omp_for);
OMP_CLAUSE_CHAIN (c) = clauses;
- OMP_PARALLEL_CLAUSES (omp_par) = finish_omp_clauses (c, false);
+ OMP_PARALLEL_CLAUSES (omp_par) = finish_omp_clauses (c, false, false);
add_stmt (omp_par);
return omp_par;
}
@@ -6471,7 +6471,10 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
case OMP_TARGET_DATA:
case OMP_TARGET_ENTER_DATA:
case OMP_TARGET_EXIT_DATA:
+ case OACC_DATA:
case OACC_HOST_DATA:
+ case OACC_PARALLEL:
+ case OACC_KERNELS:
ctx->target_firstprivatize_array_bases = true;
default:
break;
@@ -6737,10 +6740,16 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
case OMP_TARGET_DATA:
case OMP_TARGET_ENTER_DATA:
case OMP_TARGET_EXIT_DATA:
+ case OACC_DATA:
case OACC_HOST_DATA:
+ case OACC_ENTER_DATA:
+ case OACC_EXIT_DATA:
+ case OACC_UPDATE:
if (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_FIRSTPRIVATE_POINTER
|| (OMP_CLAUSE_MAP_KIND (c)
- == GOMP_MAP_FIRSTPRIVATE_REFERENCE))
+ == GOMP_MAP_FIRSTPRIVATE_REFERENCE)
+ || (!lang_GNU_Fortran () &&
+ OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_POINTER))
/* For target {,enter ,exit }data only the array slice is
mapped, but not the pointer to it. */
remove = true;
@@ -2,12 +2,12 @@ void
fun (void)
{
float *fp;
-#pragma acc parallel copy(fp[0:2],fp[0:2]) /* { dg-error "'fp' appears more than once in map clauses" } */
+#pragma acc parallel copy(fp[0:2],fp[0:2]) /* { dg-error "'fp' appears more than once in data clauses" } */
;
-#pragma acc kernels present_or_copyin(fp[3]) present_or_copyout(fp[7:4]) /* { dg-error "'fp' appears more than once in map clauses" } */
+#pragma acc kernels present_or_copyin(fp[3]) present_or_copyout(fp[7:4]) /* { dg-error "'fp' appears more than once in data clauses" } */
;
-#pragma acc data create(fp[:10]) deviceptr(fp) /* { dg-error "'fp' appears more than once in map clauses" } */
+#pragma acc data create(fp[:10]) deviceptr(fp) /* { dg-error "'fp' appears more than once in data clauses" } */
;
-#pragma acc data create(fp) present(fp) /* { dg-error "'fp' appears more than once in map clauses" } */
+#pragma acc data create(fp) present(fp) /* { dg-error "'fp' appears more than once in data clauses" } */
;
}
@@ -77,3 +77,5 @@ f (void)
#pragma acc declare present (v9) /* { dg-error "invalid use of" } */
}
+
+/* { dg-error "at file scope" "" { target c++ } 10 } */
@@ -47,7 +47,7 @@ fun2 (void)
/* { dg-error "'u' undeclared" "u undeclared" { target *-*-* } 46 } */
/* { dg-error "'fun2' is not a variable" "fun2 not a variable" { target *-*-* } 46 } */
/* { dg-error "'i' is not a pointer variable" "i not a pointer variable" { target *-*-* } 46 } */
- /* { dg-error "'fp' appears more than once in map clauses" "fp more than once" { target *-*-* } 46 } */
+ /* { dg-error "'fp' appears more than once in data clauses" "fp more than once" { target *-*-* } 46 } */
;
}
@@ -55,11 +55,11 @@ void
fun3 (void)
{
float *fp;
-#pragma acc data deviceptr(fp,fp) /* { dg-error "'fp' appears more than once in map clauses" } */
+#pragma acc data deviceptr(fp,fp) /* { dg-error "'fp' appears more than once in data clauses" } */
;
-#pragma acc parallel deviceptr(fp) deviceptr(fp) /* { dg-error "'fp' appears more than once in map clauses" } */
+#pragma acc parallel deviceptr(fp) deviceptr(fp) /* { dg-error "'fp' appears more than once in data clauses" } */
;
-#pragma acc kernels copy(fp) deviceptr(fp) /* { dg-error "'fp' appears more than once in map clauses" } */
+#pragma acc kernels copy(fp) deviceptr(fp) /* { dg-error "'fp' appears more than once in data clauses" } */
;
}
@@ -31,6 +31,6 @@ foo (void)
free (c);
}
-/* { dg-final { scan-tree-dump-times "(?n)= 0;$" 1 "optimized" } } */
-/* { dg-final { scan-tree-dump-times "(?n)= 1;$" 1 "optimized" } } */
-/* { dg-final { scan-tree-dump-times "(?n)= \\*a" 1 "optimized" } } */
+/* { dg-final { scan-tree-dump-times "(?n)= 0;$" 1 "optimized" { target c } } } */
+/* { dg-final { scan-tree-dump-times "(?n)= 1;$" 1 "optimized" { target c } } } */
+/* { dg-final { scan-tree-dump-times "(?n)= \\*a" 1 "optimized" { target c } } } */
@@ -7,4 +7,6 @@ f (char *cp)
;
}
-/* { dg-final { scan-tree-dump-times "#pragma acc parallel map\\(tofrom:\\*\\(cp \\+ 3\\) \\\[len: 5]\\) map\\(alloc:cp \\\[pointer assign, bias: 3]\\)" 1 "original" } } */
+/* { dg-final { scan-tree-dump-times "#pragma acc parallel map\\(tofrom:\\*\\(cp \\+ 3\\) \\\[len: 5]\\) map\\(alloc:cp \\\[pointer assign, bias: 3]\\)" 1 "original" { target c } } } */
+
+/* { dg-final { scan-tree-dump-times "#pragma acc parallel map\\(tofrom:\\*\\(cp \\+ 3\\) \\\[len: 5]\\) map\\(firstprivate:cp \\\[pointer assign, bias: 3]\\)" 1 "original" { target c++ } } } */
@@ -7,4 +7,6 @@ f (char *cp)
;
}
-/* { dg-final { scan-tree-dump-times "#pragma acc parallel map\\(to:\\*\\(cp \\+ 4\\) \\\[len: 6]\\) map\\(alloc:cp \\\[pointer assign, bias: 4]\\)" 1 "original" } } */
+/* { dg-final { scan-tree-dump-times "#pragma acc parallel map\\(to:\\*\\(cp \\+ 4\\) \\\[len: 6]\\) map\\(alloc:cp \\\[pointer assign, bias: 4]\\)" 1 "original" { target c } } } */
+
+/* { dg-final { scan-tree-dump-times "#pragma acc parallel map\\(to:\\*\\(cp \\+ 4\\) \\\[len: 6]\\) map\\(firstprivate:cp \\\[pointer assign, bias: 4]\\)" 1 "original" { target c++ } } } */
@@ -7,4 +7,6 @@ f (char *cp)
;
}
-/* { dg-final { scan-tree-dump-times "#pragma acc parallel map\\(from:\\*\\(cp \\+ 5\\) \\\[len: 7]\\) map\\(alloc:cp \\\[pointer assign, bias: 5]\\)" 1 "original" } } */
+/* { dg-final { scan-tree-dump-times "#pragma acc parallel map\\(from:\\*\\(cp \\+ 5\\) \\\[len: 7]\\) map\\(alloc:cp \\\[pointer assign, bias: 5]\\)" 1 "original" { target c } } } */
+
+/* { dg-final { scan-tree-dump-times "#pragma acc parallel map\\(from:\\*\\(cp \\+ 5\\) \\\[len: 7]\\) map\\(firstprivate:cp \\\[pointer assign, bias: 5]\\)" 1 "original" { target c++ } } } */
@@ -7,4 +7,6 @@ f (char *cp)
;
}
-/* { dg-final { scan-tree-dump-times "#pragma acc parallel map\\(alloc:\\*\\(cp \\+ 6\\) \\\[len: 8]\\) map\\(alloc:cp \\\[pointer assign, bias: 6]\\)" 1 "original" } } */
+/* { dg-final { scan-tree-dump-times "#pragma acc parallel map\\(alloc:\\*\\(cp \\+ 6\\) \\\[len: 8]\\) map\\(alloc:cp \\\[pointer assign, bias: 6]\\)" 1 "original" { target c } } } */
+
+/* { dg-final { scan-tree-dump-times "#pragma acc parallel map\\(alloc:\\*\\(cp \\+ 6\\) \\\[len: 8]\\) map\\(firstprivate:cp \\\[pointer assign, bias: 6]\\)" 1 "original" { target c++ } } } */
@@ -7,4 +7,6 @@ f (char *cp)
;
}
-/* { dg-final { scan-tree-dump-times "#pragma acc parallel map\\(force_present:\\*\\(cp \\+ 7\\) \\\[len: 9]\\) map\\(alloc:cp \\\[pointer assign, bias: 7]\\)" 1 "original" } } */
+/* { dg-final { scan-tree-dump-times "#pragma acc parallel map\\(force_present:\\*\\(cp \\+ 7\\) \\\[len: 9]\\) map\\(alloc:cp \\\[pointer assign, bias: 7]\\)" 1 "original" { target c } } } */
+
+/* { dg-final { scan-tree-dump-times "#pragma acc parallel map\\(force_present:\\*\\(cp \\+ 7\\) \\\[len: 9]\\) map\\(firstprivate:cp \\\[pointer assign, bias: 7]\\)" 1 "original" { target c++ } } } */
new file mode 100644
@@ -0,0 +1,39 @@
+void
+foo (int &a, int (&b)[100], int &n)
+{
+#pragma acc enter data copyin (a, b) async wait
+#pragma acc enter data create (b[20:30]) async wait
+#pragma acc enter data (a) /* { dg-error "expected '#pragma acc' clause before '\\\(' token" } */
+#pragma acc enter data create (b(1:10)) /* { dg-error "expected '\\\)' before '\\\(' token" } */
+#pragma acc exit data delete (a) if (0)
+#pragma acc exit data copyout (b) if (a)
+#pragma acc exit data delete (b)
+#pragma acc enter /* { dg-error "expected 'data' in" } */
+#pragma acc exit /* { dg-error "expected 'data' in" } */
+#pragma acc enter data /* { dg-error "has no data movement clause" } */
+#pragma acc exit data /* { dg-error "has no data movement clause" } */
+#pragma acc enter Data /* { dg-error "invalid pragma before" } */
+#pragma acc exit copyout (b) /* { dg-error "invalid pragma before" } */
+}
+
+template<typename T>
+void
+foo (T &a, T (&b)[100], T &n)
+{
+#pragma acc enter data copyin (a, b) async wait
+#pragma acc enter data create (b[20:30]) async wait
+#pragma acc enter data (a) /* { dg-error "expected '#pragma acc' clause before '\\\(' token" } */
+#pragma acc enter data create (b(1:10)) /* { dg-error "expected '\\\)' before '\\\(' token" } */
+#pragma acc exit data delete (a) if (0)
+#pragma acc exit data copyout (b) if (a)
+#pragma acc exit data delete (b)
+#pragma acc enter /* { dg-error "expected 'data' in" } */
+#pragma acc exit /* { dg-error "expected 'data' in" } */
+#pragma acc enter data /* { dg-error "has no data movement clause" } */
+#pragma acc exit data /* { dg-error "has no data movement clause" } */
+#pragma acc enter Data /* { dg-error "invalid pragma before" } */
+#pragma acc exit copyout (b) /* { dg-error "invalid pragma before" } */
+}
+
+/* { dg-error "has no data movement clause" "" { target *-*-* } 6 } */
+/* { dg-error "has no data movement clause" "" { target *-*-* } 25 } */
new file mode 100644
@@ -0,0 +1,30 @@
+void
+fun (float (&fp)[100])
+{
+ float *dptr = &fp[50];
+
+#pragma acc parallel copy(fp[0:2],fp[0:2]) /* { dg-error "'fp' appears more than once in data clauses" } */
+ ;
+#pragma acc kernels present_or_copyin(fp[3]) present_or_copyout(fp[7:4]) /* { dg-error "'fp' appears more than once in data clauses" } */
+ ;
+#pragma acc data create(fp[:10]) deviceptr(dptr)
+ ;
+#pragma acc data create(fp) present(fp) /* { dg-error "'fp' appears more than once in data clauses" } */
+ ;
+}
+
+template<typename T>
+void
+fun (T (&fp)[100])
+{
+ T *dptr = &fp[50];
+
+#pragma acc parallel copy(fp[0:2],fp[0:2]) /* { dg-error "'fp' appears more than once in data clauses" } */
+ ;
+#pragma acc kernels present_or_copyin(fp[3]) present_or_copyout(fp[7:4]) /* { dg-error "'fp' appears more than once in data clauses" } */
+ ;
+#pragma acc data create(fp[:10]) deviceptr(dptr)
+ ;
+#pragma acc data create(fp) present(fp) /* { dg-error "'fp' appears more than once in data clauses" } */
+ ;
+}
new file mode 100644
@@ -0,0 +1,72 @@
+/* { dg-require-effective-target alloca } */
+/* Integer reductions. */
+
+#define vl 32
+
+int
+main(void)
+{
+ const int n = 1000;
+ int i;
+ int result, array[n];
+ int lresult;
+
+ /* '+' reductions. */
+#pragma acc parallel vector_length (vl)
+#pragma acc loop reduction (+:result)
+ for (i = 0; i < n; i++)
+ result += array[i];
+
+ /* '*' reductions. */
+#pragma acc parallel vector_length (vl)
+#pragma acc loop reduction (*:result)
+ for (i = 0; i < n; i++)
+ result *= array[i];
+
+// result = 0;
+// vresult = 0;
+//
+// /* 'max' reductions. */
+// #pragma acc parallel vector_length (vl)
+// #pragma acc loop reduction (+:result)
+// for (i = 0; i < n; i++)
+// result = result > array[i] ? result : array[i];
+//
+// /* 'min' reductions. */
+// #pragma acc parallel vector_length (vl)
+// #pragma acc loop reduction (+:result)
+// for (i = 0; i < n; i++)
+// result = result < array[i] ? result : array[i];
+
+ /* '&' reductions. */
+#pragma acc parallel vector_length (vl)
+#pragma acc loop reduction (&:result)
+ for (i = 0; i < n; i++)
+ result &= array[i];
+
+ /* '|' reductions. */
+#pragma acc parallel vector_length (vl)
+#pragma acc loop reduction (|:result)
+ for (i = 0; i < n; i++)
+ result |= array[i];
+
+ /* '^' reductions. */
+#pragma acc parallel vector_length (vl)
+#pragma acc loop reduction (^:result)
+ for (i = 0; i < n; i++)
+ result ^= array[i];
+
+ /* '&&' reductions. */
+#pragma acc parallel vector_length (vl)
+#pragma acc loop reduction (&&:lresult)
+ for (i = 0; i < n; i++)
+ lresult = lresult && (result > array[i]);
+
+ /* '||' reductions. */
+#pragma acc parallel vector_length (vl)
+#pragma acc loop reduction (||:lresult)
+ for (i = 0; i < n; i++)
+ lresult = lresult || (result > array[i]);
+
+ return 0;
+}
new file mode 100644
@@ -0,0 +1,22 @@
+void
+f (int &i, int (&a)[10])
+{
+#pragma acc update device(i)
+#pragma acc update host(i)
+#pragma acc update self(i)
+#pragma acc update device(a[1:3])
+#pragma acc update host(a[1:3])
+#pragma acc update self(a[1:3])
+}
+
+template<typename T>
+void
+f (T &i, T (&a)[10])
+{
+#pragma acc update device(i)
+#pragma acc update host(i)
+#pragma acc update self(i)
+#pragma acc update device(a[1:3])
+#pragma acc update host(a[1:3])
+#pragma acc update self(a[1:3])
+}
new file mode 100644
@@ -0,0 +1,18 @@
+void
+fun (float (&fp)[100])
+{
+ float *dptr = &fp[50];
+
+#pragma omp target data map(tofrom:fp[0:2], fp[0:2]) /* { dg-error "'fp' appears more than once in data clauses" } */
+ ;
+}
+
+template<typename T>
+void
+fun (T (&fp)[100])
+{
+ T *dptr = &fp[50];
+
+#pragma omp target data map(tofrom:fp[0:2], fp[0:2]) /* { dg-error "'fp' appears more than once in map clauses" } */
+ ;
+}
new file mode 100644
@@ -0,0 +1,109 @@
+// Ensure that a non-scalar dummy arguments which are implicitly used inside
+// offloaded regions are properly mapped using present_or_copy.
+
+// { dg-do run }
+
+#include <cassert>
+
+const int n = 100;
+
+struct data {
+ int v;
+};
+
+void
+kernels_present (data &d, int &x)
+{
+#pragma acc kernels present (d, x) default (none)
+ {
+ d.v = x;
+ }
+}
+
+void
+parallel_present (data &d, int &x)
+{
+#pragma acc parallel present (d, x) default (none)
+ {
+ d.v = x;
+ }
+}
+
+void
+kernels_implicit (data &d, int &x)
+{
+#pragma acc kernels
+ {
+ d.v = x;
+ }
+}
+
+void
+parallel_implicit (data &d, int &x)
+{
+#pragma acc parallel
+ {
+ d.v = x;
+ }
+}
+
+void
+reference_data (data &d, int &x)
+{
+#pragma acc data copy(d, x)
+ {
+ kernels_present (d, x);
+
+#pragma acc update host(d)
+ assert (d.v == x);
+
+ x = 200;
+#pragma acc update device(x)
+
+ parallel_present (d, x);
+ }
+
+ assert (d.v = x);
+
+ x = 300;
+ kernels_implicit (d, x);
+ assert (d.v = x);
+
+ x = 400;
+ parallel_implicit (d, x);
+ assert (d.v = x);
+}
+
+int
+main ()
+{
+ data d;
+ int x = 100;
+
+#pragma acc data copy(d, x)
+ {
+ kernels_present (d, x);
+
+#pragma acc update host(d)
+ assert (d.v == x);
+
+ x = 200;
+#pragma acc update device(x)
+
+ parallel_present (d, x);
+ }
+
+ assert (d.v = x);
+
+ x = 300;
+ kernels_implicit (d, x);
+ assert (d.v = x);
+
+ x = 400;
+ parallel_implicit (d, x);
+ assert (d.v = x);
+
+ reference_data (d, x);
+
+ return 0;
+}
new file mode 100644
@@ -0,0 +1,178 @@
+/* Test data mappings on variables with reference types. */
+
+/* { dg-do run } */
+
+#include <stdlib.h>
+
+void
+test (int &N, float *(&a), float *(&b), float *(&c), float *(&d), float *(&e))
+{
+ int i;
+
+ for (i = 0; i < N; i++)
+ {
+ a[i] = 3.0;
+ b[i] = 0.0;
+ }
+
+#pragma acc enter data copyin (a[0:N]) copyin (b[0:N]) copyin (N) async
+#pragma acc parallel async wait present (N, a[0:N], b[0:N])
+#pragma acc loop
+ for (i = 0; i < N; i++)
+ b[i] = a[i];
+
+#pragma acc exit data copyout (a[0:N]) copyout (b[0:N]) wait async
+#pragma acc wait
+
+ for (i = 0; i < N; i++)
+ {
+ if (a[i] != 3.0)
+ abort ();
+
+ if (b[i] != 3.0)
+ abort ();
+ }
+
+ for (i = 0; i < N; i++)
+ {
+ a[i] = 2.0;
+ b[i] = 0.0;
+ }
+
+#pragma acc enter data copyin (a[0:N]) copyin (b[0:N]) copyin (N) async (1)
+#pragma acc parallel async (1) present (N, a[0:N], b[0:N])
+#pragma acc loop
+ for (i = 0; i < N; i++)
+ b[i] = a[i];
+
+#pragma acc update host (a[0:N], b[0:N]) wait (1) async (1)
+#pragma acc wait (1)
+
+ for (i = 0; i < N; i++)
+ {
+ if (a[i] != 2.0)
+ abort ();
+
+ if (b[i] != 2.0)
+ abort ();
+ }
+
+ for (i = 0; i < N; i++)
+ {
+ a[i] = 3.0;
+ b[i] = 0.0;
+ c[i] = 0.0;
+ d[i] = 0.0;
+ }
+
+#pragma acc update device (a[0:N], b[0:N])
+#pragma acc enter data copyin (c[0:N], d[0:N]) async (1)
+
+#pragma acc parallel present (N, a[0:N], b[0:N]) async (1) wait (1)
+#pragma acc loop
+ for (i = 0; i < N; i++)
+ b[i] = (a[i] * a[i] * a[i]) / a[i];
+
+#pragma acc parallel present (N, a[0:N], c[0:N]) async (2) wait (1)
+#pragma acc loop
+ for (i = 0; i < N; i++)
+ c[i] = (a[i] + a[i] + a[i] + a[i]) / a[i];
+
+#pragma acc parallel present (N, a[0:N], d[0:N]) async (3) wait (1)
+#pragma acc loop
+ for (i = 0; i < N; i++)
+ d[i] = ((a[i] * a[i] + a[i]) / a[i]) - a[i];
+
+#pragma acc exit data copyout (a[0:N], b[0:N], c[0:N], d[0:N]) wait (1, 2, 3) async (1)
+#pragma acc wait (1)
+
+ for (i = 0; i < N; i++)
+ {
+ if (a[i] != 3.0)
+ abort ();
+
+ if (b[i] != 9.0)
+ abort ();
+
+ if (c[i] != 4.0)
+ abort ();
+
+ if (d[i] != 1.0)
+ abort ();
+ }
+
+ for (i = 0; i < N; i++)
+ {
+ a[i] = 2.0;
+ b[i] = 0.0;
+ c[i] = 0.0;
+ d[i] = 0.0;
+ e[i] = 0.0;
+ }
+
+#pragma acc enter data copyin (a[0:N], b[0:N], c[0:N], d[0:N], e[0:N]) async (1)
+
+#pragma acc parallel loop present (N, a[0:N], b[0:N]) async (1) wait (1)
+ for (int ii = 0; ii < N; ii++)
+ b[ii] = (a[ii] * a[ii] * a[ii]) / a[ii];
+
+#pragma acc parallel loop present (N, a[0:N], c[0:N]) async (2) wait (1)
+ for (int ii = 0; ii < N; ii++)
+ c[ii] = (a[ii] + a[ii] + a[ii] + a[ii]) / a[ii];
+
+#pragma acc parallel loop present (N, a[0:N], d[0:N]) async (3) wait (1)
+ for (int ii = 0; ii < N; ii++)
+ d[ii] = ((a[ii] * a[ii] + a[ii]) / a[ii]) - a[ii];
+
+#pragma acc parallel loop present (N, a[0:N], b[0:N], c[0:N], d[0:N], e[0:N]) wait (1) async (4)
+ for (int ii = 0; ii < N; ii++)
+ e[ii] = a[ii] + b[ii] + c[ii] + d[ii];
+
+#pragma acc exit data copyout (a[0:N], b[0:N], c[0:N], d[0:N], e[0:N]) wait (1, 2, 3, 4) async (1) delete (N)
+#pragma acc wait (1)
+
+ for (i = 0; i < N; i++)
+ {
+ if (a[i] != 2.0)
+ abort ();
+
+ if (b[i] != 4.0)
+ abort ();
+
+ if (c[i] != 4.0)
+ abort ();
+
+ if (d[i] != 1.0)
+ abort ();
+
+ if (e[i] != 11.0)
+ abort ();
+ }
+}
+
+int
+main (int argc, char **argv)
+{
+ int N = 128; //1024 * 1024;
+ float *a, *b, *c, *d, *e;
+ int i;
+ int nbytes;
+
+ nbytes = N * sizeof (float);
+
+ a = (float *) malloc (nbytes);
+ b = (float *) malloc (nbytes);
+ c = (float *) malloc (nbytes);
+ d = (float *) malloc (nbytes);
+ e = (float *) malloc (nbytes);
+
+ test (N, a, b, c, d, e);
+
+ free (a);
+ free (b);
+ free (c);
+ free (d);
+ free (e);
+
+ return 0;
+}
new file mode 100644
@@ -0,0 +1,179 @@
+/* Test data mappings on variables with template types. */
+
+/* { dg-do run } */
+
+#include <stdlib.h>
+
+template<typename T1, typename T2>
+void
+test (T1 &N, T2 *(&a), T2 *(&b), T2 *(&c), T2 *(&d), T2 *(&e))
+{
+ int i;
+
+ for (i = 0; i < N; i++)
+ {
+ a[i] = 3.0;
+ b[i] = 0.0;
+ }
+
+#pragma acc enter data copyin (a[0:N]) copyin (b[0:N]) copyin (N) async
+#pragma acc parallel async wait present (N, a[0:N], b[0:N])
+#pragma acc loop
+ for (i = 0; i < N; i++)
+ b[i] = a[i];
+
+#pragma acc exit data copyout (a[0:N]) copyout (b[0:N]) wait async
+#pragma acc wait
+
+ for (i = 0; i < N; i++)
+ {
+ if (a[i] != 3.0)
+ abort ();
+
+ if (b[i] != 3.0)
+ abort ();
+ }
+
+ for (i = 0; i < N; i++)
+ {
+ a[i] = 2.0;
+ b[i] = 0.0;
+ }
+
+#pragma acc enter data copyin (a[0:N]) copyin (b[0:N]) copyin (N) async (1)
+#pragma acc parallel async (1) present (N, a[0:N], b[0:N])
+#pragma acc loop
+ for (i = 0; i < N; i++)
+ b[i] = a[i];
+
+#pragma acc update host (a[0:N], b[0:N]) wait (1) async (1)
+#pragma acc wait (1)
+
+ for (i = 0; i < N; i++)
+ {
+ if (a[i] != 2.0)
+ abort ();
+
+ if (b[i] != 2.0)
+ abort ();
+ }
+
+ for (i = 0; i < N; i++)
+ {
+ a[i] = 3.0;
+ b[i] = 0.0;
+ c[i] = 0.0;
+ d[i] = 0.0;
+ }
+
+#pragma acc update device (a[0:N], b[0:N])
+#pragma acc enter data copyin (c[0:N], d[0:N]) async (1)
+
+#pragma acc parallel present (N, a[0:N], b[0:N]) async (1) wait (1)
+#pragma acc loop
+ for (i = 0; i < N; i++)
+ b[i] = (a[i] * a[i] * a[i]) / a[i];
+
+#pragma acc parallel present (N, a[0:N], c[0:N]) async (2) wait (1)
+#pragma acc loop
+ for (i = 0; i < N; i++)
+ c[i] = (a[i] + a[i] + a[i] + a[i]) / a[i];
+
+#pragma acc parallel present (N, a[0:N], d[0:N]) async (3) wait (1)
+#pragma acc loop
+ for (i = 0; i < N; i++)
+ d[i] = ((a[i] * a[i] + a[i]) / a[i]) - a[i];
+
+#pragma acc exit data copyout (a[0:N], b[0:N], c[0:N], d[0:N]) wait (1, 2, 3) async (1)
+#pragma acc wait (1)
+
+ for (i = 0; i < N; i++)
+ {
+ if (a[i] != 3.0)
+ abort ();
+
+ if (b[i] != 9.0)
+ abort ();
+
+ if (c[i] != 4.0)
+ abort ();
+
+ if (d[i] != 1.0)
+ abort ();
+ }
+
+ for (i = 0; i < N; i++)
+ {
+ a[i] = 2.0;
+ b[i] = 0.0;
+ c[i] = 0.0;
+ d[i] = 0.0;
+ e[i] = 0.0;
+ }
+
+#pragma acc enter data copyin (a[0:N], b[0:N], c[0:N], d[0:N], e[0:N]) async (1)
+
+#pragma acc parallel loop present (N, a[0:N], b[0:N]) async (1) wait (1)
+ for (int ii = 0; ii < N; ii++)
+ b[ii] = (a[ii] * a[ii] * a[ii]) / a[ii];
+
+#pragma acc parallel loop present (N, a[0:N], c[0:N]) async (2) wait (1)
+ for (int ii = 0; ii < N; ii++)
+ c[ii] = (a[ii] + a[ii] + a[ii] + a[ii]) / a[ii];
+
+#pragma acc parallel loop present (N, a[0:N], d[0:N]) async (3) wait (1)
+ for (int ii = 0; ii < N; ii++)
+ d[ii] = ((a[ii] * a[ii] + a[ii]) / a[ii]) - a[ii];
+
+#pragma acc parallel loop present (N, a[0:N], b[0:N], c[0:N], d[0:N], e[0:N]) wait (1) async (4)
+ for (int ii = 0; ii < N; ii++)
+ e[ii] = a[ii] + b[ii] + c[ii] + d[ii];
+
+#pragma acc exit data copyout (a[0:N], b[0:N], c[0:N], d[0:N], e[0:N]) wait (1, 2, 3, 4) async (1) delete (N)
+#pragma acc wait (1)
+
+ for (i = 0; i < N; i++)
+ {
+ if (a[i] != 2.0)
+ abort ();
+
+ if (b[i] != 4.0)
+ abort ();
+
+ if (c[i] != 4.0)
+ abort ();
+
+ if (d[i] != 1.0)
+ abort ();
+
+ if (e[i] != 11.0)
+ abort ();
+ }
+}
+
+int
+main (int argc, char **argv)
+{
+ int N = 128; //1024 * 1024;
+ float *a, *b, *c, *d, *e;
+ int i;
+ int nbytes;
+
+ nbytes = N * sizeof (float);
+
+ a = (float *) malloc (nbytes);
+ b = (float *) malloc (nbytes);
+ c = (float *) malloc (nbytes);
+ d = (float *) malloc (nbytes);
+ e = (float *) malloc (nbytes);
+
+ test (N, a, b, c, d, e);
+
+ free (a);
+ free (b);
+ free (c);
+ free (d);
+ free (e);
+
+ return 0;
+}
new file mode 100644
@@ -0,0 +1,114 @@
+// Ensure that a non-scalar dummy arguments which are implicitly used inside
+// offloaded regions are properly mapped using present_or_copy.
+
+// { dg-do run }
+
+#include <cassert>
+
+const int n = 100;
+
+struct data {
+ int v;
+};
+
+template <class t1, class t2>
+void
+kernels_present (t1 &d, t2 &x)
+{
+#pragma acc kernels present (d, x) default (none)
+ {
+ d.v = x;
+ }
+}
+
+template <class t1, class t2>
+void
+parallel_present (t1 &d, t2 &x)
+{
+#pragma acc parallel present (d, x) default (none)
+ {
+ d.v = x;
+ }
+}
+
+template <class t1, class t2>
+void
+kernels_implicit (t1 &d, t2 &x)
+{
+#pragma acc kernels
+ {
+ d.v = x;
+ }
+}
+
+template <class t1, class t2>
+void
+parallel_implicit (t1 &d, t2 &x)
+{
+#pragma acc parallel
+ {
+ d.v = x;
+ }
+}
+
+template <class t1, class t2>
+void
+reference_data (t1 &d, t2 &x)
+{
+#pragma acc data copy(d, x)
+ {
+ kernels_present (d, x);
+
+#pragma acc update host(d)
+ assert (d.v == x);
+
+ x = 200;
+#pragma acc update device(x)
+
+ parallel_present (d, x);
+ }
+
+ assert (d.v = x);
+
+ x = 300;
+ kernels_implicit (d, x);
+ assert (d.v = x);
+
+ x = 400;
+ parallel_implicit (d, x);
+ assert (d.v = x);
+}
+
+int
+main ()
+{
+ data d;
+ int x = 100;
+
+#pragma acc data copy(d, x)
+ {
+ kernels_present (d, x);
+
+#pragma acc update host(d)
+ assert (d.v == x);
+
+ x = 200;
+#pragma acc update device(x)
+
+ parallel_present (d, x);
+ }
+
+ assert (d.v = x);
+
+ x = 300;
+ kernels_implicit (d, x);
+ assert (d.v = x);
+
+ x = 400;
+ parallel_implicit (d, x);
+ assert (d.v = x);
+
+ reference_data (d, x);
+
+ return 0;
+}
new file mode 100644
@@ -0,0 +1,300 @@
+/* Copy of update-1.c with self exchanged with host for #pragma acc update.
+ This exercises references types. */
+
+/* { dg-do run } */
+/* { dg-skip-if "" { *-*-* } { "*" } { "-DACC_MEM_SHARED=0" } } */
+
+#include <openacc.h>
+#include <string.h>
+#include <stdio.h>
+#include <stdlib.h>
+#include <stdbool.h>
+
+int
+test (int &N, float *(&a), float *(&b), float *(&c), float *(&d_a),
+ float *(&d_b), float *(&d_c), int &i)
+{
+
+
+ for (i = 0; i < N; i++)
+ {
+ a[i] = 3.0;
+ b[i] = 0.0;
+ }
+
+ acc_map_data (a, d_a, N * sizeof (float));
+ acc_map_data (b, d_b, N * sizeof (float));
+ acc_map_data (c, d_c, N * sizeof (float));
+
+#pragma acc update device (a[0:N], b[0:N])
+
+#pragma acc parallel present (a[0:N], b[0:N])
+ {
+ int ii;
+
+ for (ii = 0; ii < N; ii++)
+ b[ii] = a[ii];
+ }
+
+#pragma acc update self (a[0:N], b[0:N])
+
+ for (i = 0; i < N; i++)
+ {
+ if (a[i] != 3.0)
+ abort ();
+
+ if (b[i] != 3.0)
+ abort ();
+ }
+
+ if (!acc_is_present (&a[0], (N * sizeof (float))))
+ abort ();
+
+ if (!acc_is_present (&b[0], (N * sizeof (float))))
+ abort ();
+
+ for (i = 0; i < N; i++)
+ {
+ a[i] = 5.0;
+ b[i] = 1.0;
+ }
+
+#pragma acc update device (a[0:N], b[0:N])
+
+#pragma acc parallel present (a[0:N], b[0:N])
+ {
+ int ii;
+
+ for (ii = 0; ii < N; ii++)
+ b[ii] = a[ii];
+ }
+
+#pragma acc update self (a[0:N], b[0:N])
+
+ for (i = 0; i < N; i++)
+ {
+ if (a[i] != 5.0)
+ abort ();
+
+ if (b[i] != 5.0)
+ abort ();
+ }
+
+ if (!acc_is_present (&a[0], (N * sizeof (float))))
+ abort ();
+
+ if (!acc_is_present (&b[0], (N * sizeof (float))))
+ abort ();
+
+ for (i = 0; i < N; i++)
+ {
+ a[i] = 5.0;
+ b[i] = 1.0;
+ }
+
+#pragma acc update device (a[0:N], b[0:N])
+
+#pragma acc parallel present (a[0:N], b[0:N])
+ {
+ int ii;
+
+ for (ii = 0; ii < N; ii++)
+ b[ii] = a[ii];
+ }
+
+#pragma acc update host (a[0:N], b[0:N])
+
+ for (i = 0; i < N; i++)
+ {
+ if (a[i] != 5.0)
+ abort ();
+
+ if (b[i] != 5.0)
+ abort ();
+ }
+
+ if (!acc_is_present (&a[0], (N * sizeof (float))))
+ abort ();
+
+ if (!acc_is_present (&b[0], (N * sizeof (float))))
+ abort ();
+
+ for (i = 0; i < N; i++)
+ {
+ a[i] = 6.0;
+ b[i] = 0.0;
+ }
+
+#pragma acc update device (a[0:N], b[0:N])
+
+ for (i = 0; i < N; i++)
+ {
+ a[i] = 9.0;
+ }
+
+#pragma acc parallel present (a[0:N], b[0:N])
+ {
+ int ii;
+
+ for (ii = 0; ii < N; ii++)
+ b[ii] = a[ii];
+ }
+
+#pragma acc update self (a[0:N], b[0:N])
+
+ for (i = 0; i < N; i++)
+ {
+ if (a[i] != 6.0)
+ abort ();
+
+ if (b[i] != 6.0)
+ abort ();
+ }
+
+ if (!acc_is_present (&a[0], (N * sizeof (float))))
+ abort ();
+
+ if (!acc_is_present (&b[0], (N * sizeof (float))))
+ abort ();
+
+ for (i = 0; i < N; i++)
+ {
+ a[i] = 7.0;
+ b[i] = 2.0;
+ }
+
+#pragma acc update device (a[0:N], b[0:N])
+
+ for (i = 0; i < N; i++)
+ {
+ a[i] = 9.0;
+ }
+
+#pragma acc parallel present (a[0:N], b[0:N])
+ {
+ int ii;
+
+ for (ii = 0; ii < N; ii++)
+ b[ii] = a[ii];
+ }
+
+#pragma acc update self (a[0:N], b[0:N])
+
+ for (i = 0; i < N; i++)
+ {
+ if (a[i] != 7.0)
+ abort ();
+
+ if (b[i] != 7.0)
+ abort ();
+ }
+
+ for (i = 0; i < N; i++)
+ {
+ a[i] = 9.0;
+ }
+
+#pragma acc update device (a[0:N])
+
+#pragma acc parallel present (a[0:N], b[0:N])
+ {
+ int ii;
+
+ for (ii = 0; ii < N; ii++)
+ b[ii] = a[ii];
+ }
+
+#pragma acc update self (a[0:N], b[0:N])
+
+ for (i = 0; i < N; i++)
+ {
+ if (a[i] != 9.0)
+ abort ();
+
+ if (b[i] != 9.0)
+ abort ();
+ }
+
+ if (!acc_is_present (&a[0], (N * sizeof (float))))
+ abort ();
+
+ if (!acc_is_present (&b[0], (N * sizeof (float))))
+ abort ();
+
+ for (i = 0; i < N; i++)
+ {
+ a[i] = 5.0;
+ }
+
+#pragma acc update device (a[0:N])
+
+ for (i = 0; i < N; i++)
+ {
+ a[i] = 6.0;
+ }
+
+#pragma acc update device (a[0:N >> 1])
+
+#pragma acc parallel present (a[0:N], b[0:N])
+ {
+ int ii;
+
+ for (ii = 0; ii < N; ii++)
+ b[ii] = a[ii];
+ }
+
+#pragma acc update self (a[0:N], b[0:N])
+
+ for (i = 0; i < (N >> 1); i++)
+ {
+ if (a[i] != 6.0)
+ abort ();
+
+ if (b[i] != 6.0)
+ abort ();
+ }
+
+ for (i = (N >> 1); i < N; i++)
+ {
+ if (a[i] != 5.0)
+ abort ();
+
+ if (b[i] != 5.0)
+ abort ();
+ }
+
+ if (!acc_is_present (&a[0], (N * sizeof (float))))
+ abort ();
+
+ if (!acc_is_present (&b[0], (N * sizeof (float))))
+ abort ();
+}
+
+int
+main (int argc, char **argv)
+{
+ int N = 8;
+ float *a, *b, *c;
+ float *d_a, *d_b, *d_c;
+ int i;
+
+ a = (float *) malloc (N * sizeof (float));
+ b = (float *) malloc (N * sizeof (float));
+ c = (float *) malloc (N * sizeof (float));
+
+ d_a = (float *) acc_malloc (N * sizeof (float));
+ d_b = (float *) acc_malloc (N * sizeof (float));
+ d_c = (float *) acc_malloc (N * sizeof (float));
+
+ test (N, a, b, c, d_a, d_b, d_c, i);
+
+ acc_free (d_a);
+ acc_free (d_b);
+ acc_free (d_c);
+
+ free (a);
+ free (b);
+ free (c);
+
+ return 0;
+}
new file mode 100644
@@ -0,0 +1,301 @@
+/* Copy of update-1.c with self exchanged with host for #pragma acc update.
+ This exercises templates. */
+
+/* { dg-do run } */
+/* { dg-skip-if "" { *-*-* } { "*" } { "-DACC_MEM_SHARED=0" } } */
+
+#include <openacc.h>
+#include <string.h>
+#include <stdio.h>
+#include <stdlib.h>
+#include <stdbool.h>
+
+template<typename T1, typename T2>
+int
+test (T1 &N, T2 *(&a), T2 *(&b), T2 *(&c), T2 *(&d_a), T2 *(&d_b),
+ T2 *(&d_c), T1 &i)
+{
+
+
+ for (i = 0; i < N; i++)
+ {
+ a[i] = 3.0;
+ b[i] = 0.0;
+ }
+
+ acc_map_data (a, d_a, N * sizeof (float));
+ acc_map_data (b, d_b, N * sizeof (float));
+ acc_map_data (c, d_c, N * sizeof (float));
+
+#pragma acc update device (a[0:N], b[0:N])
+
+#pragma acc parallel present (a[0:N], b[0:N])
+ {
+ int ii;
+
+ for (ii = 0; ii < N; ii++)
+ b[ii] = a[ii];
+ }
+
+#pragma acc update self (a[0:N], b[0:N])
+
+ for (i = 0; i < N; i++)
+ {
+ if (a[i] != 3.0)
+ abort ();
+
+ if (b[i] != 3.0)
+ abort ();
+ }
+
+ if (!acc_is_present (&a[0], (N * sizeof (float))))
+ abort ();
+
+ if (!acc_is_present (&b[0], (N * sizeof (float))))
+ abort ();
+
+ for (i = 0; i < N; i++)
+ {
+ a[i] = 5.0;
+ b[i] = 1.0;
+ }
+
+#pragma acc update device (a[0:N], b[0:N])
+
+#pragma acc parallel present (a[0:N], b[0:N])
+ {
+ int ii;
+
+ for (ii = 0; ii < N; ii++)
+ b[ii] = a[ii];
+ }
+
+#pragma acc update self (a[0:N], b[0:N])
+
+ for (i = 0; i < N; i++)
+ {
+ if (a[i] != 5.0)
+ abort ();
+
+ if (b[i] != 5.0)
+ abort ();
+ }
+
+ if (!acc_is_present (&a[0], (N * sizeof (float))))
+ abort ();
+
+ if (!acc_is_present (&b[0], (N * sizeof (float))))
+ abort ();
+
+ for (i = 0; i < N; i++)
+ {
+ a[i] = 5.0;
+ b[i] = 1.0;
+ }
+
+#pragma acc update device (a[0:N], b[0:N])
+
+#pragma acc parallel present (a[0:N], b[0:N])
+ {
+ int ii;
+
+ for (ii = 0; ii < N; ii++)
+ b[ii] = a[ii];
+ }
+
+#pragma acc update host (a[0:N], b[0:N])
+
+ for (i = 0; i < N; i++)
+ {
+ if (a[i] != 5.0)
+ abort ();
+
+ if (b[i] != 5.0)
+ abort ();
+ }
+
+ if (!acc_is_present (&a[0], (N * sizeof (float))))
+ abort ();
+
+ if (!acc_is_present (&b[0], (N * sizeof (float))))
+ abort ();
+
+ for (i = 0; i < N; i++)
+ {
+ a[i] = 6.0;
+ b[i] = 0.0;
+ }
+
+#pragma acc update device (a[0:N], b[0:N])
+
+ for (i = 0; i < N; i++)
+ {
+ a[i] = 9.0;
+ }
+
+#pragma acc parallel present (a[0:N], b[0:N])
+ {
+ int ii;
+
+ for (ii = 0; ii < N; ii++)
+ b[ii] = a[ii];
+ }
+
+#pragma acc update self (a[0:N], b[0:N])
+
+ for (i = 0; i < N; i++)
+ {
+ if (a[i] != 6.0)
+ abort ();
+
+ if (b[i] != 6.0)
+ abort ();
+ }
+
+ if (!acc_is_present (&a[0], (N * sizeof (float))))
+ abort ();
+
+ if (!acc_is_present (&b[0], (N * sizeof (float))))
+ abort ();
+
+ for (i = 0; i < N; i++)
+ {
+ a[i] = 7.0;
+ b[i] = 2.0;
+ }
+
+#pragma acc update device (a[0:N], b[0:N])
+
+ for (i = 0; i < N; i++)
+ {
+ a[i] = 9.0;
+ }
+
+#pragma acc parallel present (a[0:N], b[0:N])
+ {
+ int ii;
+
+ for (ii = 0; ii < N; ii++)
+ b[ii] = a[ii];
+ }
+
+#pragma acc update self (a[0:N], b[0:N])
+
+ for (i = 0; i < N; i++)
+ {
+ if (a[i] != 7.0)
+ abort ();
+
+ if (b[i] != 7.0)
+ abort ();
+ }
+
+ for (i = 0; i < N; i++)
+ {
+ a[i] = 9.0;
+ }
+
+#pragma acc update device (a[0:N])
+
+#pragma acc parallel present (a[0:N], b[0:N])
+ {
+ int ii;
+
+ for (ii = 0; ii < N; ii++)
+ b[ii] = a[ii];
+ }
+
+#pragma acc update self (a[0:N], b[0:N])
+
+ for (i = 0; i < N; i++)
+ {
+ if (a[i] != 9.0)
+ abort ();
+
+ if (b[i] != 9.0)
+ abort ();
+ }
+
+ if (!acc_is_present (&a[0], (N * sizeof (float))))
+ abort ();
+
+ if (!acc_is_present (&b[0], (N * sizeof (float))))
+ abort ();
+
+ for (i = 0; i < N; i++)
+ {
+ a[i] = 5.0;
+ }
+
+#pragma acc update device (a[0:N])
+
+ for (i = 0; i < N; i++)
+ {
+ a[i] = 6.0;
+ }
+
+#pragma acc update device (a[0:N >> 1])
+
+#pragma acc parallel present (a[0:N], b[0:N])
+ {
+ int ii;
+
+ for (ii = 0; ii < N; ii++)
+ b[ii] = a[ii];
+ }
+
+#pragma acc update self (a[0:N], b[0:N])
+
+ for (i = 0; i < (N >> 1); i++)
+ {
+ if (a[i] != 6.0)
+ abort ();
+
+ if (b[i] != 6.0)
+ abort ();
+ }
+
+ for (i = (N >> 1); i < N; i++)
+ {
+ if (a[i] != 5.0)
+ abort ();
+
+ if (b[i] != 5.0)
+ abort ();
+ }
+
+ if (!acc_is_present (&a[0], (N * sizeof (float))))
+ abort ();
+
+ if (!acc_is_present (&b[0], (N * sizeof (float))))
+ abort ();
+}
+
+int
+main (int argc, char **argv)
+{
+ int N = 8;
+ float *a, *b, *c;
+ float *d_a, *d_b, *d_c;
+ int i;
+
+ a = (float *) malloc (N * sizeof (float));
+ b = (float *) malloc (N * sizeof (float));
+ c = (float *) malloc (N * sizeof (float));
+
+ d_a = (float *) acc_malloc (N * sizeof (float));
+ d_b = (float *) acc_malloc (N * sizeof (float));
+ d_c = (float *) acc_malloc (N * sizeof (float));
+
+ test (N, a, b, c, d_a, d_b, d_c, i);
+
+ acc_free (d_a);
+ acc_free (d_b);
+ acc_free (d_c);
+
+ free (a);
+ free (b);
+ free (c);
+
+ return 0;
+}
@@ -1,3 +1,5 @@
+/* Test acc_set_cuda_stream integration with async and wait. */
+
/* { dg-do run { target openacc_nvidia_accel_selected } } */
/* { dg-additional-options "-lcuda" } */
@@ -37,7 +39,7 @@ main (int argc, char **argv)
#pragma acc data copy (a[0:N]) copy (b[0:N]) copyin (N)
{
-#pragma acc parallel async
+#pragma acc parallel async present (a[0:N], b[0:N], N)
{
int ii;
@@ -67,7 +69,7 @@ main (int argc, char **argv)
#pragma acc data copy (a[0:N]) copy (b[0:N]) copyin (N)
{
-#pragma acc parallel async (1)
+#pragma acc parallel async (1) present (a[0:N], b[0:N], N)
{
int ii;
@@ -99,7 +101,7 @@ main (int argc, char **argv)
#pragma acc data copy (a[0:N]) copy (b[0:N]) copy (c[0:N]) copy (d[0:N]) copyin (N)
{
-#pragma acc parallel async (1)
+#pragma acc parallel async (1) present (a[0:N], b[0:N], N)
{
int ii;
@@ -107,7 +109,7 @@ main (int argc, char **argv)
b[ii] = (a[ii] * a[ii] * a[ii]) / a[ii];
}
-#pragma acc parallel async (1)
+#pragma acc parallel async (1) present (a[0:N], c[0:N], N)
{
int ii;
@@ -116,7 +118,7 @@ main (int argc, char **argv)
}
-#pragma acc parallel async (1)
+#pragma acc parallel async (1) present (a[0:N], d[0:N], N)
{
int ii;
@@ -155,7 +157,7 @@ main (int argc, char **argv)
#pragma acc data copy (a[0:N], b[0:N], c[0:N], d[0:N], e[0:N]) copyin (N)
{
-#pragma acc parallel async (1)
+#pragma acc parallel async (1) present (a[0:N], b[0:N], N)
{
int ii;
@@ -163,7 +165,7 @@ main (int argc, char **argv)
b[ii] = (a[ii] * a[ii] * a[ii]) / a[ii];
}
-#pragma acc parallel async (1)
+#pragma acc parallel async (1) present (a[0:N], c[0:N], N)
{
int ii;
@@ -171,7 +173,7 @@ main (int argc, char **argv)
c[ii] = (a[ii] + a[ii] + a[ii] + a[ii]) / a[ii];
}
-#pragma acc parallel async (1)
+#pragma acc parallel async (1) present (a[0:N], d[0:N], N)
{
int ii;
@@ -179,7 +181,8 @@ main (int argc, char **argv)
d[ii] = ((a[ii] * a[ii] + a[ii]) / a[ii]) - a[ii];
}
-#pragma acc parallel wait (1) async (1)
+#pragma acc parallel wait (1) async (1) present (a[0:N], b[0:N], c[0:N]) \
+ present (d[0:N], e[0:N], N)
{
int ii;
@@ -228,7 +231,7 @@ main (int argc, char **argv)
#pragma acc data copy (a[0:N], b[0:N]) copyin (N)
{
-#pragma acc parallel async (1)
+#pragma acc parallel async (1) present (a[0:N], b[0:N], N)
{
int ii;
@@ -257,10 +260,11 @@ main (int argc, char **argv)
d[i] = 0.0;
}
-#pragma acc data copy (a[0:N]) copy (b[0:N]) copy (c[0:N]) copy (d[0:N]) copyin (N)
+#pragma acc data copy (a[0:N]) copy (b[0:N]) copy (c[0:N]) copy (d[0:N]) \
+ copyin (N)
{
-#pragma acc parallel async (1)
+#pragma acc parallel async (1) present (a[0:N], b[0:N], N)
{
int ii;
@@ -268,7 +272,7 @@ main (int argc, char **argv)
b[ii] = (a[ii] * a[ii] * a[ii]) / a[ii];
}
-#pragma acc parallel async (1)
+#pragma acc parallel async (1) present (a[0:N], c[0:N], N)
{
int ii;
@@ -276,7 +280,7 @@ main (int argc, char **argv)
c[ii] = (a[ii] + a[ii] + a[ii] + a[ii]) / a[ii];
}
-#pragma acc parallel async (1)
+#pragma acc parallel async (1) present (a[0:N], d[0:N], N)
{
int ii;
@@ -315,7 +319,7 @@ main (int argc, char **argv)
#pragma acc data copy (a[0:N], b[0:N], c[0:N], d[0:N], e[0:N]) copyin (N)
{
-#pragma acc parallel async (1)
+#pragma acc parallel async (1) present (a[0:N], b[0:N], N)
{
int ii;
@@ -323,7 +327,7 @@ main (int argc, char **argv)
b[ii] = (a[ii] * a[ii] * a[ii]) / a[ii];
}
-#pragma acc parallel async (1)
+#pragma acc parallel async (1) present (a[0:N], c[0:N], N)
{
int ii;
@@ -331,7 +335,7 @@ main (int argc, char **argv)
c[ii] = (a[ii] + a[ii] + a[ii] + a[ii]) / a[ii];
}
-#pragma acc parallel async (1)
+#pragma acc parallel async (1) present (a[0:N], d[0:N], N)
{
int ii;
@@ -339,7 +343,8 @@ main (int argc, char **argv)
d[ii] = ((a[ii] * a[ii] + a[ii]) / a[ii]) - a[ii];
}
-#pragma acc parallel wait (1) async (1)
+#pragma acc parallel wait (1) async (1) present (a[0:N], b[0:N], c[0:N]) \
+ present (d[0:N], e[0:N], N)
{
int ii;
@@ -381,7 +386,7 @@ main (int argc, char **argv)
#pragma acc data copyin (a[0:N], b[0:N], c[0:N]) copyin (N)
{
-#pragma acc parallel async (1)
+#pragma acc parallel async (1) present (a[0:N], b[0:N], N)
{
int ii;
@@ -389,7 +394,7 @@ main (int argc, char **argv)
b[ii] = (a[ii] * a[ii] * a[ii]) / a[ii];
}
-#pragma acc parallel async (1)
+#pragma acc parallel async (1) present (a[0:N], c[0:N], N)
{
int ii;
@@ -426,7 +431,7 @@ main (int argc, char **argv)
#pragma acc data copyin (a[0:N], b[0:N], c[0:N]) copyin (N)
{
-#pragma acc parallel async (1)
+#pragma acc parallel async (1) present (a[0:N], b[0:N], N)
{
int ii;
@@ -434,7 +439,7 @@ main (int argc, char **argv)
b[ii] = (a[ii] * a[ii] * a[ii]) / a[ii];
}
-#pragma acc parallel async (1)
+#pragma acc parallel async (1) present (a[0:N], c[0:N], N)
{
int ii;
@@ -462,5 +467,11 @@ main (int argc, char **argv)
acc_shutdown (acc_device_nvidia);
+ free (a);
+ free (b);
+ free (c);
+ free (d);
+ free (e);
+
return 0;
}
@@ -1,3 +1,5 @@
+/* Test "acc enter data" and "acc exit data". */
+
/* { dg-do run } */
#include <stdlib.h>
@@ -25,7 +27,7 @@ main (int argc, char **argv)
}
#pragma acc enter data copyin (a[0:N]) copyin (b[0:N]) copyin (N) async
-#pragma acc parallel async wait
+#pragma acc parallel async wait present (a[0:N], b[0:N], N)
#pragma acc loop
for (i = 0; i < N; i++)
b[i] = a[i];
@@ -41,7 +43,7 @@ main (int argc, char **argv)
if (b[i] != 3.0)
abort ();
}
-
+ return 0;
for (i = 0; i < N; i++)
{
a[i] = 2.0;
@@ -49,7 +51,7 @@ main (int argc, char **argv)
}
#pragma acc enter data copyin (a[0:N]) copyin (b[0:N]) copyin (N) async (1)
-#pragma acc parallel async (1)
+#pragma acc parallel async (1) present (a[0:N], b[0:N], N)
#pragma acc loop
for (i = 0; i < N; i++)
b[i] = a[i];
@@ -74,24 +76,26 @@ main (int argc, char **argv)
d[i] = 0.0;
}
-#pragma acc enter data copyin (a[0:N]) copyin (b[0:N]) copyin (c[0:N]) copyin (d[0:N]) copyin (N) async (1)
+#pragma acc enter data copyin (a[0:N]) copyin (b[0:N]) copyin (c[0:N]) \
+ copyin (d[0:N]) copyin (N) async (1)
-#pragma acc parallel async (1) wait (1)
+#pragma acc parallel async (1) wait (1) present (a[0:N], b[0:N], N)
#pragma acc loop
for (i = 0; i < N; i++)
b[i] = (a[i] * a[i] * a[i]) / a[i];
-#pragma acc parallel async (2) wait (1)
+#pragma acc parallel async (2) wait (1) present (a[0:N], c[0:N], N)
#pragma acc loop
for (i = 0; i < N; i++)
c[i] = (a[i] + a[i] + a[i] + a[i]) / a[i];
-#pragma acc parallel async (3) wait (1)
+#pragma acc parallel async (3) wait (1) present (a[0:N], d[0:N], N)
#pragma acc loop
for (i = 0; i < N; i++)
d[i] = ((a[i] * a[i] + a[i]) / a[i]) - a[i];
-#pragma acc exit data copyout (a[0:N]) copyout (b[0:N]) copyout (c[0:N]) copyout (d[0:N]) wait (1, 2, 3) async (1)
+#pragma acc exit data copyout (a[0:N]) copyout (b[0:N]) copyout (c[0:N]) \
+ copyout (d[0:N]) wait (1, 2, 3) async (1)
#pragma acc wait (1)
for (i = 0; i < N; i++)
@@ -118,28 +122,30 @@ main (int argc, char **argv)
e[i] = 0.0;
}
-#pragma acc enter data copyin (a[0:N]) copyin (b[0:N]) copyin (c[0:N]) copyin (d[0:N]) copyin (e[0:N]) copyin (N) async (1)
+#pragma acc enter data copyin (a[0:N]) copyin (b[0:N]) copyin (c[0:N]) \
+ copyin (d[0:N]) copyin (e[0:N]) copyin (N) async (1)
-#pragma acc parallel async (1) wait (1)
+#pragma acc parallel async (1) wait (1) present (a[0:N], b[0:N], N)
for (int ii = 0; ii < N; ii++)
b[ii] = (a[ii] * a[ii] * a[ii]) / a[ii];
-#pragma acc parallel async (2) wait (1)
+#pragma acc parallel async (2) wait (1) present (a[0:N], c[0:N], N)
for (int ii = 0; ii < N; ii++)
c[ii] = (a[ii] + a[ii] + a[ii] + a[ii]) / a[ii];
-#pragma acc parallel async (3) wait (1)
+#pragma acc parallel async (3) wait (1) present (a[0:N], d[0:N], N)
for (int ii = 0; ii < N; ii++)
d[ii] = ((a[ii] * a[ii] + a[ii]) / a[ii]) - a[ii];
-#pragma acc parallel wait (1) async (4)
+#pragma acc parallel wait (1) async (4) present (a[0:N], b[0:N], c[0:N]) \
+ present (d[0:N], e[0:N], N)
for (int ii = 0; ii < N; ii++)
e[ii] = a[ii] + b[ii] + c[ii] + d[ii];
-#pragma acc exit data copyout (a[0:N]) copyout (b[0:N]) copyout (c[0:N]) copyout (d[0:N]) copyout (e[0:N]) wait (1, 2, 3, 4) async (1)
+#pragma acc exit data copyout (a[0:N]) copyout (b[0:N]) copyout (c[0:N]) \
+ copyout (d[0:N]) copyout (e[0:N]) wait (1, 2, 3, 4) async (1) delete (N)
#pragma acc wait (1)
-
for (i = 0; i < N; i++)
{
if (a[i] != 2.0)
@@ -158,5 +164,11 @@ main (int argc, char **argv)
abort ();
}
+ free (a);
+ free (b);
+ free (c);
+ free (d);
+ free (e);
+
return 0;
}
@@ -1,3 +1,5 @@
+/* Test "acc enter data", "acc exit data" and "acc update". */
+
/* { dg-do run } */
#include <stdlib.h>
@@ -25,7 +27,7 @@ main (int argc, char **argv)
}
#pragma acc enter data copyin (a[0:N]) copyin (b[0:N]) copyin (N) async
-#pragma acc parallel async wait
+#pragma acc parallel async wait present (a[0:N], b[0:N], N)
#pragma acc loop
for (i = 0; i < N; i++)
b[i] = a[i];
@@ -49,7 +51,7 @@ main (int argc, char **argv)
}
#pragma acc update device (a[0:N], b[0:N]) async (1)
-#pragma acc parallel async (1)
+#pragma acc parallel async (1) present (a[0:N], b[0:N], N)
#pragma acc loop
for (i = 0; i < N; i++)
b[i] = a[i];
@@ -78,17 +80,17 @@ main (int argc, char **argv)
#pragma acc update device (b[0:N]) async (2)
#pragma acc enter data copyin (c[0:N], d[0:N]) async (3)
-#pragma acc parallel async (1) wait (1,2)
+#pragma acc parallel async (1) wait (1,2) present (a[0:N], b[0:N], N)
#pragma acc loop
for (i = 0; i < N; i++)
b[i] = (a[i] * a[i] * a[i]) / a[i];
-#pragma acc parallel async (2) wait (1,3)
+#pragma acc parallel async (2) wait (1,3) present (a[0:N], c[0:N], N)
#pragma acc loop
for (i = 0; i < N; i++)
c[i] = (a[i] + a[i] + a[i] + a[i]) / a[i];
-#pragma acc parallel async (3) wait (1,3)
+#pragma acc parallel async (3) wait (1,3) present (a[0:N], d[0:N], N)
#pragma acc loop
for (i = 0; i < N; i++)
d[i] = ((a[i] * a[i] + a[i]) / a[i]) - a[i];
@@ -123,27 +125,28 @@ main (int argc, char **argv)
#pragma acc update device (a[0:N], b[0:N], c[0:N], d[0:N]) async (1)
#pragma acc enter data copyin (e[0:N]) async (5)
-#pragma acc parallel async (1) wait (1)
+#pragma acc parallel async (1) wait (1) present (a[0:N], b[0:N], N)
for (int ii = 0; ii < N; ii++)
b[ii] = (a[ii] * a[ii] * a[ii]) / a[ii];
-#pragma acc parallel async (2) wait (1)
+#pragma acc parallel async (2) wait (1) present (a[0:N], c[0:N], N)
for (int ii = 0; ii < N; ii++)
c[ii] = (a[ii] + a[ii] + a[ii] + a[ii]) / a[ii];
-#pragma acc parallel async (3) wait (1)
+#pragma acc parallel async (3) wait (1) present (a[0:N], d[0:N], N)
for (int ii = 0; ii < N; ii++)
d[ii] = ((a[ii] * a[ii] + a[ii]) / a[ii]) - a[ii];
-#pragma acc parallel wait (1,5) async (4)
+#pragma acc parallel wait (1,5) async (4) present (a[0:N], b[0:N], c[0:N]) \
+ present (d[0:N], e[0:N], N)
for (int ii = 0; ii < N; ii++)
e[ii] = a[ii] + b[ii] + c[ii] + d[ii];
-#pragma acc exit data copyout (a[0:N]) copyout (b[0:N]) copyout (c[0:N]) copyout (d[0:N]) copyout (e[0:N]) wait (1, 2, 3, 4) async (1)
+#pragma acc exit data copyout (a[0:N]) copyout (b[0:N]) copyout (c[0:N]) \
+ copyout (d[0:N]) copyout (e[0:N]) wait (1, 2, 3, 4) async (1)
#pragma acc exit data delete (N)
#pragma acc wait (1)
-
for (i = 0; i < N; i++)
{
if (a[i] != 2.0)
@@ -162,5 +165,11 @@ main (int argc, char **argv)
abort ();
}
+ free (a);
+ free (b);
+ free (c);
+ free (d);
+ free (e);
+
return 0;
}
@@ -457,7 +457,7 @@ main(int argc, char **argv)
#pragma acc data copyin(a[0:N]) copyout(b[0:N]) if(1)
{
-#pragma acc parallel present(a[0:N])
+#pragma acc parallel present(a[0:N], b[0:N])
{
int ii;
@@ -1,3 +1,5 @@
+/* Test implicit data clauses inside data regions. */
+
/* { dg-do run } */
/* { dg-skip-if "" { *-*-* } { "*" } { "-DACC_MEM_SHARED=0" } } */
@@ -11,20 +13,16 @@ int
main (int argc, char **argv)
{
int N = 8;
- float *a, *b, *c, *d;
+ float a[N], b[N], c[N], *d;
int i;
- a = (float *) malloc (N * sizeof (float));
- b = (float *) malloc (N * sizeof (float));
- c = (float *) malloc (N * sizeof (float));
-
for (i = 0; i < N; i++)
{
a[i] = 3.0;
b[i] = 0.0;
}
-#pragma acc data copyin (a[0:N]) copyout (b[0:N])
+#pragma acc data copyin (a) copyout (b)
{
#pragma acc parallel
{
@@ -53,7 +51,7 @@ main (int argc, char **argv)
b[i] = 1.0;
}
-#pragma acc data copyin (a[0:N]) copyout (b[0:N])
+#pragma acc data copyin (a) copyout (b)
{
#pragma acc parallel
{
@@ -89,7 +87,7 @@ main (int argc, char **argv)
a[i] = 9.0;
}
-#pragma acc data present_or_copyin (a[0:N]) copyout (b[0:N])
+#pragma acc data present_or_copyin (a) copyout (b)
{
#pragma acc parallel
{
@@ -120,7 +118,7 @@ main (int argc, char **argv)
b[i] = 0.0;
}
-#pragma acc data copyin (a[0:N]) present_or_copyout (b[0:N])
+#pragma acc data copyin (a) present_or_copyout (b)
{
#pragma acc parallel
{
@@ -151,7 +149,7 @@ main (int argc, char **argv)
d = (float *) acc_copyin (&b[0], N * sizeof (float));
-#pragma acc data copyin (a[0:N]) present_or_copyout (b[0:N])
+#pragma acc data copyin (a) present_or_copyout (b)
{
#pragma acc parallel
{
@@ -188,7 +186,7 @@ main (int argc, char **argv)
b[i] = 4.0;
}
-#pragma acc data copy (a[0:N]) copyout (b[0:N])
+#pragma acc data copy (a) copyout (b)
{
#pragma acc parallel
{
@@ -223,7 +221,7 @@ main (int argc, char **argv)
b[i] = 7.0;
}
-#pragma acc data present_or_copy (a[0:N]) present_or_copy (b[0:N])
+#pragma acc data present_or_copy (a) present_or_copy (b)
{
#pragma acc parallel
{
@@ -261,7 +259,7 @@ main (int argc, char **argv)
d = (float *) acc_copyin (&a[0], N * sizeof (float));
d = (float *) acc_copyin (&b[0], N * sizeof (float));
-#pragma acc data present_or_copy (a[0:N]) present_or_copy (b[0:N])
+#pragma acc data present_or_copy (a) present_or_copy (b)
{
#pragma acc parallel
{
@@ -305,7 +303,7 @@ main (int argc, char **argv)
b[i] = 7.0;
}
-#pragma acc data copyin (a[0:N]) create (c[0:N]) copyout (b[0:N])
+#pragma acc data copyin (a) create (c) copyout (b)
{
#pragma acc parallel
{
@@ -343,7 +341,7 @@ main (int argc, char **argv)
b[i] = 8.0;
}
-#pragma acc data copyin (a[0:N]) present_or_create (c[0:N]) copyout (b[0:N])
+#pragma acc data copyin (a) present_or_create (c) copyout (b)
{
#pragma acc parallel
{
@@ -384,7 +382,7 @@ main (int argc, char **argv)
d = (float *) acc_malloc (N * sizeof (float));
acc_map_data (c, d, N * sizeof (float));
-#pragma acc data copyin (a[0:N]) present_or_create (c[0:N]) copyout (b[0:N])
+#pragma acc data copyin (a) present_or_create (c) copyout (b)
{
#pragma acc parallel
{
@@ -431,7 +429,7 @@ main (int argc, char **argv)
d = (float *) acc_malloc (N * sizeof (float));
acc_map_data (c, d, N * sizeof (float));
-#pragma acc data copyin (a[0:N]) present (c[0:N]) copyout (b[0:N])
+#pragma acc data copyin (a) present (c) copyout (b)
{
#pragma acc parallel
{
@@ -488,7 +486,7 @@ main (int argc, char **argv)
if (!acc_is_present (a, (N * sizeof (float))))
abort ();
-#pragma acc data present (a[0:N]) present (c[0:N]) present (b[0:N])
+#pragma acc data present (a) present (c) present (b)
{
#pragma acc parallel
{
@@ -543,7 +541,7 @@ main (int argc, char **argv)
d = (float *) acc_malloc (N * sizeof (float));
-#pragma acc parallel copyin (a[0:N]) deviceptr (d) copyout (b[0:N])
+#pragma acc parallel copyin (a) deviceptr (d) copyout (b)
{
int ii;
@@ -584,7 +582,7 @@ main (int argc, char **argv)
a[i] = 9.0;
}
-#pragma acc data pcopyin (a[0:N]) copyout (b[0:N])
+#pragma acc data pcopyin (a) copyout (b)
{
#pragma acc parallel
{
@@ -615,7 +613,7 @@ main (int argc, char **argv)
b[i] = 0.0;
}
-#pragma acc data copyin (a[0:N]) pcopyout (b[0:N])
+#pragma acc data copyin (a) pcopyout (b)
{
#pragma acc parallel
{
@@ -644,7 +642,7 @@ main (int argc, char **argv)
b[i] = 7.0;
}
-#pragma acc data copyin (a[0:N]) pcreate (c[0:N]) copyout (b[0:N])
+#pragma acc data copyin (a) pcreate (c) copyout (b)
{
#pragma acc parallel
{
@@ -1,3 +1,5 @@
+/* Test both explicit and implicitly present data. */
+
/* { dg-do run } */
/* { dg-skip-if "" { *-*-* } { "*" } { "-DACC_MEM_SHARED=0" } } */
@@ -8,11 +10,10 @@ int
main (int argc, char **argv)
{
int N = 8;
- float *a, *b;
+ float *a, b[N];
int i;
a = (float *) malloc (N * sizeof (float));
- b = (float *) malloc (N * sizeof (float));
for (i = 0; i < N; i++)
{
@@ -20,7 +21,7 @@ main (int argc, char **argv)
b[i] = 0.0;
}
-#pragma acc data copyin(a[0:N]) copyout(b[0:N])
+#pragma acc data copyin(a[0:N]) copy (b)
{
#pragma acc parallel present(a[0:N])
@@ -44,5 +45,7 @@ main (int argc, char **argv)
abort ();
}
+ free (a);
+
return 0;
}