diff --git a/gcc/c-family/c-common.h b/gcc/c-family/c-common.h index bb38e6c76a4..3eb909a2946 100644 --- a/gcc/c-family/c-common.h +++ b/gcc/c-family/c-common.h @@ -1221,6 +1221,7 @@ extern enum omp_clause_defaultmap_kind c_omp_predetermined_mapping (tree); extern tree c_omp_check_context_selector (location_t, tree); extern void c_omp_mark_declare_variant (location_t, tree, tree); extern const char *c_omp_map_clause_name (tree, bool); +extern void c_omp_adjust_map_clauses (tree, bool); /* Return next tree in the chain for chain_next walking of tree nodes. */ static inline tree diff --git a/gcc/c-family/c-omp.c b/gcc/c-family/c-omp.c index d7cff0f4cca..275c6afabe1 100644 --- a/gcc/c-family/c-omp.c +++ b/gcc/c-family/c-omp.c @@ -2579,3 +2579,92 @@ c_omp_map_clause_name (tree clause, bool oacc) } return omp_clause_code_name[OMP_CLAUSE_CODE (clause)]; } + +/* Adjust map clauses after normal clause parsing, mainly to turn specific + base-pointer map cases into attach/detach and mark them addressable. */ +void +c_omp_adjust_map_clauses (tree clauses, bool is_target) +{ + if (!is_target) + { + /* If this is not a target construct, just turn firstprivate pointers + into attach/detach, the runtime will check and do the rest. */ + + for (tree c = clauses; c; c = OMP_CLAUSE_CHAIN (c)) + if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP + && OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_FIRSTPRIVATE_POINTER + && DECL_P (OMP_CLAUSE_DECL (c)) + && POINTER_TYPE_P (TREE_TYPE (OMP_CLAUSE_DECL (c)))) + { + tree ptr = OMP_CLAUSE_DECL (c); + OMP_CLAUSE_SET_MAP_KIND (c, GOMP_MAP_ATTACH_DETACH); + c_common_mark_addressable_vec (ptr); + } + return; + } + + struct map_clause + { + tree clause; + bool firstprivate_ptr_p; + bool decl_mapped; + bool omp_declare_target; + map_clause (void) : clause (NULL_TREE), firstprivate_ptr_p (false), + decl_mapped (false), omp_declare_target (false) { } + }; + + hash_map maps; + + for (tree c = clauses; c; c = OMP_CLAUSE_CHAIN (c)) + if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP + && DECL_P (OMP_CLAUSE_DECL (c))) + { + /* If this is for a target construct, the firstprivate pointer + is changed to attach/detach if either is true: + (1) the base-pointer is mapped in this same construct, or + (2) the base-pointer is a variable place on the device by + "declare target" directives. + + Here we iterate through all map clauses collecting these cases, + and merge them with a hash_map to process below. */ + + if (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_FIRSTPRIVATE_POINTER + && POINTER_TYPE_P (TREE_TYPE (OMP_CLAUSE_DECL (c)))) + { + tree ptr = OMP_CLAUSE_DECL (c); + map_clause &mc = maps.get_or_insert (ptr); + if (mc.clause == NULL_TREE) + mc.clause = c; + mc.firstprivate_ptr_p = true; + + if (is_global_var (ptr) + && lookup_attribute ("omp declare target", + DECL_ATTRIBUTES (ptr))) + mc.omp_declare_target = true; + } + else if (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ALLOC + || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_TO + || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_FROM + || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_TOFROM + || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ALWAYS_TO + || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ALWAYS_FROM + || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ALWAYS_TOFROM) + { + map_clause &mc = maps.get_or_insert (OMP_CLAUSE_DECL (c)); + mc.decl_mapped = true; + } + } + + for (hash_map::iterator i = maps.begin (); + i != maps.end (); ++i) + { + map_clause &mc = (*i).second; + + if (mc.firstprivate_ptr_p + && (mc.decl_mapped || mc.omp_declare_target)) + { + OMP_CLAUSE_SET_MAP_KIND (mc.clause, GOMP_MAP_ATTACH_DETACH); + c_common_mark_addressable_vec (OMP_CLAUSE_DECL (mc.clause)); + } + } +} diff --git a/gcc/c/c-parser.c b/gcc/c/c-parser.c index b6a7ef4c92b..ab7b0bbc29f 100644 --- a/gcc/c/c-parser.c +++ b/gcc/c/c-parser.c @@ -19470,6 +19470,7 @@ c_parser_omp_target_data (location_t loc, c_parser *parser, bool *if_p) tree clauses = c_parser_omp_all_clauses (parser, OMP_TARGET_DATA_CLAUSE_MASK, "#pragma omp target data"); + c_omp_adjust_map_clauses (clauses, false); int map_seen = 0; for (tree *pc = &clauses; *pc;) { @@ -19487,6 +19488,7 @@ c_parser_omp_target_data (location_t loc, c_parser *parser, bool *if_p) break; case GOMP_MAP_FIRSTPRIVATE_POINTER: case GOMP_MAP_ALWAYS_POINTER: + case GOMP_MAP_ATTACH_DETACH: break; default: map_seen |= 1; @@ -19610,6 +19612,7 @@ c_parser_omp_target_enter_data (location_t loc, c_parser *parser, tree clauses = c_parser_omp_all_clauses (parser, OMP_TARGET_ENTER_DATA_CLAUSE_MASK, "#pragma omp target enter data"); + c_omp_adjust_map_clauses (clauses, false); int map_seen = 0; for (tree *pc = &clauses; *pc;) { @@ -19623,6 +19626,7 @@ c_parser_omp_target_enter_data (location_t loc, c_parser *parser, break; case GOMP_MAP_FIRSTPRIVATE_POINTER: case GOMP_MAP_ALWAYS_POINTER: + case GOMP_MAP_ATTACH_DETACH: break; default: map_seen |= 1; @@ -19694,7 +19698,7 @@ c_parser_omp_target_exit_data (location_t loc, c_parser *parser, tree clauses = c_parser_omp_all_clauses (parser, OMP_TARGET_EXIT_DATA_CLAUSE_MASK, "#pragma omp target exit data"); - + c_omp_adjust_map_clauses (clauses, false); int map_seen = 0; for (tree *pc = &clauses; *pc;) { @@ -19709,6 +19713,7 @@ c_parser_omp_target_exit_data (location_t loc, c_parser *parser, break; case GOMP_MAP_FIRSTPRIVATE_POINTER: case GOMP_MAP_ALWAYS_POINTER: + case GOMP_MAP_ATTACH_DETACH: break; default: map_seen |= 1; @@ -19918,6 +19923,8 @@ c_parser_omp_target (c_parser *parser, enum pragma_context context, bool *if_p) OMP_TARGET_CLAUSES (stmt) = c_parser_omp_all_clauses (parser, OMP_TARGET_CLAUSE_MASK, "#pragma omp target"); + c_omp_adjust_map_clauses (OMP_TARGET_CLAUSES (stmt), true); + pc = &OMP_TARGET_CLAUSES (stmt); keep_next_level (); block = c_begin_compound_stmt (true); @@ -19942,6 +19949,7 @@ check_clauses: case GOMP_MAP_ALLOC: case GOMP_MAP_FIRSTPRIVATE_POINTER: case GOMP_MAP_ALWAYS_POINTER: + case GOMP_MAP_ATTACH_DETACH: break; default: error_at (OMP_CLAUSE_LOCATION (*pc), diff --git a/gcc/c/c-typeck.c b/gcc/c/c-typeck.c index 459090e227d..ada6662fca7 100644 --- a/gcc/c/c-typeck.c +++ b/gcc/c/c-typeck.c @@ -13501,11 +13501,7 @@ handle_omp_array_sections (tree c, enum c_omp_region_type ort) if (ort != C_ORT_OMP && ort != C_ORT_ACC) OMP_CLAUSE_SET_MAP_KIND (c2, GOMP_MAP_POINTER); else if (TREE_CODE (t) == COMPONENT_REF) - { - gomp_map_kind k = (ort == C_ORT_ACC) ? GOMP_MAP_ATTACH_DETACH - : GOMP_MAP_ALWAYS_POINTER; - OMP_CLAUSE_SET_MAP_KIND (c2, k); - } + OMP_CLAUSE_SET_MAP_KIND (c2, GOMP_MAP_ATTACH_DETACH); else OMP_CLAUSE_SET_MAP_KIND (c2, GOMP_MAP_FIRSTPRIVATE_POINTER); if (OMP_CLAUSE_MAP_KIND (c2) != GOMP_MAP_FIRSTPRIVATE_POINTER @@ -14604,7 +14600,9 @@ c_finish_omp_clauses (tree clauses, enum c_omp_region_type ort) break; if (VAR_P (t) || TREE_CODE (t) == PARM_DECL) { - if (bitmap_bit_p (&map_field_head, DECL_UID (t))) + if (bitmap_bit_p (&map_field_head, DECL_UID (t)) + || (ort == C_ORT_OMP + && bitmap_bit_p (&map_head, DECL_UID (t)))) break; } } @@ -14673,7 +14671,9 @@ c_finish_omp_clauses (tree clauses, enum c_omp_region_type ort) else bitmap_set_bit (&generic_head, DECL_UID (t)); } - else if (bitmap_bit_p (&map_head, DECL_UID (t))) + else if (bitmap_bit_p (&map_head, DECL_UID (t)) + && (ort != C_ORT_OMP + || !bitmap_bit_p (&map_field_head, DECL_UID (t)))) { if (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_MAP) error_at (OMP_CLAUSE_LOCATION (c), @@ -14687,7 +14687,13 @@ c_finish_omp_clauses (tree clauses, enum c_omp_region_type ort) remove = true; } else if (bitmap_bit_p (&generic_head, DECL_UID (t)) - || bitmap_bit_p (&firstprivate_head, DECL_UID (t))) + && ort == C_ORT_ACC) + { + error_at (OMP_CLAUSE_LOCATION (c), + "%qD appears more than once in data clauses", t); + remove = true; + } + else if (bitmap_bit_p (&firstprivate_head, DECL_UID (t))) { if (ort == C_ORT_ACC) error_at (OMP_CLAUSE_LOCATION (c), diff --git a/gcc/cp/parser.c b/gcc/cp/parser.c index 7ec7d42773c..8527a7d0478 100644 --- a/gcc/cp/parser.c +++ b/gcc/cp/parser.c @@ -40510,6 +40510,7 @@ cp_parser_omp_target_data (cp_parser *parser, cp_token *pragma_tok, bool *if_p) tree clauses = cp_parser_omp_all_clauses (parser, OMP_TARGET_DATA_CLAUSE_MASK, "#pragma omp target data", pragma_tok); + c_omp_adjust_map_clauses (clauses, false); int map_seen = 0; for (tree *pc = &clauses; *pc;) { @@ -40528,6 +40529,7 @@ cp_parser_omp_target_data (cp_parser *parser, cp_token *pragma_tok, bool *if_p) case GOMP_MAP_FIRSTPRIVATE_POINTER: case GOMP_MAP_FIRSTPRIVATE_REFERENCE: case GOMP_MAP_ALWAYS_POINTER: + case GOMP_MAP_ATTACH_DETACH: break; default: map_seen |= 1; @@ -40611,6 +40613,7 @@ cp_parser_omp_target_enter_data (cp_parser *parser, cp_token *pragma_tok, tree clauses = cp_parser_omp_all_clauses (parser, OMP_TARGET_ENTER_DATA_CLAUSE_MASK, "#pragma omp target enter data", pragma_tok); + c_omp_adjust_map_clauses (clauses, false); int map_seen = 0; for (tree *pc = &clauses; *pc;) { @@ -40625,6 +40628,7 @@ cp_parser_omp_target_enter_data (cp_parser *parser, cp_token *pragma_tok, case GOMP_MAP_FIRSTPRIVATE_POINTER: case GOMP_MAP_FIRSTPRIVATE_REFERENCE: case GOMP_MAP_ALWAYS_POINTER: + case GOMP_MAP_ATTACH_DETACH: break; default: map_seen |= 1; @@ -40699,6 +40703,7 @@ cp_parser_omp_target_exit_data (cp_parser *parser, cp_token *pragma_tok, tree clauses = cp_parser_omp_all_clauses (parser, OMP_TARGET_EXIT_DATA_CLAUSE_MASK, "#pragma omp target exit data", pragma_tok); + c_omp_adjust_map_clauses (clauses, false); int map_seen = 0; for (tree *pc = &clauses; *pc;) { @@ -40714,6 +40719,7 @@ cp_parser_omp_target_exit_data (cp_parser *parser, cp_token *pragma_tok, case GOMP_MAP_FIRSTPRIVATE_POINTER: case GOMP_MAP_FIRSTPRIVATE_REFERENCE: case GOMP_MAP_ALWAYS_POINTER: + case GOMP_MAP_ATTACH_DETACH: break; default: map_seen |= 1; @@ -40962,6 +40968,8 @@ cp_parser_omp_target (cp_parser *parser, cp_token *pragma_tok, OMP_TARGET_CLAUSES (stmt) = cp_parser_omp_all_clauses (parser, OMP_TARGET_CLAUSE_MASK, "#pragma omp target", pragma_tok); + c_omp_adjust_map_clauses (OMP_TARGET_CLAUSES (stmt), true); + pc = &OMP_TARGET_CLAUSES (stmt); keep_next_level (true); OMP_TARGET_BODY (stmt) = cp_parser_omp_structured_block (parser, if_p); @@ -40985,6 +40993,7 @@ check_clauses: case GOMP_MAP_FIRSTPRIVATE_POINTER: case GOMP_MAP_FIRSTPRIVATE_REFERENCE: case GOMP_MAP_ALWAYS_POINTER: + case GOMP_MAP_ATTACH_DETACH: break; default: error_at (OMP_CLAUSE_LOCATION (*pc), diff --git a/gcc/cp/semantics.c b/gcc/cp/semantics.c index 1e42cd799c2..a8dc769db34 100644 --- a/gcc/cp/semantics.c +++ b/gcc/cp/semantics.c @@ -5382,11 +5382,7 @@ handle_omp_array_sections (tree c, enum c_omp_region_type ort) if ((ort & C_ORT_OMP_DECLARE_SIMD) != C_ORT_OMP && ort != C_ORT_ACC) OMP_CLAUSE_SET_MAP_KIND (c2, GOMP_MAP_POINTER); else if (TREE_CODE (t) == COMPONENT_REF) - { - gomp_map_kind k = (ort == C_ORT_ACC) ? GOMP_MAP_ATTACH_DETACH - : GOMP_MAP_ALWAYS_POINTER; - OMP_CLAUSE_SET_MAP_KIND (c2, k); - } + OMP_CLAUSE_SET_MAP_KIND (c2, GOMP_MAP_ATTACH_DETACH); else if (REFERENCE_REF_P (t) && TREE_CODE (TREE_OPERAND (t, 0)) == COMPONENT_REF) { @@ -5424,8 +5420,12 @@ handle_omp_array_sections (tree c, enum c_omp_region_type ort) OMP_CLAUSE_MAP); OMP_CLAUSE_SET_MAP_KIND (c3, OMP_CLAUSE_MAP_KIND (c2)); OMP_CLAUSE_DECL (c3) = ptr; - if (OMP_CLAUSE_MAP_KIND (c2) == GOMP_MAP_ALWAYS_POINTER) - OMP_CLAUSE_DECL (c2) = build_simple_mem_ref (ptr); + if (OMP_CLAUSE_MAP_KIND (c2) == GOMP_MAP_ALWAYS_POINTER + || OMP_CLAUSE_MAP_KIND (c2) == GOMP_MAP_ATTACH_DETACH) + { + OMP_CLAUSE_DECL (c2) = build_simple_mem_ref (ptr); + OMP_CLAUSE_SET_MAP_KIND (c2, GOMP_MAP_ALWAYS_POINTER); + } else OMP_CLAUSE_DECL (c2) = convert_from_reference (ptr); OMP_CLAUSE_SIZE (c3) = size_zero_node; @@ -7411,7 +7411,7 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort) t = TREE_OPERAND (t, 0); OMP_CLAUSE_DECL (c) = t; } - if (ort == C_ORT_ACC + if ((ort == C_ORT_ACC || ort == C_ORT_OMP) && TREE_CODE (t) == COMPONENT_REF && TREE_CODE (TREE_OPERAND (t, 0)) == INDIRECT_REF) t = TREE_OPERAND (TREE_OPERAND (t, 0), 0); @@ -7457,7 +7457,9 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort) t = TREE_OPERAND (t, 0); if (VAR_P (t) || TREE_CODE (t) == PARM_DECL) { - if (bitmap_bit_p (&map_field_head, DECL_UID (t))) + if (bitmap_bit_p (&map_field_head, DECL_UID (t)) + || (ort == C_ORT_OMP + && bitmap_bit_p (&map_head, DECL_UID (t)))) goto handle_map_references; } } @@ -7551,13 +7553,12 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort) bitmap_set_bit (&generic_head, DECL_UID (t)); } else if (bitmap_bit_p (&map_head, DECL_UID (t)) - && (ort != C_ORT_ACC - || !bitmap_bit_p (&map_field_head, DECL_UID (t)))) + && !bitmap_bit_p (&map_field_head, DECL_UID (t))) { if (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_MAP) error_at (OMP_CLAUSE_LOCATION (c), "%qD appears more than once in motion clauses", t); - if (ort == C_ORT_ACC) + else if (ort == C_ORT_ACC) error_at (OMP_CLAUSE_LOCATION (c), "%qD appears more than once in data clauses", t); else @@ -7566,7 +7567,13 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort) remove = true; } else if (bitmap_bit_p (&generic_head, DECL_UID (t)) - || bitmap_bit_p (&firstprivate_head, DECL_UID (t))) + && ort == C_ORT_ACC) + { + error_at (OMP_CLAUSE_LOCATION (c), + "%qD appears more than once in data clauses", t); + remove = true; + } + else if (bitmap_bit_p (&firstprivate_head, DECL_UID (t))) { if (ort == C_ORT_ACC) error_at (OMP_CLAUSE_LOCATION (c), @@ -7602,17 +7609,14 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort) && (OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_FIRSTPRIVATE_REFERENCE) && (OMP_CLAUSE_MAP_KIND (c) - != GOMP_MAP_ALWAYS_POINTER)) + != GOMP_MAP_ALWAYS_POINTER) + && (OMP_CLAUSE_MAP_KIND (c) + != GOMP_MAP_ATTACH_DETACH)) { tree c2 = build_omp_clause (OMP_CLAUSE_LOCATION (c), OMP_CLAUSE_MAP); if (TREE_CODE (t) == COMPONENT_REF) - { - gomp_map_kind k - = (ort == C_ORT_ACC) ? GOMP_MAP_ATTACH_DETACH - : GOMP_MAP_ALWAYS_POINTER; - OMP_CLAUSE_SET_MAP_KIND (c2, k); - } + OMP_CLAUSE_SET_MAP_KIND (c2, GOMP_MAP_ATTACH_DETACH); else OMP_CLAUSE_SET_MAP_KIND (c2, GOMP_MAP_FIRSTPRIVATE_REFERENCE);