From 9e6280242225587be256fdb80c41327736238e77 Mon Sep 17 00:00:00 2001 From: Chung-Lin Tang Date: Tue, 10 Nov 2020 03:36:58 -0800 Subject: [PATCH] openmp: Implement OpenMP 5.0 base-pointer attachement and clause ordering This patch implements some parts of the target variable mapping changes specified in OpenMP 5.0, including base-pointer attachment/detachment behavior for array section list-items in map clauses, and ordering of map clauses according to map kind. 2020-11-10 Chung-Lin Tang gcc/c-family/ChangeLog: * c-common.h (c_omp_adjust_map_clauses): New declaration. * c-omp.c (struct map_clause): Helper type for c_omp_adjust_map_clauses. (c_omp_adjust_map_clauses): New function. gcc/c/ChangeLog: * c-parser.c (c_parser_omp_target_data): Add use of new c_omp_adjust_map_clauses function. Add GOMP_MAP_ATTACH_DETACH as handled map clause kind. (c_parser_omp_target_enter_data): Likewise. (c_parser_omp_target_exit_data): Likewise. (c_parser_omp_target): Likewise. * c-typeck.c (handle_omp_array_sections): Adjust COMPONENT_REF case to use GOMP_MAP_ATTACH_DETACH map kind for C_ORT_OMP region type. (c_finish_omp_clauses): Adjust bitmap checks to allow struct decl and same struct field access to co-exist on OpenMP construct. gcc/cp/ChangeLog: * parser.c (cp_parser_omp_target_data): Add use of new c_omp_adjust_map_clauses function. Add GOMP_MAP_ATTACH_DETACH as handled map clause kind. (cp_parser_omp_target_enter_data): Likewise. (cp_parser_omp_target_exit_data): Likewise. (cp_parser_omp_target): Likewise. * semantics.c (handle_omp_array_sections): Adjust COMPONENT_REF case to use GOMP_MAP_ATTACH_DETACH map kind for C_ORT_OMP region type. Fix interaction between reference case and attach/detach. (finish_omp_clauses): Adjust bitmap checks to allow struct decl and same struct field access to co-exist on OpenMP construct. gcc/ChangeLog: * gimplify.c (is_or_contains_p): New static helper function. (omp_target_reorder_clauses): New function. (gimplify_scan_omp_clauses): Add use of omp_target_reorder_clauses to reorder clause list according to OpenMP 5.0 rules. Add handling of GOMP_MAP_ATTACH_DETACH for OpenMP cases. * omp-low.c (is_omp_target): New static helper function. (scan_sharing_clauses): Add scan phase handling of GOMP_MAP_ATTACH/DETACH for OpenMP cases. (lower_omp_target): Add lowering handling of GOMP_MAP_ATTACH/DETACH for OpenMP cases. gcc/testsuite/ChangeLog: * c-c++-common/gomp/clauses-2.c: Remove dg-error cases now valid. * gfortran.dg/gomp/map-2.f90: Likewise. * c-c++-common/gomp/map-5.c: New testcase. libgomp/ChangeLog: * libgomp.h (enum gomp_map_vars_kind): Adjust enum values to be bit-flag usable. * oacc-mem.c (acc_map_data): Adjust gomp_map_vars argument flags to 'GOMP_MAP_VARS_OPENACC | GOMP_MAP_VARS_ENTER_DATA'. (goacc_enter_datum): Likewise for call to gomp_map_vars_async. (goacc_enter_data_internal): Likewise. * target.c (gomp_map_vars_internal): Change checks of GOMP_MAP_VARS_ENTER_DATA to use bit-and (&). Adjust use of gomp_attach_pointer for OpenMP cases. (gomp_exit_data): Add handling of GOMP_MAP_DETACH. (GOMP_target_enter_exit_data): Add handling of GOMP_MAP_ATTACH. * testsuite/libgomp.c-c++-common/ptr-attach-1.c: New testcase. --- gcc/c-family/c-common.h | 1 + gcc/c-family/c-omp.c | 90 ++++++ gcc/c/c-parser.c | 10 +- gcc/c/c-typeck.c | 22 +- gcc/cp/parser.c | 9 + gcc/cp/semantics.c | 44 +-- gcc/gimplify.c | 272 +++++++++++++++--- gcc/omp-low.c | 90 +++++- gcc/testsuite/c-c++-common/gomp/clauses-2.c | 20 +- gcc/testsuite/c-c++-common/gomp/map-5.c | 24 ++ gcc/testsuite/gfortran.dg/gomp/map-2.f90 | 2 +- libgomp/libgomp.h | 8 +- libgomp/oacc-mem.c | 9 +- libgomp/target.c | 36 ++- .../libgomp.c-c++-common/ptr-attach-1.c | 82 ++++++ 15 files changed, 616 insertions(+), 103 deletions(-) create mode 100644 gcc/testsuite/c-c++-common/gomp/map-5.c create mode 100644 libgomp/testsuite/libgomp.c-c++-common/ptr-attach-1.c diff --git a/gcc/c-family/c-common.h b/gcc/c-family/c-common.h index bfcc279609c..b80db230a6d 100644 --- a/gcc/c-family/c-common.h +++ b/gcc/c-family/c-common.h @@ -1224,6 +1224,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 da5564be5ba..84572116c76 100644 --- a/gcc/c-family/c-omp.c +++ b/gcc/c-family/c-omp.c @@ -2771,3 +2771,93 @@ c_omp_map_clause_name (tree clause, bool oacc) } return omp_clause_code_name[OMP_CLAUSE_CODE (clause)]; } + +/* Used to merge map clause information in c_omp_adjust_map_clauses. */ +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) { } +}; + +/* 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; + } + + 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 ecc3d2119fa..377914cad16 100644 --- a/gcc/c/c-parser.c +++ b/gcc/c/c-parser.c @@ -19511,6 +19511,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;) { @@ -19528,6 +19529,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; @@ -19651,6 +19653,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;) { @@ -19664,6 +19667,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; @@ -19735,7 +19739,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;) { @@ -19750,6 +19754,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; @@ -19960,6 +19965,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); @@ -19984,6 +19991,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 96840377d90..df1dad468df 100644 --- a/gcc/c/c-typeck.c +++ b/gcc/c/c-typeck.c @@ -13584,11 +13584,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 @@ -14711,7 +14707,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; } } @@ -14780,7 +14778,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), @@ -14794,7 +14794,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 4c819ea1c5d..9c08c0e46a2 100644 --- a/gcc/cp/parser.c +++ b/gcc/cp/parser.c @@ -40785,6 +40785,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;) { @@ -40803,6 +40804,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; @@ -40886,6 +40888,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;) { @@ -40900,6 +40903,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; @@ -40974,6 +40978,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;) { @@ -40989,6 +40994,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; @@ -41238,6 +41244,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); @@ -41261,6 +41269,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 a550db69488..33d715edaec 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; @@ -7486,7 +7486,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); @@ -7532,7 +7532,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; } } @@ -7626,13 +7628,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 @@ -7641,7 +7642,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), @@ -7677,17 +7684,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); diff --git a/gcc/gimplify.c b/gcc/gimplify.c index aa3b914f6e5..b2c623be456 100644 --- a/gcc/gimplify.c +++ b/gcc/gimplify.c @@ -8364,6 +8364,113 @@ extract_base_bit_offset (tree base, tree *base_ref, poly_int64 *bitposp, return base; } +/* Returns true if EXPR is or contains (as a sub-component) BASE_PTR. */ + +static bool +is_or_contains_p (tree expr, tree base_ptr) +{ + while (expr != base_ptr) + if (TREE_CODE (base_ptr) == COMPONENT_REF) + base_ptr = TREE_OPERAND (base_ptr, 0); + else + break; + return expr == base_ptr; +} + +/* Implement OpenMP 5.x map ordering rules for target directives. There are + several rules, and with some level of ambiguity, hopefully we can at least + collect the complexity here in one place. */ + +static void +omp_target_reorder_clauses (tree *list_p) +{ + /* Collect refs to alloc/release/delete maps. */ + auto_vec ard; + tree *cp = list_p; + while (*cp != NULL_TREE) + if (OMP_CLAUSE_CODE (*cp) == OMP_CLAUSE_MAP + && (OMP_CLAUSE_MAP_KIND (*cp) == GOMP_MAP_ALLOC + || OMP_CLAUSE_MAP_KIND (*cp) == GOMP_MAP_RELEASE + || OMP_CLAUSE_MAP_KIND (*cp) == GOMP_MAP_DELETE)) + { + /* Unlink cp and push to ard. */ + tree c = *cp; + tree nc = OMP_CLAUSE_CHAIN (c); + *cp = nc; + ard.safe_push (c); + + /* Any associated pointer type maps should also move along. */ + while (*cp != NULL_TREE + && OMP_CLAUSE_CODE (*cp) == OMP_CLAUSE_MAP + && (OMP_CLAUSE_MAP_KIND (*cp) == GOMP_MAP_FIRSTPRIVATE_REFERENCE + || OMP_CLAUSE_MAP_KIND (*cp) == GOMP_MAP_FIRSTPRIVATE_POINTER + || OMP_CLAUSE_MAP_KIND (*cp) == GOMP_MAP_ATTACH_DETACH + || OMP_CLAUSE_MAP_KIND (*cp) == GOMP_MAP_POINTER + || OMP_CLAUSE_MAP_KIND (*cp) == GOMP_MAP_ALWAYS_POINTER + || OMP_CLAUSE_MAP_KIND (*cp) == GOMP_MAP_TO_PSET)) + { + c = *cp; + nc = OMP_CLAUSE_CHAIN (c); + *cp = nc; + ard.safe_push (c); + } + } + else + cp = &OMP_CLAUSE_CHAIN (*cp); + + /* Link alloc/release/delete maps to the end of list. */ + for (unsigned int i = 0; i < ard.length (); i++) + { + *cp = ard[i]; + cp = &OMP_CLAUSE_CHAIN (ard[i]); + } + *cp = NULL_TREE; + + /* OpenMP 5.0 requires that pointer variables are mapped before + its use as a base-pointer. */ + auto_vec atf; + for (tree *cp = list_p; *cp; cp = &OMP_CLAUSE_CHAIN (*cp)) + if (OMP_CLAUSE_CODE (*cp) == OMP_CLAUSE_MAP) + { + /* Collect alloc, to, from, to/from clause tree pointers. */ + gomp_map_kind k = OMP_CLAUSE_MAP_KIND (*cp); + if (k == GOMP_MAP_ALLOC + || k == GOMP_MAP_TO + || k == GOMP_MAP_FROM + || k == GOMP_MAP_TOFROM + || k == GOMP_MAP_ALWAYS_TO + || k == GOMP_MAP_ALWAYS_FROM + || k == GOMP_MAP_ALWAYS_TOFROM) + atf.safe_push (cp); + } + + for (unsigned int i = 0; i < atf.length (); i++) + if (atf[i]) + { + tree *cp = atf[i]; + tree decl = OMP_CLAUSE_DECL (*cp); + if (TREE_CODE (decl) == INDIRECT_REF || TREE_CODE (decl) == MEM_REF) + { + tree base_ptr = TREE_OPERAND (decl, 0); + STRIP_TYPE_NOPS (base_ptr); + for (unsigned int j = i + 1; j < atf.length (); j++) + { + tree *cp2 = atf[j]; + tree decl2 = OMP_CLAUSE_DECL (*cp2); + if (is_or_contains_p (decl2, base_ptr)) + { + /* Move *cp2 to before *cp. */ + tree c = *cp2; + *cp2 = OMP_CLAUSE_CHAIN (c); + OMP_CLAUSE_CHAIN (c) = *cp; + *cp = c; + atf[j] = NULL; + } + } + } + } +} + /* Scan the OMP clauses in *LIST_P, installing mappings into a new and previous omp contexts. */ @@ -8405,6 +8512,12 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p, break; } + if (code == OMP_TARGET + || code == OMP_TARGET_DATA + || code == OMP_TARGET_ENTER_DATA + || code == OMP_TARGET_EXIT_DATA) + omp_target_reorder_clauses (list_p); + while ((c = *list_p) != NULL) { bool remove = false; @@ -8845,15 +8958,18 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p, } else if ((OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_FIRSTPRIVATE_POINTER || (OMP_CLAUSE_MAP_KIND (c) - == GOMP_MAP_FIRSTPRIVATE_REFERENCE)) + == GOMP_MAP_FIRSTPRIVATE_REFERENCE) + || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ATTACH_DETACH) && TREE_CODE (OMP_CLAUSE_SIZE (c)) != INTEGER_CST) { OMP_CLAUSE_SIZE (c) = get_initialized_tmp_var (OMP_CLAUSE_SIZE (c), pre_p, NULL, false); - omp_add_variable (ctx, OMP_CLAUSE_SIZE (c), - GOVD_FIRSTPRIVATE | GOVD_SEEN); + if ((region_type & ORT_TARGET) != 0) + omp_add_variable (ctx, OMP_CLAUSE_SIZE (c), + GOVD_FIRSTPRIVATE | GOVD_SEEN); } + if (!DECL_P (decl)) { tree d = decl, *pd; @@ -8878,7 +8994,7 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p, bool indir_p = false; tree orig_decl = decl; tree decl_ref = NULL_TREE; - if ((region_type & ORT_ACC) != 0 + if ((region_type & (ORT_ACC | ORT_TARGET | ORT_TARGET_DATA)) != 0 && TREE_CODE (*pd) == COMPONENT_REF && OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ATTACH_DETACH && code != OACC_UPDATE) @@ -8886,9 +9002,11 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p, while (TREE_CODE (decl) == COMPONENT_REF) { decl = TREE_OPERAND (decl, 0); - if ((TREE_CODE (decl) == MEM_REF - && integer_zerop (TREE_OPERAND (decl, 1))) - || INDIRECT_REF_P (decl)) + if (((TREE_CODE (decl) == MEM_REF + && integer_zerop (TREE_OPERAND (decl, 1))) + || INDIRECT_REF_P (decl)) + && (TREE_CODE (TREE_TYPE (TREE_OPERAND (decl, 0))) + == POINTER_TYPE)) { indir_p = true; decl = TREE_OPERAND (decl, 0); @@ -8915,8 +9033,9 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p, } if (decl != orig_decl && DECL_P (decl) && indir_p) { - gomp_map_kind k = (code == OACC_EXIT_DATA) ? GOMP_MAP_DETACH - : GOMP_MAP_ATTACH; + gomp_map_kind k + = ((code == OACC_EXIT_DATA || code == OMP_TARGET_EXIT_DATA) + ? GOMP_MAP_DETACH : GOMP_MAP_ATTACH); /* We have a dereference of a struct member. Make this an attach/detach operation, and ensure the base pointer is mapped as a FIRSTPRIVATE_POINTER. */ @@ -8925,6 +9044,7 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p, tree next_clause = OMP_CLAUSE_CHAIN (c); if (k == GOMP_MAP_ATTACH && code != OACC_ENTER_DATA + && code != OMP_TARGET_ENTER_DATA && (!next_clause || (OMP_CLAUSE_CODE (next_clause) != OMP_CLAUSE_MAP) || (OMP_CLAUSE_MAP_KIND (next_clause) @@ -8972,17 +9092,12 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p, if (code == OACC_UPDATE && OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ATTACH_DETACH) OMP_CLAUSE_SET_MAP_KIND (c, GOMP_MAP_ALWAYS_POINTER); - if (gimplify_expr (pd, pre_p, NULL, is_gimple_lvalue, fb_lvalue) - == GS_ERROR) - { - remove = true; - break; - } if (DECL_P (decl) && OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_TO_PSET && OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_ATTACH && OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_DETACH - && code != OACC_UPDATE) + && code != OACC_UPDATE + && code != OMP_TARGET_UPDATE) { if (error_operand_p (decl)) { @@ -9044,15 +9159,19 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p, bool has_attachments = false; /* For OpenACC, pointers in structs should trigger an attach action. */ - if (attach_detach && (region_type & ORT_ACC) != 0) + if (attach_detach + && ((region_type & (ORT_ACC | ORT_TARGET | ORT_TARGET_DATA)) + || code == OMP_TARGET_ENTER_DATA + || code == OMP_TARGET_EXIT_DATA)) + { /* Turn a GOMP_MAP_ATTACH_DETACH clause into a GOMP_MAP_ATTACH or GOMP_MAP_DETACH clause after we have detected a case that needs a GOMP_MAP_STRUCT mapping added. */ gomp_map_kind k - = (code == OACC_EXIT_DATA) ? GOMP_MAP_DETACH - : GOMP_MAP_ATTACH; + = ((code == OACC_EXIT_DATA || code == OMP_TARGET_EXIT_DATA) + ? GOMP_MAP_DETACH : GOMP_MAP_ATTACH); OMP_CLAUSE_SET_MAP_KIND (c, k); has_attachments = true; } @@ -9148,33 +9267,38 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p, break; if (scp) continue; - tree d1 = OMP_CLAUSE_DECL (*sc); - tree d2 = OMP_CLAUSE_DECL (c); - while (TREE_CODE (d1) == ARRAY_REF) - d1 = TREE_OPERAND (d1, 0); - while (TREE_CODE (d2) == ARRAY_REF) - d2 = TREE_OPERAND (d2, 0); - if (TREE_CODE (d1) == INDIRECT_REF) - d1 = TREE_OPERAND (d1, 0); - if (TREE_CODE (d2) == INDIRECT_REF) - d2 = TREE_OPERAND (d2, 0); - while (TREE_CODE (d1) == COMPONENT_REF) - if (TREE_CODE (d2) == COMPONENT_REF - && TREE_OPERAND (d1, 1) - == TREE_OPERAND (d2, 1)) - { - d1 = TREE_OPERAND (d1, 0); - d2 = TREE_OPERAND (d2, 0); - } - else - break; - if (d1 == d2) + if ((region_type & ORT_ACC) != 0) { - error_at (OMP_CLAUSE_LOCATION (c), - "%qE appears more than once in map " - "clauses", OMP_CLAUSE_DECL (c)); - remove = true; - break; + /* This duplicate checking code is currently only + enabled for OpenACC. */ + tree d1 = OMP_CLAUSE_DECL (*sc); + tree d2 = OMP_CLAUSE_DECL (c); + while (TREE_CODE (d1) == ARRAY_REF) + d1 = TREE_OPERAND (d1, 0); + while (TREE_CODE (d2) == ARRAY_REF) + d2 = TREE_OPERAND (d2, 0); + if (TREE_CODE (d1) == INDIRECT_REF) + d1 = TREE_OPERAND (d1, 0); + if (TREE_CODE (d2) == INDIRECT_REF) + d2 = TREE_OPERAND (d2, 0); + while (TREE_CODE (d1) == COMPONENT_REF) + if (TREE_CODE (d2) == COMPONENT_REF + && TREE_OPERAND (d1, 1) + == TREE_OPERAND (d2, 1)) + { + d1 = TREE_OPERAND (d1, 0); + d2 = TREE_OPERAND (d2, 0); + } + else + break; + if (d1 == d2) + { + error_at (OMP_CLAUSE_LOCATION (c), + "%qE appears more than once in map " + "clauses", OMP_CLAUSE_DECL (c)); + remove = true; + break; + } } if (maybe_lt (offset1, offsetn) || (known_eq (offset1, offsetn) @@ -9220,6 +9344,14 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p, } } } + + if (gimplify_expr (pd, pre_p, NULL, is_gimple_lvalue, fb_lvalue) + == GS_ERROR) + { + remove = true; + break; + } + if (!remove && OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_ALWAYS_POINTER && OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_ATTACH_DETACH @@ -9236,10 +9368,60 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p, break; } + else + { + /* DECL_P (decl) == true */ + tree *sc; + if (struct_map_to_clause + && (sc = struct_map_to_clause->get (decl)) != NULL + && OMP_CLAUSE_MAP_KIND (*sc) == GOMP_MAP_STRUCT + && decl == OMP_CLAUSE_DECL (*sc)) + { + /* We have found a map of the whole structure after a + leading GOMP_MAP_STRUCT has been created, so refill the + leading clause into a map of the whole structure + variable, and remove the current one. + TODO: we should be able to remove some maps of the + following structure element maps if they are of + compatible TO/FROM/ALLOC type. */ + OMP_CLAUSE_SET_MAP_KIND (*sc, OMP_CLAUSE_MAP_KIND (c)); + OMP_CLAUSE_SIZE (*sc) = unshare_expr (OMP_CLAUSE_SIZE (c)); + remove = true; + break; + } + } flags = GOVD_MAP | GOVD_EXPLICIT; if (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ALWAYS_TO || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ALWAYS_TOFROM) flags |= GOVD_MAP_ALWAYS_TO; + + if ((code == OMP_TARGET + || code == OMP_TARGET_DATA + || code == OMP_TARGET_ENTER_DATA + || code == OMP_TARGET_EXIT_DATA) + && OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ATTACH_DETACH) + { + for (struct gimplify_omp_ctx *octx = outer_ctx; octx; + octx = octx->outer_context) + { + splay_tree_node n + = splay_tree_lookup (octx->variables, + (splay_tree_key) OMP_CLAUSE_DECL (c)); + /* If this is contained in an outer OpenMP region as a + firstprivate value, remove the attach/detach. */ + if (n && (n->value & GOVD_FIRSTPRIVATE)) + { + OMP_CLAUSE_SET_MAP_KIND (c, GOMP_MAP_FIRSTPRIVATE_POINTER); + goto do_add; + } + } + + enum gomp_map_kind map_kind = (code == OMP_TARGET_EXIT_DATA + ? GOMP_MAP_DETACH + : GOMP_MAP_ATTACH); + OMP_CLAUSE_SET_MAP_KIND (c, map_kind); + } + goto do_add; case OMP_CLAUSE_DEPEND: diff --git a/gcc/omp-low.c b/gcc/omp-low.c index ea9008b61c4..447d7dbc92a 100644 --- a/gcc/omp-low.c +++ b/gcc/omp-low.c @@ -214,6 +214,21 @@ is_oacc_kernels (omp_context *ctx) == GF_OMP_TARGET_KIND_OACC_KERNELS)); } +/* Return true if STMT corresponds to an OpenMP target region. */ +static bool +is_omp_target (gimple *stmt) +{ + if (gimple_code (stmt) == GIMPLE_OMP_TARGET) + { + int kind = gimple_omp_target_kind (stmt); + return (kind == GF_OMP_TARGET_KIND_REGION + || kind == GF_OMP_TARGET_KIND_DATA + || kind == GF_OMP_TARGET_KIND_ENTER_DATA + || kind == GF_OMP_TARGET_KIND_EXIT_DATA); + } + return false; +} + /* If DECL is the artificial dummy VAR_DECL created for non-static data member privatization, return the underlying "this" parameter, otherwise return NULL. */ @@ -1346,7 +1361,9 @@ scan_sharing_clauses (tree clauses, omp_context *ctx) && DECL_P (decl) && ((OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_FIRSTPRIVATE_POINTER && (OMP_CLAUSE_MAP_KIND (c) - != GOMP_MAP_FIRSTPRIVATE_REFERENCE)) + != GOMP_MAP_FIRSTPRIVATE_REFERENCE) + && OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_ATTACH + && OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_DETACH) || TREE_CODE (TREE_TYPE (decl)) == ARRAY_TYPE) && OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_ALWAYS_TO && OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_ALWAYS_FROM @@ -1367,6 +1384,40 @@ scan_sharing_clauses (tree clauses, omp_context *ctx) && !OMP_CLAUSE_MAP_ZERO_BIAS_ARRAY_SECTION (c)) break; } + if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP + && DECL_P (decl) + && (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ATTACH + || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_DETACH) + && is_omp_target (ctx->stmt)) + { + /* If this is an offloaded region, an attach operation should + only exist when the pointer variable is mapped in a prior + clause. */ + if (is_gimple_omp_offloaded (ctx->stmt)) + gcc_assert + (maybe_lookup_decl (decl, ctx) + || (is_global_var (maybe_lookup_decl_in_outer_ctx (decl, ctx)) + && lookup_attribute ("omp declare target", + DECL_ATTRIBUTES (decl)))); + + /* By itself, attach/detach is generated as part of pointer + variable mapping and should not create new variables in the + offloaded region, however sender refs for it must be created + for its address to be passed to the runtime. */ + tree field + = build_decl (OMP_CLAUSE_LOCATION (c), + FIELD_DECL, NULL_TREE, ptr_type_node); + SET_DECL_ALIGN (field, TYPE_ALIGN (ptr_type_node)); + insert_field_into_struct (ctx->record_type, field); + /* To not clash with a map of the pointer variable itself, + attach/detach maps have their field looked up by the *clause* + tree expression, not the decl. */ + gcc_assert (!splay_tree_lookup (ctx->field_map, + (splay_tree_key) c)); + splay_tree_insert (ctx->field_map, (splay_tree_key) c, + (splay_tree_value) field); + break; + } if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP && (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_FIRSTPRIVATE_POINTER || (OMP_CLAUSE_MAP_KIND (c) @@ -1607,6 +1658,11 @@ scan_sharing_clauses (tree clauses, omp_context *ctx) && is_global_var (maybe_lookup_decl_in_outer_ctx (decl, ctx)) && varpool_node::get_create (decl)->offloadable) break; + if ((OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ATTACH + || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_DETACH) + && is_omp_target (ctx->stmt) + && !is_gimple_omp_offloaded (ctx->stmt)) + break; if (DECL_P (decl)) { if ((OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_POINTER @@ -11471,6 +11527,8 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx) case GOMP_MAP_FIRSTPRIVATE_REFERENCE: case GOMP_MAP_STRUCT: case GOMP_MAP_ALWAYS_POINTER: + case GOMP_MAP_ATTACH: + case GOMP_MAP_DETACH: break; case GOMP_MAP_IF_PRESENT: case GOMP_MAP_FORCE_ALLOC: @@ -11481,8 +11539,6 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx) case GOMP_MAP_FORCE_DEVICEPTR: case GOMP_MAP_DEVICE_RESIDENT: case GOMP_MAP_LINK: - case GOMP_MAP_ATTACH: - case GOMP_MAP_DETACH: case GOMP_MAP_FORCE_DETACH: gcc_assert (is_gimple_omp_oacc (stmt)); break; @@ -11537,6 +11593,16 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx) continue; } + if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP + && (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ATTACH + || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_DETACH) + && is_omp_target (stmt)) + { + gcc_assert (maybe_lookup_field (c, ctx)); + map_cnt++; + continue; + } + if (!maybe_lookup_field (var, ctx)) continue; @@ -11769,14 +11835,28 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx) gcc_assert (DECL_P (ovar2)); ovar = ovar2; } - if (!maybe_lookup_field (ovar, ctx)) + if (!maybe_lookup_field (ovar, ctx) + && !(OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP + && (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ATTACH + || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_DETACH))) continue; } talign = TYPE_ALIGN_UNIT (TREE_TYPE (ovar)); if (DECL_P (ovar) && DECL_ALIGN_UNIT (ovar) > talign) talign = DECL_ALIGN_UNIT (ovar); - if (nc) + + if (nc + && OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP + && (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ATTACH + || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_DETACH) + && is_omp_target (stmt)) + { + var = lookup_decl_in_outer_ctx (ovar, ctx); + x = build_sender_ref (c, ctx); + gimplify_assign (x, build_fold_addr_expr (var), &ilist); + } + else if (nc) { var = lookup_decl_in_outer_ctx (ovar, ctx); x = build_sender_ref (ovar, ctx); diff --git a/gcc/testsuite/c-c++-common/gomp/clauses-2.c b/gcc/testsuite/c-c++-common/gomp/clauses-2.c index ded1d74ccde..bbc8fb4e32b 100644 --- a/gcc/testsuite/c-c++-common/gomp/clauses-2.c +++ b/gcc/testsuite/c-c++-common/gomp/clauses-2.c @@ -13,35 +13,35 @@ foo (int *p, int q, struct S t, int i, int j, int k, int l) bar (p); #pragma omp target map (p[0]) map (p) /* { dg-error "appears both in data and map clauses" } */ bar (p); - #pragma omp target map (p) , map (p[0]) /* { dg-error "appears both in data and map clauses" } */ + #pragma omp target map (p) , map (p[0]) bar (p); #pragma omp target map (q) map (q) /* { dg-error "appears more than once in map clauses" } */ bar (&q); #pragma omp target map (p[0]) map (p[0]) /* { dg-error "appears more than once in data clauses" } */ bar (p); - #pragma omp target map (t) map (t.r) /* { dg-error "appears more than once in map clauses" } */ + #pragma omp target map (t) map (t.r) bar (&t.r); - #pragma omp target map (t.r) map (t) /* { dg-error "appears more than once in map clauses" } */ + #pragma omp target map (t.r) map (t) bar (&t.r); - #pragma omp target map (t.r) map (t.r) /* { dg-error "appears more than once in map clauses" } */ + #pragma omp target map (t.r) map (t.r) bar (&t.r); #pragma omp target firstprivate (t), map (t.r) /* { dg-error "appears both in data and map clauses" } */ bar (&t.r); #pragma omp target map (t.r) firstprivate (t) /* { dg-error "appears both in data and map clauses" } */ bar (&t.r); - #pragma omp target map (t.s[0]) map (t) /* { dg-error "appears more than once in map clauses" } */ + #pragma omp target map (t.s[0]) map (t) bar (t.s); - #pragma omp target map (t) map(t.s[0]) /* { dg-error "appears more than once in map clauses" } */ + #pragma omp target map (t) map(t.s[0]) bar (t.s); #pragma omp target firstprivate (t) map (t.s[0]) /* { dg-error "appears both in data and map clauses" } */ bar (t.s); #pragma omp target map (t.s[0]) firstprivate (t) /* { dg-error "appears both in data and map clauses" } */ bar (t.s); - #pragma omp target map (t.s[0]) map (t.s[2]) /* { dg-error "appears more than once in map clauses" } */ + #pragma omp target map (t.s[0]) map (t.s[2]) bar (t.s); - #pragma omp target map (t.t[0:2]) map (t.t[4:6]) /* { dg-error "appears more than once in map clauses" } */ + #pragma omp target map (t.t[0:2]) map (t.t[4:6]) bar (t.t); - #pragma omp target map (t.t[i:j]) map (t.t[k:l]) /* { dg-error "appears more than once in map clauses" } */ + #pragma omp target map (t.t[i:j]) map (t.t[k:l]) bar (t.t); #pragma omp target map (t.s[0]) map (t.r) bar (t.s); @@ -50,5 +50,5 @@ foo (int *p, int q, struct S t, int i, int j, int k, int l) #pragma omp target map (t.r) map (t) map (t.s[0]) firstprivate (t) /* { dg-error "appears both in data and map clauses" } */ bar (t.s); #pragma omp target map (t) map (t.r) firstprivate (t) map (t.s[0]) /* { dg-error "appears both in data and map clauses" } */ - bar (t.s); /* { dg-error "appears more than once in map clauses" "" { target *-*-* } .-1 } */ + bar (t.s); } diff --git a/gcc/testsuite/c-c++-common/gomp/map-5.c b/gcc/testsuite/c-c++-common/gomp/map-5.c new file mode 100644 index 00000000000..1d9d9252864 --- /dev/null +++ b/gcc/testsuite/c-c++-common/gomp/map-5.c @@ -0,0 +1,24 @@ +/* { dg-do compile } */ +/* { dg-additional-options "-fdump-tree-gimple" } */ + +void foo (void) +{ + /* Basic test to ensure to,from,tofrom is ordered before alloc,release,delete clauses. */ + int a, b, c; + #pragma omp target enter data map(alloc:a) map(to:b) map(alloc:c) + #pragma omp target exit data map(from:a) map(release:b) map(from:c) + + #pragma omp target map(alloc:a) map(tofrom:b) map(alloc:c) + a = b = c = 1; + + #pragma omp target enter data map(to:a) map(alloc:b) map(to:c) + #pragma omp target exit data map(from:a) map(delete:b) map(from:c) +} + +/* { dg-final { scan-tree-dump "pragma omp target enter data map\\(to:.* map\\(alloc:.* map\\(alloc:.*" "gimple" } } */ +/* { dg-final { scan-tree-dump "pragma omp target exit data map\\(from:.* map\\(from:.* map\\(release:.*" "gimple" } } */ + +/* { dg-final { scan-tree-dump "pragma omp target num_teams.* map\\(tofrom:.* map\\(alloc:.* map\\(alloc:.*" "gimple" } } */ + +/* { dg-final { scan-tree-dump "pragma omp target enter data map\\(to:.* map\\(to:.* map\\(alloc:.*" "gimple" } } */ +/* { dg-final { scan-tree-dump "pragma omp target exit data map\\(from:.* map\\(from:.* map\\(delete:.*" "gimple" } } */ diff --git a/gcc/testsuite/gfortran.dg/gomp/map-2.f90 b/gcc/testsuite/gfortran.dg/gomp/map-2.f90 index 73c4f5a87d0..79bab726dea 100644 --- a/gcc/testsuite/gfortran.dg/gomp/map-2.f90 +++ b/gcc/testsuite/gfortran.dg/gomp/map-2.f90 @@ -2,5 +2,5 @@ type t integer :: i end type t type(t) v -!$omp target enter data map(to:v%i, v%i) ! { dg-error "appears more than once in map clauses" } +!$omp target enter data map(to:v%i, v%i) end diff --git a/libgomp/libgomp.h b/libgomp/libgomp.h index da7ac037dcd..0cc3f4d406b 100644 --- a/libgomp/libgomp.h +++ b/libgomp/libgomp.h @@ -1162,10 +1162,10 @@ struct gomp_device_descr /* Kind of the pragma, for which gomp_map_vars () is called. */ enum gomp_map_vars_kind { - GOMP_MAP_VARS_OPENACC, - GOMP_MAP_VARS_TARGET, - GOMP_MAP_VARS_DATA, - GOMP_MAP_VARS_ENTER_DATA + GOMP_MAP_VARS_OPENACC = 1, + GOMP_MAP_VARS_TARGET = 2, + GOMP_MAP_VARS_DATA = 4, + GOMP_MAP_VARS_ENTER_DATA = 8 }; extern void gomp_acc_declare_allocate (bool, size_t, void **, size_t *, diff --git a/libgomp/oacc-mem.c b/libgomp/oacc-mem.c index 65757ab2ffc..4c8f0e0828e 100644 --- a/libgomp/oacc-mem.c +++ b/libgomp/oacc-mem.c @@ -403,7 +403,8 @@ acc_map_data (void *h, void *d, size_t s) struct target_mem_desc *tgt = gomp_map_vars (acc_dev, mapnum, &hostaddrs, &devaddrs, &sizes, - &kinds, true, GOMP_MAP_VARS_ENTER_DATA); + &kinds, true, + GOMP_MAP_VARS_OPENACC | GOMP_MAP_VARS_ENTER_DATA); assert (tgt); assert (tgt->list_count == 1); splay_tree_key n = tgt->list[0].key; @@ -572,7 +573,8 @@ goacc_enter_datum (void **hostaddrs, size_t *sizes, void *kinds, int async) struct target_mem_desc *tgt = gomp_map_vars_async (acc_dev, aq, mapnum, hostaddrs, NULL, sizes, - kinds, true, GOMP_MAP_VARS_ENTER_DATA); + kinds, true, (GOMP_MAP_VARS_OPENACC + | GOMP_MAP_VARS_ENTER_DATA)); assert (tgt); assert (tgt->list_count == 1); n = tgt->list[0].key; @@ -1202,7 +1204,8 @@ goacc_enter_data_internal (struct gomp_device_descr *acc_dev, size_t mapnum, struct target_mem_desc *tgt = gomp_map_vars_async (acc_dev, aq, groupnum, &hostaddrs[i], NULL, &sizes[i], &kinds[i], true, - GOMP_MAP_VARS_ENTER_DATA); + (GOMP_MAP_VARS_OPENACC + | GOMP_MAP_VARS_ENTER_DATA)); assert (tgt); gomp_mutex_lock (&acc_dev->lock); diff --git a/libgomp/target.c b/libgomp/target.c index 3432a835369..6152f58e13d 100644 --- a/libgomp/target.c +++ b/libgomp/target.c @@ -683,7 +683,7 @@ gomp_map_vars_internal (struct gomp_device_descr *devicep, struct target_mem_desc *tgt = gomp_malloc (sizeof (*tgt) + sizeof (tgt->list[0]) * mapnum); tgt->list_count = mapnum; - tgt->refcount = pragma_kind == GOMP_MAP_VARS_ENTER_DATA ? 0 : 1; + tgt->refcount = (pragma_kind & GOMP_MAP_VARS_ENTER_DATA) ? 0 : 1; tgt->device_descr = devicep; tgt->prev = NULL; struct gomp_coalesce_buf cbuf, *cbufp = NULL; @@ -1212,15 +1212,16 @@ gomp_map_vars_internal (struct gomp_device_descr *devicep, /* OpenACC 'attach'/'detach' doesn't affect structured/dynamic reference counts ('n->refcount', 'n->dynamic_refcount'). */ + + gomp_attach_pointer (devicep, aq, mem_map, n, + (uintptr_t) hostaddrs[i], sizes[i], + cbufp); } - else + else if ((pragma_kind & GOMP_MAP_VARS_OPENACC) != 0) { gomp_mutex_unlock (&devicep->lock); gomp_fatal ("outer struct not mapped for attach"); } - gomp_attach_pointer (devicep, aq, mem_map, n, - (uintptr_t) hostaddrs[i], sizes[i], - cbufp); continue; } default: @@ -1415,7 +1416,7 @@ gomp_map_vars_internal (struct gomp_device_descr *devicep, /* If the variable from "omp target enter data" map-list was already mapped, tgt is not needed. Otherwise tgt will be freed by gomp_unmap_vars or gomp_exit_data. */ - if (pragma_kind == GOMP_MAP_VARS_ENTER_DATA && tgt->refcount == 0) + if ((pragma_kind & GOMP_MAP_VARS_ENTER_DATA) && tgt->refcount == 0) { free (tgt); tgt = NULL; @@ -2475,6 +2476,19 @@ gomp_exit_data (struct gomp_device_descr *devicep, size_t mapnum, return; } + for (i = 0; i < mapnum; i++) + if ((kinds[i] & typemask) == GOMP_MAP_DETACH) + { + struct splay_tree_key_s cur_node; + cur_node.host_start = (uintptr_t) hostaddrs[i]; + cur_node.host_end = cur_node.host_start + sizeof (void *); + splay_tree_key n = splay_tree_lookup (&devicep->mem_map, &cur_node); + + if (n) + gomp_detach_pointer (devicep, NULL, n, (uintptr_t) hostaddrs[i], + false, NULL); + } + for (i = 0; i < mapnum; i++) { struct splay_tree_key_s cur_node; @@ -2512,7 +2526,9 @@ gomp_exit_data (struct gomp_device_descr *devicep, size_t mapnum, cur_node.host_end - cur_node.host_start); if (k->refcount == 0) gomp_remove_var (devicep, k); + break; + case GOMP_MAP_DETACH: break; default: gomp_mutex_unlock (&devicep->lock); @@ -2621,6 +2637,14 @@ GOMP_target_enter_exit_data (int device, size_t mapnum, void **hostaddrs, &kinds[i], true, GOMP_MAP_VARS_ENTER_DATA); i += j - i - 1; } + else if (i + 1 < mapnum && (kinds[i + 1] & 0xff) == GOMP_MAP_ATTACH) + { + /* An attach operation must be processed together with the mapped + base-pointer list item. */ + gomp_map_vars (devicep, 2, &hostaddrs[i], NULL, &sizes[i], &kinds[i], + true, GOMP_MAP_VARS_ENTER_DATA); + i += 1; + } else gomp_map_vars (devicep, 1, &hostaddrs[i], NULL, &sizes[i], &kinds[i], true, GOMP_MAP_VARS_ENTER_DATA); diff --git a/libgomp/testsuite/libgomp.c-c++-common/ptr-attach-1.c b/libgomp/testsuite/libgomp.c-c++-common/ptr-attach-1.c new file mode 100644 index 00000000000..e7deec6e006 --- /dev/null +++ b/libgomp/testsuite/libgomp.c-c++-common/ptr-attach-1.c @@ -0,0 +1,82 @@ +#include + +struct S +{ + int a, b; + int *ptr; + int c, d; +}; +typedef struct S S; + +#pragma omp declare target +int *gp; +#pragma omp end declare target + +#define N 10 +int main (void) +{ + /* Test to see if pointer attachment works, for scalar pointers, + and pointer fields in structures. */ + + int *ptr = (int *) malloc (sizeof (int) * N); + int *orig_ptr = ptr; + + #pragma omp target map (ptr, ptr[:N]) + { + for (int i = 0; i < N; i++) + ptr[i] = N - i; + } + + if (ptr != orig_ptr) + abort (); + + for (int i = 0; i < N; i++) + if (ptr[i] != N - i) + abort (); + + S s = { 0 }; + s.ptr = ptr; + #pragma omp target map (s, s.ptr[:N]) + { + for (int i = 0; i < N; i++) + s.ptr[i] = i; + + s.a = 1; + s.b = 2; + } + + if (s.ptr != ptr) + abort (); + + for (int i = 0; i < N; i++) + if (s.ptr[i] != i) + abort (); + + if (s.a != 1 || s.b != 2 || s.c != 0 || s.d != 0) + abort (); + + gp = (int *) malloc (sizeof (int) * N); + orig_ptr = gp; + + for (int i = 0; i < N; i++) + gp[i] = i - 1; + + #pragma omp target map (gp[:N]) + { + for (int i = 0; i < N; i++) + gp[i] += 1; + } + + if (gp != orig_ptr) + abort (); + + for (int i = 0; i < N; i++) + if (gp[i] != i) + abort (); + + free (ptr); + free (gp); + + return 0; +} +