This patch fixes some cases for OpenACC and OpenMP where map clauses were being expanded (adding firstprivate_pointer, attach/detach nodes, and so forth) unnecessarily, after the "OpenMP/OpenACC: Rework clause expansion and nested struct handling" patch (approved but not yet committed):
https://gcc.gnu.org/pipermail/gcc-patches/2022-October/603792.html This is done by introducing a C_ORT_ACC_TARGET region type for OpenACC compute regions to help distinguish them from non-compute regions that need different handling, and by passing the region type through to the clause expansion functions. The patch also fixes clause expansion for OpenMP TO/FROM clauses, which need to dereference references but not have any additional mapping nodes. (These cases showed up due to the gimplification changes in the C++ "declare mapper" patch, but logically belong next to the earlier patch named above.) 2022-11-30 Julian Brown <jul...@codesourcery.com> gcc/ * c-family/c-common.h (c_omp_region_type): Add C_ORT_ACC_TARGET. (c_omp_address_inspector): Pass c_omp_region_type instead of "target" bool. * c-family/c-omp.cc (c_omp_address_inspector::expand_array_base): Adjust clause expansion for OpenACC and non-map (OpenMP to/from) clauses. (c_omp_address_inspector::expand_component_selector): Use c_omp_region_type parameter. Don't expand OpenMP to/from clauses. (c_omp_address_inspector::expand_map_clause): Take ORT parameter, pass to expand_array_base, etc. gcc/c/ * c-parser.cc (c_parser_oacc_all_clauses): Add TARGET parameter. Use to select region type for c_finish_omp_clauses call. (c_parser_oacc_loop): Update calls to c_parser_oacc_all_clauses. (c_parser_oacc_compute): Likewise. * c-typeck.cc (handle_omp_array_sctions_1): Update for C_ORT_ACC_TARGET addition and ai.expand_map_clause signature change. (c_finish_omp_clauses): Likewise. gcc/cp/ * parser.cc (cp_parser_oacc_all_clauses): Add TARGET parameter. Use to select region type for finish_omp_clauses call. (cp_parser_oacc_declare): Update call to cp_parser_oacc_all_clauses. (cp_parser_oacc_loop): Update calls to cp_parser_oacc_all_clauses. (cp_parser_oacc_compute): Likewise. * pt.cc (tsubst_expr): Use C_ORT_ACC_TARGET for call to tsubst_omp_clauses for compute regions. * semantics.cc (handle_omp_array_sections_1): Update for C_ORT_ACC_TARGET addition and ai.expand_map_clause signature change. (finish_omp_clauses): Likewise. --- gcc/c-family/c-common.h | 10 +++-- gcc/c-family/c-omp.cc | 90 ++++++++++++++++++++++++++++++++++++----- gcc/c/c-parser.cc | 15 ++++--- gcc/c/c-typeck.cc | 39 ++++++++---------- gcc/cp/parser.cc | 15 ++++--- gcc/cp/pt.cc | 4 +- gcc/cp/semantics.cc | 47 ++++++++++----------- 7 files changed, 144 insertions(+), 76 deletions(-) diff --git a/gcc/c-family/c-common.h b/gcc/c-family/c-common.h index d935d4b3d7d9..06674e769bd4 100644 --- a/gcc/c-family/c-common.h +++ b/gcc/c-family/c-common.h @@ -1245,7 +1245,8 @@ enum c_omp_region_type C_ORT_DECLARE_SIMD = 1 << 2, C_ORT_TARGET = 1 << 3, C_ORT_OMP_DECLARE_SIMD = C_ORT_OMP | C_ORT_DECLARE_SIMD, - C_ORT_OMP_TARGET = C_ORT_OMP | C_ORT_TARGET + C_ORT_OMP_TARGET = C_ORT_OMP | C_ORT_TARGET, + C_ORT_ACC_TARGET = C_ORT_ACC | C_ORT_TARGET }; extern tree c_finish_omp_master (location_t, tree); @@ -1345,10 +1346,11 @@ public: bool maybe_zero_length_array_section (tree); tree expand_array_base (tree, vec<omp_addr_token *> &, tree, unsigned *, - bool, bool); + c_omp_region_type, bool); tree expand_component_selector (tree, vec<omp_addr_token *> &, tree, - unsigned *, bool); - tree expand_map_clause (tree, tree, vec<omp_addr_token *> &, bool); + unsigned *, c_omp_region_type); + tree expand_map_clause (tree, tree, vec<omp_addr_token *> &, + c_omp_region_type); }; enum c_omp_directive_kind { diff --git a/gcc/c-family/c-omp.cc b/gcc/c-family/c-omp.cc index d32c2a977304..74c01d8f2a52 100644 --- a/gcc/c-family/c-omp.cc +++ b/gcc/c-family/c-omp.cc @@ -3370,7 +3370,8 @@ tree c_omp_address_inspector::expand_array_base (tree c, vec<omp_addr_token *> &addr_tokens, tree expr, unsigned *idx, - bool target, bool decl_p) + c_omp_region_type ort, + bool decl_p) { using namespace omp_addr_tokenizer; location_t loc = OMP_CLAUSE_LOCATION (c); @@ -3380,14 +3381,26 @@ c_omp_address_inspector::expand_array_base (tree c, && is_global_var (decl) && lookup_attribute ("omp declare target", DECL_ATTRIBUTES (decl))); + bool map_p = OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP; bool implicit_p = (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP && OMP_CLAUSE_MAP_IMPLICIT (c)); bool chain_p = omp_access_chain_p (addr_tokens, i + 1); tree c2 = NULL_TREE, c3 = NULL_TREE; unsigned consume_tokens = 2; + bool target = (ort & C_ORT_TARGET) != 0; + bool openmp = (ort & C_ORT_OMP) != 0; gcc_assert (i == 0); + if (!openmp + && OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP + && (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ATTACH + || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_DETACH)) + { + *idx = ++i; + return c; + } + switch (addr_tokens[i + 1]->u.access_kind) { case ACCESS_DIRECT: @@ -3397,11 +3410,19 @@ c_omp_address_inspector::expand_array_base (tree c, case ACCESS_REF: { - /* Copy the referenced object. */ + /* Copy the referenced object. Note that we do this even for !MAP_P + clauses. */ tree obj = convert_from_reference (addr_tokens[i + 1]->expr); OMP_CLAUSE_DECL (c) = obj; OMP_CLAUSE_SIZE (c) = TYPE_SIZE_UNIT (TREE_TYPE (obj)); + if (!map_p) + { + if (decl_p) + c_common_mark_addressable_vec (addr_tokens[i + 1]->expr); + break; + } + /* If we have a reference to a pointer, avoid using FIRSTPRIVATE_REFERENCE here in case the pointer is modified in the offload region (we can only do that if the pointer does not point @@ -3441,6 +3462,13 @@ c_omp_address_inspector::expand_array_base (tree c, case ACCESS_INDEXED_REF_TO_ARRAY: { + if (!map_p) + { + if (decl_p) + c_common_mark_addressable_vec (addr_tokens[i + 1]->expr); + break; + } + tree virtual_origin = convert_from_reference (addr_tokens[i + 1]->expr); virtual_origin = build_fold_addr_expr (virtual_origin); @@ -3467,6 +3495,13 @@ c_omp_address_inspector::expand_array_base (tree c, case ACCESS_INDEXED_ARRAY: { + if (!map_p) + { + if (decl_p) + c_common_mark_addressable_vec (addr_tokens[i + 1]->expr); + break; + } + /* The code handling "firstprivatize_array_bases" in gimplify.cc is relevant here. What do we need to create for arrays at this stage? (This condition doesn't feel quite right. FIXME?) */ @@ -3501,6 +3536,13 @@ c_omp_address_inspector::expand_array_base (tree c, case ACCESS_POINTER: case ACCESS_POINTER_OFFSET: { + if (!map_p) + { + if (decl_p) + c_common_mark_addressable_vec (addr_tokens[i + 1]->expr); + break; + } + unsigned last_access = i + 1; tree virtual_origin; @@ -3524,7 +3566,11 @@ c_omp_address_inspector::expand_array_base (tree c, addr_tokens[last_access]->expr); tree data_addr = omp_accessed_addr (addr_tokens, last_access, expr); c2 = build_omp_clause (loc, OMP_CLAUSE_MAP); - if (decl_p && target && !chain_p && !declare_target_p) + /* For OpenACC, use FIRSTPRIVATE_POINTER for decls even on non-compute + regions (e.g. "acc data" constructs). It'll be removed anyway in + gimplify.cc, but doing it this way maintains diagnostic + behaviour. */ + if (decl_p && (target || !openmp) && !chain_p && !declare_target_p) OMP_CLAUSE_SET_MAP_KIND (c2, GOMP_MAP_FIRSTPRIVATE_POINTER); else { @@ -3544,6 +3590,13 @@ c_omp_address_inspector::expand_array_base (tree c, case ACCESS_REF_TO_POINTER: case ACCESS_REF_TO_POINTER_OFFSET: { + if (!map_p) + { + if (decl_p) + c_common_mark_addressable_vec (addr_tokens[i + 1]->expr); + break; + } + unsigned last_access = i + 1; tree virtual_origin; @@ -3618,7 +3671,7 @@ c_omp_address_inspector::expand_array_base (tree c, i += consume_tokens; *idx = i; - if (target && chain_p) + if (target && chain_p && map_p) return omp_expand_access_chain (c, expr, addr_tokens, idx); else if (chain_p) while (*idx < addr_tokens.length () @@ -3635,13 +3688,15 @@ c_omp_address_inspector::expand_component_selector (tree c, vec<omp_addr_token *> &addr_tokens, tree expr, unsigned *idx, - bool target) + c_omp_region_type ort) { using namespace omp_addr_tokenizer; location_t loc = OMP_CLAUSE_LOCATION (c); unsigned i = *idx; tree c2 = NULL_TREE, c3 = NULL_TREE; bool chain_p = omp_access_chain_p (addr_tokens, i + 1); + bool map_p = OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP; + bool target = (ort & C_ORT_TARGET) != 0; switch (addr_tokens[i + 1]->u.access_kind) { @@ -3651,11 +3706,15 @@ c_omp_address_inspector::expand_component_selector (tree c, case ACCESS_REF: { - /* Copy the referenced object. */ + /* Copy the referenced object. Note that we also do this for !MAP_P + clauses. */ tree obj = convert_from_reference (addr_tokens[i + 1]->expr); OMP_CLAUSE_DECL (c) = obj; OMP_CLAUSE_SIZE (c) = TYPE_SIZE_UNIT (TREE_TYPE (obj)); + if (!map_p) + break; + c2 = build_omp_clause (loc, OMP_CLAUSE_MAP); OMP_CLAUSE_SET_MAP_KIND (c2, GOMP_MAP_ATTACH_DETACH); OMP_CLAUSE_DECL (c2) = addr_tokens[i + 1]->expr; @@ -3665,6 +3724,9 @@ c_omp_address_inspector::expand_component_selector (tree c, case ACCESS_INDEXED_REF_TO_ARRAY: { + if (!map_p) + break; + tree virtual_origin = convert_from_reference (addr_tokens[i + 1]->expr); virtual_origin = build_fold_addr_expr (virtual_origin); @@ -3686,6 +3748,9 @@ c_omp_address_inspector::expand_component_selector (tree c, case ACCESS_POINTER: case ACCESS_POINTER_OFFSET: { + if (!map_p) + break; + tree virtual_origin = fold_convert_loc (loc, ptrdiff_type_node, addr_tokens[i + 1]->expr); @@ -3705,6 +3770,9 @@ c_omp_address_inspector::expand_component_selector (tree c, case ACCESS_REF_TO_POINTER: case ACCESS_REF_TO_POINTER_OFFSET: { + if (!map_p) + break; + tree ptr = convert_from_reference (addr_tokens[i + 1]->expr); tree virtual_origin = fold_convert_loc (loc, ptrdiff_type_node, ptr); @@ -3750,7 +3818,7 @@ c_omp_address_inspector::expand_component_selector (tree c, i += 2; *idx = i; - if (target && chain_p) + if (target && chain_p && map_p) return omp_expand_access_chain (c, expr, addr_tokens, idx); else if (chain_p) while (*idx < addr_tokens.length () @@ -3766,7 +3834,7 @@ c_omp_address_inspector::expand_component_selector (tree c, tree c_omp_address_inspector::expand_map_clause (tree c, tree expr, vec<omp_addr_token *> &addr_tokens, - bool target) + c_omp_region_type ort) { using namespace omp_addr_tokenizer; unsigned i, length = addr_tokens.length (); @@ -3780,7 +3848,7 @@ c_omp_address_inspector::expand_map_clause (tree c, tree expr, && addr_tokens[i]->u.structure_base_kind == BASE_DECL && addr_tokens[i + 1]->type == ACCESS_METHOD) { - c = expand_array_base (c, addr_tokens, expr, &i, target, true); + c = expand_array_base (c, addr_tokens, expr, &i, ort, true); if (c == error_mark_node) return error_mark_node; } @@ -3789,7 +3857,7 @@ c_omp_address_inspector::expand_map_clause (tree c, tree expr, && addr_tokens[i]->u.structure_base_kind == BASE_ARBITRARY_EXPR && addr_tokens[i + 1]->type == ACCESS_METHOD) { - c = expand_array_base (c, addr_tokens, expr, &i, target, false); + c = expand_array_base (c, addr_tokens, expr, &i, ort, false); if (c == error_mark_node) return error_mark_node; } @@ -3825,7 +3893,7 @@ c_omp_address_inspector::expand_map_clause (tree c, tree expr, && addr_tokens[i]->type == COMPONENT_SELECTOR && addr_tokens[i + 1]->type == ACCESS_METHOD) { - c = expand_component_selector (c, addr_tokens, expr, &i, target); + c = expand_component_selector (c, addr_tokens, expr, &i, ort); /* We used 'expr', so these must have been the last tokens. */ gcc_assert (i == length); if (c == error_mark_node) diff --git a/gcc/c/c-parser.cc b/gcc/c/c-parser.cc index 7d6960fffbb6..15e2fd7962af 100644 --- a/gcc/c/c-parser.cc +++ b/gcc/c/c-parser.cc @@ -17514,7 +17514,8 @@ c_parser_omp_clause_detach (c_parser *parser, tree list) static tree c_parser_oacc_all_clauses (c_parser *parser, omp_clause_mask mask, - const char *where, bool finish_p = true) + const char *where, bool finish_p = true, + bool target = false) { tree clauses = NULL; bool first = true; @@ -17715,7 +17716,8 @@ 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, C_ORT_ACC); + return c_finish_omp_clauses (clauses, target ? C_ORT_ACC_TARGET + : C_ORT_ACC); return clauses; } @@ -18440,12 +18442,13 @@ c_parser_oacc_loop (location_t loc, c_parser *parser, char *p_name, mask |= OACC_LOOP_CLAUSE_MASK; tree clauses = c_parser_oacc_all_clauses (parser, mask, p_name, - cclauses == NULL); + /*finish_p=*/cclauses == NULL, + /*target=*/is_parallel); if (cclauses) { clauses = c_oacc_split_loop_clauses (clauses, cclauses, is_parallel); if (*cclauses) - *cclauses = c_finish_omp_clauses (*cclauses, C_ORT_ACC); + *cclauses = c_finish_omp_clauses (*cclauses, C_ORT_ACC_TARGET); if (clauses) clauses = c_finish_omp_clauses (clauses, C_ORT_ACC); } @@ -18570,7 +18573,9 @@ c_parser_oacc_compute (location_t loc, c_parser *parser, } } - tree clauses = c_parser_oacc_all_clauses (parser, mask, p_name); + tree clauses = c_parser_oacc_all_clauses (parser, mask, p_name, + /*finish_p=*/true, + /*target=*/true); tree block = c_begin_omp_parallel (); add_stmt (c_parser_omp_structured_block (parser, if_p)); diff --git a/gcc/c/c-typeck.cc b/gcc/c/c-typeck.cc index ca32d8cd5d21..a4f3b855f6e7 100644 --- a/gcc/c/c-typeck.cc +++ b/gcc/c/c-typeck.cc @@ -13634,6 +13634,7 @@ handle_omp_array_sections_1 (tree c, tree t, vec<tree> &types, enum c_omp_region_type ort) { tree ret, low_bound, length, type; + bool openacc = (ort & C_ORT_ACC) != 0; if (TREE_CODE (t) != TREE_LIST) { if (error_operand_p (t)) @@ -13753,7 +13754,7 @@ handle_omp_array_sections_1 (tree c, tree t, vec<tree> &types, { error_at (OMP_CLAUSE_LOCATION (c), "expected single pointer in %qs clause", - user_omp_clause_code_name (c, ort == C_ORT_ACC)); + user_omp_clause_code_name (c, openacc)); return error_mark_node; } } @@ -14201,9 +14202,7 @@ handle_omp_array_sections (tree c, enum c_omp_region_type ort) c_omp_address_inspector ai (OMP_CLAUSE_LOCATION (c), t); - tree nc = ai.expand_map_clause (c, first, addr_tokens, - (ort == C_ORT_OMP_TARGET - || ort == C_ORT_ACC)); + tree nc = ai.expand_map_clause (c, first, addr_tokens, ort); if (nc != error_mark_node) { if (ai.maybe_zero_length_array_section (c)) @@ -14486,6 +14485,7 @@ c_finish_omp_clauses (tree clauses, enum c_omp_region_type ort) bool allocate_seen = false; bool implicit_moved = false; bool target_in_reduction_seen = false; + bool openacc = (ort & C_ORT_ACC) != 0; bitmap_obstack_initialize (NULL); bitmap_initialize (&generic_head, &bitmap_default_obstack); @@ -14501,7 +14501,7 @@ c_finish_omp_clauses (tree clauses, enum c_omp_region_type ort) bitmap_initialize (&oacc_reduction_head, &bitmap_default_obstack); bitmap_initialize (&is_on_device_head, &bitmap_default_obstack); - if (ort & C_ORT_ACC) + if (openacc) for (c = clauses; c; c = OMP_CLAUSE_CHAIN (c)) if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_ASYNC) { @@ -14895,8 +14895,7 @@ c_finish_omp_clauses (tree clauses, enum c_omp_region_type ort) omp_clause_code_name[OMP_CLAUSE_CODE (c)]); remove = true; } - else if ((ort == C_ORT_ACC - && OMP_CLAUSE_CODE (c) == OMP_CLAUSE_REDUCTION) + else if ((openacc && OMP_CLAUSE_CODE (c) == OMP_CLAUSE_REDUCTION) || (ort == C_ORT_OMP && (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_USE_DEVICE_PTR || (OMP_CLAUSE_CODE (c) @@ -14919,7 +14918,7 @@ c_finish_omp_clauses (tree clauses, enum c_omp_region_type ort) if (bitmap_bit_p (&oacc_reduction_head, DECL_UID (t))) { error_at (OMP_CLAUSE_LOCATION (c), - ort == C_ORT_ACC + openacc ? "%qD appears more than once in reduction clauses" : "%qD appears more than once in data clauses", t); @@ -14942,7 +14941,7 @@ c_finish_omp_clauses (tree clauses, enum c_omp_region_type ort) || OMP_CLAUSE_CODE (c) == OMP_CLAUSE_IS_DEVICE_PTR) && bitmap_bit_p (&map_head, DECL_UID (t))) { - if (ort == C_ORT_ACC) + if (openacc) error_at (OMP_CLAUSE_LOCATION (c), "%qD appears more than once in data clauses", t); else @@ -15010,7 +15009,7 @@ c_finish_omp_clauses (tree clauses, enum c_omp_region_type ort) else if (bitmap_bit_p (&map_head, DECL_UID (t)) || bitmap_bit_p (&map_field_head, DECL_UID (t))) { - if (ort == C_ORT_ACC) + if (openacc) error_at (OMP_CLAUSE_LOCATION (c), "%qD appears more than once in data clauses", t); else if (OMP_CLAUSE_FIRSTPRIVATE_IMPLICIT (c) @@ -15357,7 +15356,7 @@ c_finish_omp_clauses (tree clauses, enum c_omp_region_type ort) error_at (OMP_CLAUSE_LOCATION (c), "%qD appears more than once in motion " "clauses", rt); - else if (ort == C_ORT_ACC) + else if (openacc) error_at (OMP_CLAUSE_LOCATION (c), "%qD appears more than once in data " "clauses", rt); @@ -15445,8 +15444,7 @@ c_finish_omp_clauses (tree clauses, enum c_omp_region_type ort) to be written. */ if (addr_tokens[0]->type == STRUCTURE_BASE && (bitmap_bit_p (&map_field_head, DECL_UID (t)) - || (ort != C_ORT_ACC - && bitmap_bit_p (&map_head, DECL_UID (t))))) + || (!openacc && bitmap_bit_p (&map_head, DECL_UID (t))))) goto skip_decl_checks; if (!VAR_P (t) && TREE_CODE (t) != PARM_DECL) @@ -15512,7 +15510,7 @@ c_finish_omp_clauses (tree clauses, enum c_omp_region_type ort) } else if (bitmap_bit_p (&map_head, DECL_UID (t)) && !bitmap_bit_p (&map_field_head, DECL_UID (t)) - && ort == C_ORT_ACC) + && openacc) { error_at (OMP_CLAUSE_LOCATION (c), "%qD appears more than once in data clauses", t); @@ -15527,7 +15525,7 @@ c_finish_omp_clauses (tree clauses, enum c_omp_region_type ort) if (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_MAP) error_at (OMP_CLAUSE_LOCATION (c), "%qD appears more than once in motion clauses", t); - else if (ort == C_ORT_ACC) + else if (openacc) error_at (OMP_CLAUSE_LOCATION (c), "%qD appears more than once in data clauses", t); else @@ -15535,8 +15533,7 @@ c_finish_omp_clauses (tree clauses, enum c_omp_region_type ort) "%qD appears more than once in map clauses", t); remove = true; } - else if (ort == C_ORT_ACC - && bitmap_bit_p (&generic_head, DECL_UID (t))) + else if (openacc && bitmap_bit_p (&generic_head, DECL_UID (t))) { error_at (OMP_CLAUSE_LOCATION (c), "%qD appears more than once in data clauses", t); @@ -15545,7 +15542,7 @@ c_finish_omp_clauses (tree clauses, enum c_omp_region_type ort) else if (bitmap_bit_p (&firstprivate_head, DECL_UID (t)) || bitmap_bit_p (&is_on_device_head, DECL_UID (t))) { - if (ort == C_ORT_ACC) + if (openacc) error_at (OMP_CLAUSE_LOCATION (c), "%qD appears more than once in data clauses", t); else @@ -15578,9 +15575,7 @@ c_finish_omp_clauses (tree clauses, enum c_omp_region_type ort) grp_start_p = pc; grp_sentinel = OMP_CLAUSE_CHAIN (c); tree nc = ai.expand_map_clause (c, OMP_CLAUSE_DECL (c), - addr_tokens, - (ort == C_ORT_OMP_TARGET - || ort == C_ORT_ACC)); + addr_tokens, ort); if (nc != error_mark_node) c = nc; } @@ -15661,7 +15656,7 @@ c_finish_omp_clauses (tree clauses, enum c_omp_region_type ort) if (TREE_CODE (TREE_TYPE (t)) != POINTER_TYPE) { if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_USE_DEVICE_PTR - && ort != C_ORT_ACC) + && !openacc) { error_at (OMP_CLAUSE_LOCATION (c), "%qs variable is not a pointer", diff --git a/gcc/cp/parser.cc b/gcc/cp/parser.cc index bfd8aeae39f6..3905d561152b 100644 --- a/gcc/cp/parser.cc +++ b/gcc/cp/parser.cc @@ -40764,7 +40764,7 @@ cp_parser_oacc_clause_async (cp_parser *parser, tree list) static tree cp_parser_oacc_all_clauses (cp_parser *parser, omp_clause_mask mask, const char *where, cp_token *pragma_tok, - bool finish_p = true) + bool finish_p = true, bool target = false) { tree clauses = NULL; bool first = true; @@ -40974,7 +40974,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, C_ORT_ACC); + return finish_omp_clauses (clauses, target ? C_ORT_ACC_TARGET : C_ORT_ACC); return clauses; } @@ -45587,7 +45587,7 @@ cp_parser_oacc_declare (cp_parser *parser, cp_token *pragma_tok) bool found_in_scope = global_bindings_p (); clauses = cp_parser_oacc_all_clauses (parser, OACC_DECLARE_CLAUSE_MASK, - "#pragma acc declare", pragma_tok, true); + "#pragma acc declare", pragma_tok); if (omp_find_clause (clauses, OMP_CLAUSE_MAP) == NULL_TREE) @@ -45835,12 +45835,13 @@ cp_parser_oacc_loop (cp_parser *parser, cp_token *pragma_tok, char *p_name, mask |= OACC_LOOP_CLAUSE_MASK; tree clauses = cp_parser_oacc_all_clauses (parser, mask, p_name, pragma_tok, - cclauses == NULL); + /*finish_p=*/cclauses == NULL, + /*target=*/is_parallel); if (cclauses) { clauses = c_oacc_split_loop_clauses (clauses, cclauses, is_parallel); if (*cclauses) - *cclauses = finish_omp_clauses (*cclauses, C_ORT_ACC); + *cclauses = finish_omp_clauses (*cclauses, C_ORT_ACC_TARGET); if (clauses) clauses = finish_omp_clauses (clauses, C_ORT_ACC); } @@ -45965,7 +45966,9 @@ cp_parser_oacc_compute (cp_parser *parser, cp_token *pragma_tok, } } - tree clauses = cp_parser_oacc_all_clauses (parser, mask, p_name, pragma_tok); + tree clauses = cp_parser_oacc_all_clauses (parser, mask, p_name, pragma_tok, + /*finish_p=*/true, + /*target=*/true); tree block = begin_omp_parallel (); unsigned int save = cp_parser_begin_omp_structured_block (parser); diff --git a/gcc/cp/pt.cc b/gcc/cp/pt.cc index e68c74913f5d..63772bd64e73 100644 --- a/gcc/cp/pt.cc +++ b/gcc/cp/pt.cc @@ -19258,8 +19258,8 @@ tsubst_expr (tree t, tree args, tsubst_flags_t complain, tree in_decl) case OACC_KERNELS: case OACC_PARALLEL: case OACC_SERIAL: - tmp = tsubst_omp_clauses (OMP_CLAUSES (t), C_ORT_ACC, args, complain, - in_decl); + tmp = tsubst_omp_clauses (OMP_CLAUSES (t), C_ORT_ACC_TARGET, args, + complain, in_decl); stmt = begin_omp_parallel (); RECUR (OMP_BODY (t)); finish_omp_construct (TREE_CODE (t), stmt, tmp); diff --git a/gcc/cp/semantics.cc b/gcc/cp/semantics.cc index d6cf1cee6be2..f8f8299bbf64 100644 --- a/gcc/cp/semantics.cc +++ b/gcc/cp/semantics.cc @@ -5167,6 +5167,7 @@ handle_omp_array_sections_1 (tree c, tree t, vec<tree> &types, enum c_omp_region_type ort) { tree ret, low_bound, length, type; + bool openacc = (ort & C_ORT_ACC) != 0; if (TREE_CODE (t) != TREE_LIST) { if (error_operand_p (t)) @@ -5283,7 +5284,7 @@ handle_omp_array_sections_1 (tree c, tree t, vec<tree> &types, { error_at (OMP_CLAUSE_LOCATION (c), "expected single pointer in %qs clause", - user_omp_clause_code_name (c, ort == C_ORT_ACC)); + user_omp_clause_code_name (c, openacc)); return error_mark_node; } } @@ -5772,9 +5773,7 @@ handle_omp_array_sections (tree &c, enum c_omp_region_type ort) cp_omp_address_inspector ai (OMP_CLAUSE_LOCATION (c), t); - tree nc = ai.expand_map_clause (c, first, addr_tokens, - (ort == C_ORT_OMP_TARGET - || ort == C_ORT_ACC)); + tree nc = ai.expand_map_clause (c, first, addr_tokens, ort); if (nc != error_mark_node) { using namespace omp_addr_tokenizer; @@ -6721,6 +6720,7 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort) bitmap_head oacc_reduction_head, is_on_device_head; tree c, t, *pc; tree safelen = NULL_TREE; + bool openacc = (ort & C_ORT_ACC) != 0; bool branch_seen = false; bool copyprivate_seen = false; bool ordered_seen = false; @@ -6753,7 +6753,7 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort) bitmap_initialize (&oacc_reduction_head, &bitmap_default_obstack); bitmap_initialize (&is_on_device_head, &bitmap_default_obstack); - if (ort & C_ORT_ACC) + if (openacc) for (c = clauses; c; c = OMP_CLAUSE_CHAIN (c)) if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_ASYNC) { @@ -7002,7 +7002,7 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort) t = OMP_CLAUSE_DECL (c); check_dup_generic_t: if (t == current_class_ptr - && ((ort != C_ORT_OMP_DECLARE_SIMD && ort != C_ORT_ACC) + && ((ort != C_ORT_OMP_DECLARE_SIMD && !openacc) || (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_LINEAR && OMP_CLAUSE_CODE (c) != OMP_CLAUSE_UNIFORM))) { @@ -7027,7 +7027,7 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort) omp_clause_code_name[OMP_CLAUSE_CODE (c)]); remove = true; } - else if ((ort == C_ORT_ACC + else if ((openacc && OMP_CLAUSE_CODE (c) == OMP_CLAUSE_REDUCTION) || (ort == C_ORT_OMP && (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_USE_DEVICE_PTR @@ -7051,7 +7051,7 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort) if (bitmap_bit_p (&oacc_reduction_head, DECL_UID (t))) { error_at (OMP_CLAUSE_LOCATION (c), - ort == C_ORT_ACC + openacc ? "%qD appears more than once in reduction clauses" : "%qD appears more than once in data clauses", t); @@ -7074,7 +7074,7 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort) || OMP_CLAUSE_CODE (c) == OMP_CLAUSE_IS_DEVICE_PTR) && bitmap_bit_p (&map_head, DECL_UID (t))) { - if (ort == C_ORT_ACC) + if (openacc) error_at (OMP_CLAUSE_LOCATION (c), "%qD appears more than once in data clauses", t); else @@ -7136,7 +7136,7 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort) omp_note_field_privatization (t, OMP_CLAUSE_DECL (c)); else t = OMP_CLAUSE_DECL (c); - if (ort != C_ORT_ACC && t == current_class_ptr) + if (!openacc && t == current_class_ptr) { error_at (OMP_CLAUSE_LOCATION (c), "%<this%> allowed in OpenMP only in %<declare simd%>" @@ -7175,7 +7175,7 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort) else if (bitmap_bit_p (&map_head, DECL_UID (t)) || bitmap_bit_p (&map_field_head, DECL_UID (t))) { - if (ort == C_ORT_ACC) + if (openacc) error_at (OMP_CLAUSE_LOCATION (c), "%qD appears more than once in data clauses", t); else if (OMP_CLAUSE_FIRSTPRIVATE_IMPLICIT (c) @@ -7196,7 +7196,7 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort) omp_note_field_privatization (t, OMP_CLAUSE_DECL (c)); else t = OMP_CLAUSE_DECL (c); - if (ort != C_ORT_ACC && t == current_class_ptr) + if (!openacc && t == current_class_ptr) { error_at (OMP_CLAUSE_LOCATION (c), "%<this%> allowed in OpenMP only in %<declare simd%>" @@ -8076,7 +8076,7 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort) error_at (OMP_CLAUSE_LOCATION (c), "%qD appears more than once in motion" " clauses", rt); - else if (ort == C_ORT_ACC) + else if (openacc) error_at (OMP_CLAUSE_LOCATION (c), "%qD appears more than once in data" " clauses", rt); @@ -8173,8 +8173,7 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort) to be written. */ if (addr_tokens[0]->type == STRUCTURE_BASE && (bitmap_bit_p (&map_field_head, DECL_UID (t)) - || (ort != C_ORT_ACC - && bitmap_bit_p (&map_head, DECL_UID (t))))) + || (!openacc && bitmap_bit_p (&map_head, DECL_UID (t))))) goto skip_decl_checks; if (!processing_template_decl && TREE_CODE (t) == FIELD_DECL) @@ -8270,7 +8269,7 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort) } else if (bitmap_bit_p (&map_head, DECL_UID (t)) && !bitmap_bit_p (&map_field_head, DECL_UID (t)) - && ort == C_ORT_ACC) + && openacc) { error_at (OMP_CLAUSE_LOCATION (c), "%qD appears more than once in data clauses", t); @@ -8289,7 +8288,7 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort) if (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_MAP) error_at (OMP_CLAUSE_LOCATION (c), "%qD appears more than once in motion clauses", t); - else if (ort == C_ORT_ACC) + else if (openacc) error_at (OMP_CLAUSE_LOCATION (c), "%qD appears more than once in data clauses", t); else @@ -8297,8 +8296,7 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort) "%qD appears more than once in map clauses", t); remove = true; } - else if (ort == C_ORT_ACC - && bitmap_bit_p (&generic_head, DECL_UID (t))) + else if (openacc && bitmap_bit_p (&generic_head, DECL_UID (t))) { error_at (OMP_CLAUSE_LOCATION (c), "%qD appears more than once in data clauses", t); @@ -8307,7 +8305,7 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort) else if (bitmap_bit_p (&firstprivate_head, DECL_UID (t)) || bitmap_bit_p (&is_on_device_head, DECL_UID (t))) { - if (ort == C_ORT_ACC) + if (openacc) error_at (OMP_CLAUSE_LOCATION (c), "%qD appears more than once in data clauses", t); else @@ -8331,7 +8329,7 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort) } skip_decl_checks: - /* If we call omp_expand_map_clause in handle_omp_array_sections, + /* If we call ai.expand_map_clause in handle_omp_array_sections, the containing loop (here) iterates through the new nodes created by that expansion. Avoid expanding those again (just by checking the node type). */ @@ -8339,8 +8337,7 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort) && !processing_template_decl && ort != C_ORT_DECLARE_SIMD && (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_MAP - || ((OMP_CLAUSE_MAP_KIND (c) - != GOMP_MAP_FIRSTPRIVATE_POINTER) + || (OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_FIRSTPRIVATE_POINTER && (OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_FIRSTPRIVATE_REFERENCE) && OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_ALWAYS_POINTER @@ -8349,9 +8346,7 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort) grp_start_p = pc; grp_sentinel = OMP_CLAUSE_CHAIN (c); tree nc = ai.expand_map_clause (c, OMP_CLAUSE_DECL (c), - addr_tokens, - (ort == C_ORT_OMP_TARGET - || ort == C_ORT_ACC)); + addr_tokens, ort); if (nc != error_mark_node) c = nc; } -- 2.29.2
