This is the mail archive of the
gcc-patches@gcc.gnu.org
mailing list for the GCC project.
[patch] fix openacc data clause errors in c/c++ for PR70688
- From: Cesar Philippidis <cesar at codesourcery dot com>
- To: "gcc-patches at gcc dot gnu dot org" <gcc-patches at gcc dot gnu dot org>, Jakub Jelinek <jakub at redhat dot com>
- Date: Fri, 22 Apr 2016 07:07:20 -0700
- Subject: [patch] fix openacc data clause errors in c/c++ for PR70688
- Authentication-results: sourceware.org; auth=none
This patch makes the c and c++ FEs aware that acc data clauses are not
omp maps. I've also taught those FEs that reduction clauses are not data
clauses in openacc, which fixes PR70688.
I don't really like how *finish_omp_clauses has default arguments for
declare_simd, is_cilk and not is_oacc, but I went with that because it
seems like the thing to do here. Maybe it would be better if all of
those arguments were consolidated into a single bitmask or enum?
Is this OK for trunk and possibly gcc-6.2?
Cesar
2016-04-21 Cesar Philippidis <cesar@codesourcery.com>
gcc/c/
PR c/70688
* c-parser.c (c_parser_oacc_all_clauses): Update call to
c_finish_omp_clauses.
(c_parser_oacc_cache): Likewise.
(c_parser_oacc_loop): Likewise.
* c-tree.h (c_finish_omp_clauses): Add bool argument.
* c-typeck.c (c_finish_omp_clauses): Add bool is_oacc argument.
Use it to report OpenACC-specific diagnostics.
gcc/cp/
* cp-tree.h (finish_omp_clauses): Add bool argument.
* parser.c (cp_parser_oacc_all_clauses): Update call to
finish_omp_clauses.
(cp_parser_oacc_cache): Likewise.
(cp_parser_oacc_loop): Likewise.
* pt.c (tsubst_attribute): Add bool is_oacc arguments. Propagate it
to finish_omp_clauses.
(tsubst_omp_clauses): Update calls to tsubst_omp_clauses.
(tsubst_expr): Update calls to tsubst_omp_clauses.
* semantics.c (finish_omp_clauses): Add bool is_oacc argument.
Use it to report OpenACC-specific diagnostics.
gcc/testsuite/
* c-c++-common/goacc/data-clause-duplicate-1.c: Update expected errors.
* c-c++-common/goacc/deviceptr-1.c: Likewise.
* c-c++-common/goacc/kernels-alias-ipa-pta-3.c: Likewise.
* c-c++-common/goacc/reduction-5.c: Likewise.
* c-c++-common/goacc/pr70688.c: New test.
* g++.dg/goacc/data-1.C: New test.
* g++.dg/goacc/data-2.C: New test.
* g++.dg/gomp/template-data.C: New test.
diff --git a/gcc/c/c-parser.c b/gcc/c/c-parser.c
index bdd669d..11ba146 100644
--- a/gcc/c/c-parser.c
+++ b/gcc/c/c-parser.c
@@ -13177,7 +13177,7 @@ c_parser_oacc_all_clauses (c_parser *parser, omp_clause_mask mask,
c_parser_skip_to_pragma_eol (parser);
if (finish_p)
- return c_finish_omp_clauses (clauses, false);
+ return c_finish_omp_clauses (clauses, false, false, false, true);
return clauses;
}
@@ -13497,7 +13497,7 @@ c_parser_oacc_cache (location_t loc, c_parser *parser)
tree stmt, clauses;
clauses = c_parser_omp_var_list_parens (parser, OMP_CLAUSE__CACHE_, NULL);
- clauses = c_finish_omp_clauses (clauses, false);
+ clauses = c_finish_omp_clauses (clauses, false, false, false, true);
c_parser_skip_to_pragma_eol (parser);
@@ -13831,9 +13831,9 @@ c_parser_oacc_loop (location_t loc, c_parser *parser, char *p_name,
{
clauses = c_oacc_split_loop_clauses (clauses, cclauses);
if (*cclauses)
- *cclauses = c_finish_omp_clauses (*cclauses, false);
+ *cclauses = c_finish_omp_clauses (*cclauses, false, false, false, true);
if (clauses)
- clauses = c_finish_omp_clauses (clauses, false);
+ clauses = c_finish_omp_clauses (clauses, false, false, false, true);
}
tree block = c_begin_compound_stmt (true);
diff --git a/gcc/c/c-tree.h b/gcc/c/c-tree.h
index 4633182..0fa8a69 100644
--- a/gcc/c/c-tree.h
+++ b/gcc/c/c-tree.h
@@ -661,7 +661,8 @@ extern tree c_begin_omp_task (void);
extern tree c_finish_omp_task (location_t, tree, tree);
extern void c_finish_omp_cancel (location_t, tree);
extern void c_finish_omp_cancellation_point (location_t, tree);
-extern tree c_finish_omp_clauses (tree, bool, bool = false, bool = false);
+extern tree c_finish_omp_clauses (tree, bool, bool = false, bool = false,
+ bool = false);
extern tree c_build_va_arg (location_t, tree, location_t, tree);
extern tree c_finish_transaction (location_t, tree, int);
extern bool c_tree_equal (tree, tree);
diff --git a/gcc/c/c-typeck.c b/gcc/c/c-typeck.c
index 58c2139..6731fd0 100644
--- a/gcc/c/c-typeck.c
+++ b/gcc/c/c-typeck.c
@@ -12497,10 +12497,10 @@ c_find_omp_placeholder_r (tree *tp, int *, void *data)
tree
c_finish_omp_clauses (tree clauses, bool is_omp, bool declare_simd,
- bool is_cilk)
+ bool is_cilk, bool is_oacc)
{
bitmap_head generic_head, firstprivate_head, lastprivate_head;
- bitmap_head aligned_head, map_head, map_field_head;
+ bitmap_head aligned_head, map_head, map_field_head, oacc_reduction_head;
tree c, t, type, *pc;
tree simdlen = NULL_TREE, safelen = NULL_TREE;
bool branch_seen = false;
@@ -12517,6 +12517,7 @@ c_finish_omp_clauses (tree clauses, bool is_omp, bool declare_simd,
bitmap_initialize (&aligned_head, &bitmap_default_obstack);
bitmap_initialize (&map_head, &bitmap_default_obstack);
bitmap_initialize (&map_field_head, &bitmap_default_obstack);
+ bitmap_initialize (&oacc_reduction_head, &bitmap_default_obstack);
for (pc = &clauses, c = clauses; c ; c = *pc)
{
@@ -12854,6 +12855,16 @@ c_finish_omp_clauses (tree clauses, bool is_omp, bool declare_simd,
omp_clause_code_name[OMP_CLAUSE_CODE (c)]);
remove = true;
}
+ else if (is_oacc && OMP_CLAUSE_CODE (c) == OMP_CLAUSE_REDUCTION)
+ {
+ if (bitmap_bit_p (&oacc_reduction_head, DECL_UID (t)))
+ {
+ error ("%qD appears more than once in reduction clauses", t);
+ remove = true;
+ }
+ else
+ bitmap_set_bit (&oacc_reduction_head, DECL_UID (t));
+ }
else if (bitmap_bit_p (&generic_head, DECL_UID (t))
|| bitmap_bit_p (&firstprivate_head, DECL_UID (t))
|| bitmap_bit_p (&lastprivate_head, DECL_UID (t)))
@@ -13145,8 +13156,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))
diff --git a/gcc/cp/cp-tree.h b/gcc/cp/cp-tree.h
index ec92718..e28a16a 100644
--- a/gcc/cp/cp-tree.h
+++ b/gcc/cp/cp-tree.h
@@ -6397,7 +6397,7 @@ 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,
- bool = false);
+ bool = false, bool = false);
extern tree push_omp_privatization_clauses (bool);
extern void pop_omp_privatization_clauses (tree);
extern void save_omp_privatization_clauses (vec<tree> &);
diff --git a/gcc/cp/parser.c b/gcc/cp/parser.c
index feb8de7..d46610e 100644
--- a/gcc/cp/parser.c
+++ b/gcc/cp/parser.c
@@ -32259,7 +32259,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, false, false, false, true);
return clauses;
}
@@ -35090,7 +35090,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, false, false, false, true);
cp_parser_require_pragma_eol (parser, cp_lexer_peek_token (parser->lexer));
@@ -35415,9 +35415,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)
- *cclauses = finish_omp_clauses (*cclauses, false);
+ *cclauses = finish_omp_clauses (*cclauses, false, false, false, true);
if (clauses)
- clauses = finish_omp_clauses (clauses, false);
+ clauses = finish_omp_clauses (clauses, false, false, false, true);
}
tree block = begin_omp_structured_block ();
diff --git a/gcc/cp/pt.c b/gcc/cp/pt.c
index e18422f..1c14ed6 100644
--- a/gcc/cp/pt.c
+++ b/gcc/cp/pt.c
@@ -9563,7 +9563,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. */
@@ -9583,7 +9584,7 @@ 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);
tree parms = DECL_ARGUMENTS (*decl_p);
@@ -14536,7 +14537,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;
@@ -14749,7 +14751,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, allow_fields, false, false, is_oacc);
if (linear_no_step)
for (nc = new_clauses; nc; nc = OMP_CLAUSE_CHAIN (nc))
if (nc == linear_no_step)
@@ -15452,8 +15454,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);
@@ -15462,7 +15464,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 ();
@@ -15476,7 +15478,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);
@@ -15498,9 +15500,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)));
@@ -15557,7 +15559,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);
@@ -15572,9 +15574,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 ();
@@ -15619,8 +15620,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;
@@ -15629,7 +15630,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);
@@ -15638,8 +15639,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);
@@ -15647,7 +15648,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);
diff --git a/gcc/cp/semantics.c b/gcc/cp/semantics.c
index 2365a73..2f5e305 100644
--- a/gcc/cp/semantics.c
+++ b/gcc/cp/semantics.c
@@ -5794,10 +5794,10 @@ cp_finish_omp_clause_depend_sink (tree sink_clause)
tree
finish_omp_clauses (tree clauses, bool allow_fields, bool declare_simd,
- bool is_cilk)
+ bool is_cilk, bool is_oacc)
{
bitmap_head generic_head, firstprivate_head, lastprivate_head;
- bitmap_head aligned_head, map_head, map_field_head;
+ bitmap_head aligned_head, map_head, map_field_head, oacc_reduction_head;
tree c, t, *pc;
tree safelen = NULL_TREE;
bool branch_seen = false;
@@ -5811,6 +5811,7 @@ finish_omp_clauses (tree clauses, bool allow_fields, bool declare_simd,
bitmap_initialize (&aligned_head, &bitmap_default_obstack);
bitmap_initialize (&map_head, &bitmap_default_obstack);
bitmap_initialize (&map_field_head, &bitmap_default_obstack);
+ bitmap_initialize (&oacc_reduction_head, &bitmap_default_obstack);
for (pc = &clauses, c = clauses; c ; c = *pc)
{
@@ -6040,6 +6041,16 @@ finish_omp_clauses (tree clauses, bool allow_fields, bool declare_simd,
omp_clause_code_name[OMP_CLAUSE_CODE (c)]);
remove = true;
}
+ else if (is_oacc && OMP_CLAUSE_CODE (c) == OMP_CLAUSE_REDUCTION)
+ {
+ if (bitmap_bit_p (&oacc_reduction_head, DECL_UID (t)))
+ {
+ error ("%qD appears more than once in reduction clauses", t);
+ remove = true;
+ }
+ else
+ bitmap_set_bit (&oacc_reduction_head, DECL_UID (t));
+ }
else if (bitmap_bit_p (&generic_head, DECL_UID (t))
|| bitmap_bit_p (&firstprivate_head, DECL_UID (t))
|| bitmap_bit_p (&lastprivate_head, DECL_UID (t)))
@@ -6050,7 +6061,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
@@ -6076,7 +6090,7 @@ 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");
@@ -6102,7 +6116,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
@@ -6612,6 +6629,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);
@@ -6623,6 +6643,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;
}
@@ -6699,7 +6740,7 @@ 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");
@@ -6748,7 +6789,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
@@ -6758,6 +6802,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;
@@ -6765,7 +6811,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
diff --git a/gcc/testsuite/c-c++-common/goacc/data-clause-duplicate-1.c b/gcc/testsuite/c-c++-common/goacc/data-clause-duplicate-1.c
index 7a1cf68..6245beb 100644
--- a/gcc/testsuite/c-c++-common/goacc/data-clause-duplicate-1.c
+++ b/gcc/testsuite/c-c++-common/goacc/data-clause-duplicate-1.c
@@ -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" } */
;
}
diff --git a/gcc/testsuite/c-c++-common/goacc/deviceptr-1.c b/gcc/testsuite/c-c++-common/goacc/deviceptr-1.c
index 08ddb10..3aa0e8a 100644
--- a/gcc/testsuite/c-c++-common/goacc/deviceptr-1.c
+++ b/gcc/testsuite/c-c++-common/goacc/deviceptr-1.c
@@ -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" } */
;
}
diff --git a/gcc/testsuite/c-c++-common/goacc/kernels-alias-ipa-pta-3.c b/gcc/testsuite/c-c++-common/goacc/kernels-alias-ipa-pta-3.c
index 1eb56eb..97d8f52 100644
--- a/gcc/testsuite/c-c++-common/goacc/kernels-alias-ipa-pta-3.c
+++ b/gcc/testsuite/c-c++-common/goacc/kernels-alias-ipa-pta-3.c
@@ -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 } } } */
diff --git a/gcc/testsuite/c-c++-common/goacc/pr70688.c b/gcc/testsuite/c-c++-common/goacc/pr70688.c
new file mode 100644
index 0000000..5a23665
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/goacc/pr70688.c
@@ -0,0 +1,48 @@
+const int n = 100;
+
+int
+private_reduction ()
+{
+ int i, r;
+
+ #pragma acc parallel
+ #pragma acc loop private (r) reduction (+:r)
+ for (i = 0; i < 100; i++)
+ r += 10;
+
+ return r;
+}
+
+int
+parallel_reduction ()
+{
+ int sum = 0;
+ int dummy = 0;
+
+#pragma acc data copy (dummy)
+ {
+#pragma acc parallel num_gangs (10) copy (sum) reduction (+:sum)
+ {
+ int v = 5;
+ sum += 10 + v;
+ }
+ }
+
+ return sum;
+}
+
+int
+main ()
+{
+ int i, s = 0;
+
+#pragma acc parallel num_gangs (10) copy (s) reduction (+:s)
+ for (i = 0; i < n; i++)
+ s += i+1;
+
+#pragma acc parallel num_gangs (10) reduction (+:s) copy (s)
+ for (i = 0; i < n; i++)
+ s += i+1;
+
+ return 0;
+}
diff --git a/gcc/testsuite/c-c++-common/goacc/reduction-5.c b/gcc/testsuite/c-c++-common/goacc/reduction-5.c
index 74daad3..dfdbab9 100644
--- a/gcc/testsuite/c-c++-common/goacc/reduction-5.c
+++ b/gcc/testsuite/c-c++-common/goacc/reduction-5.c
@@ -7,9 +7,9 @@ main(void)
{
int v1;
-#pragma acc parallel reduction(+:v1) private(v1) /* { dg-error "appears more than once in data clauses" } */
+#pragma acc parallel reduction(+:v1) private(v1) /* { dg-error "invalid private reduction" } */
;
-#pragma acc parallel reduction(+:v1) firstprivate(v1) /* { dg-error "appears more than once in data clauses" } */
+#pragma acc parallel reduction(+:v1) firstprivate(v1) /* { dg-error "invalid private reduction" } */
;
return 0;
diff --git a/gcc/testsuite/g++.dg/goacc/data-1.C b/gcc/testsuite/g++.dg/goacc/data-1.C
new file mode 100644
index 0000000..54676dc
--- /dev/null
+++ b/gcc/testsuite/g++.dg/goacc/data-1.C
@@ -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 } */
diff --git a/gcc/testsuite/g++.dg/goacc/data-2.C b/gcc/testsuite/g++.dg/goacc/data-2.C
new file mode 100644
index 0000000..efa002d
--- /dev/null
+++ b/gcc/testsuite/g++.dg/goacc/data-2.C
@@ -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" } */
+ ;
+}
diff --git a/gcc/testsuite/g++.dg/gomp/template-data.C b/gcc/testsuite/g++.dg/gomp/template-data.C
new file mode 100644
index 0000000..0be14d4
--- /dev/null
+++ b/gcc/testsuite/g++.dg/gomp/template-data.C
@@ -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" } */
+ ;
+}