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