From: Chung-Lin Tang Date: Tue, 2 Feb 2021 12:31:37 +0000 (+0800) Subject: OpenMP 5.0: map this[:1] in C++ non-static member functions (PR 92120) X-Git-Url: http://git.ipfire.org/cgi-bin/gitweb.cgi?a=commitdiff_plain;h=36a1ebdb650657c10be1280b547e68c1833006f4;p=thirdparty%2Fgcc.git OpenMP 5.0: map this[:1] in C++ non-static member functions (PR 92120) This is a merge of: https://gcc.gnu.org/pipermail/gcc-patches/2020-November/558975.html This patch creates automatic mapping of map(this[:1]) and pointer members as zero-length array sections, as specified by the OpenMP 5.0 specification. This may possibly reverted/updated when a final patch is approved for mainline. 2021-02-02 Chung-Lin Tang PR middle-end/92120 gcc/cp/ChangeLog: * cp-tree.h (finish_omp_target): New declaration. (set_omp_target_this_expr): Likewise. * lambda.c (lambda_expr_this_capture): Add call to set_omp_target_this_expr. * parser.c (cp_parser_omp_target): Factor out code, change to call finish_omp_target, add re-initing call to set_omp_target_this_expr. * semantics.c (omp_target_this_expr): New static variable. (omp_target_ptr_members_accessed): New static hash_map for tracking accessed non-static pointer-type members. (finish_non_static_data_member): Add call to set_omp_target_this_expr. Add recording of non-static pointer-type members access. (finish_this_expr): Add call to set_omp_target_this_expr. (set_omp_target_this_expr): New function to set omp_target_this_expr. (finish_omp_target): New function with code merged from cp_parser_omp_target, plus code to implement this[:1] and __closure map clauses for OpenMP. (handle_omp_array_sections_1): Move code to peel of '*' for reference-based COMPONENT_REFs before FIELD_DECL transforming. (finish_omp_clauses): Handle 'A->member' case in map clauses. gcc/ChangeLog: * omp-low.c (lower_omp_target): Handle GOMP_MAP_ATTACH_ZERO_LENGTH_ARRAY_SECTION, and GOMP_MAP_POINTER_TO_ZERO_LENGTH_ARRAY_SECTION map kinds. * tree-pretty-print.c (dump_omp_clause): Likewise. gcc/testsuite/ChangeLog: * g++.dg/gomp/target-this-1.C: New testcase. * g++.dg/gomp/target-this-2.C: New testcase. * g++.dg/gomp/target-this-3.C: New testcase. * g++.dg/gomp/target-this-4.C: New testcase. include/ChangeLog: * gomp-constants.h (enum gomp_map_kind): Add GOMP_MAP_ATTACH_ZERO_LENGTH_ARRAY_SECTION, and GOMP_MAP_POINTER_TO_ZERO_LENGTH_ARRAY_SECTION map kinds. (GOMP_MAP_POINTER_P): Include GOMP_MAP_POINTER_TO_ZERO_LENGTH_ARRAY_SECTION. libgomp/ChangeLog: * libgomp.h (gomp_attach_pointer): Add bool parameter. * oacc-mem.c (acc_attach_async): Update call to gomp_attach_pointer. (goacc_enter_data_internal): Likewise. * target.c (gomp_map_vars_existing): Update assert condition to include GOMP_MAP_ATTACH_ZERO_LENGTH_ARRAY_SECTION. (gomp_map_pointer): Add 'bool allow_zero_length_array_sections' parameter, add support for mapping a pointer with NULL target. (gomp_attach_pointer): Add 'bool allow_zero_length_array_sections' parameter, add support for attaching a pointer with NULL target. (gomp_map_vars_internal): Update calls to gomp_map_pointer and gomp_attach_pointer, add handling for GOMP_MAP_ATTACH_ZERO_LENGTH_ARRAY_SECTION, and GOMP_MAP_POINTER_TO_ZERO_LENGTH_ARRAY_SECTION cases. * testsuite/libgomp.c++/target-this-1.C: New testcase. * testsuite/libgomp.c++/target-this-2.C: New testcase. * testsuite/libgomp.c++/target-this-3.C: New testcase. * testsuite/libgomp.c++/target-this-4.C: New testcase. --- diff --git a/gcc/cp/cp-tree.h b/gcc/cp/cp-tree.h index af7df8687a18..b77bdc380a05 100644 --- a/gcc/cp/cp-tree.h +++ b/gcc/cp/cp-tree.h @@ -7315,6 +7315,8 @@ extern void record_null_lambda_scope (tree); extern void finish_lambda_scope (void); extern tree start_lambda_function (tree fn, tree lambda_expr); extern void finish_lambda_function (tree body); +extern tree finish_omp_target (location_t, tree, tree, bool); +extern void set_omp_target_this_expr (tree); /* in tree.c */ extern int cp_tree_operand_length (const_tree); diff --git a/gcc/cp/lambda.c b/gcc/cp/lambda.c index b55c2f85d27e..9ecf0dbed0c7 100644 --- a/gcc/cp/lambda.c +++ b/gcc/cp/lambda.c @@ -842,6 +842,9 @@ lambda_expr_this_capture (tree lambda, int add_capture_p) type cast (_expr.cast_ 5.4) to the type of 'this'. [ The cast ensures that the transformed expression is an rvalue. ] */ result = rvalue (result); + + /* Acknowledge to OpenMP target that 'this' was referenced. */ + set_omp_target_this_expr (result); } return result; diff --git a/gcc/cp/parser.c b/gcc/cp/parser.c index cea6ed0b2c19..6eb228099bc7 100644 --- a/gcc/cp/parser.c +++ b/gcc/cp/parser.c @@ -40717,8 +40717,6 @@ static bool cp_parser_omp_target (cp_parser *parser, cp_token *pragma_tok, enum pragma_context context, bool *if_p) { - tree *pc = NULL, stmt; - if (flag_openmp) omp_requires_mask = (enum omp_requires) (omp_requires_mask | OMP_REQUIRES_TARGET_USED); @@ -40771,6 +40769,7 @@ cp_parser_omp_target (cp_parser *parser, cp_token *pragma_tok, keep_next_level (true); tree sb = begin_omp_structured_block (), ret; unsigned save = cp_parser_begin_omp_structured_block (parser); + set_omp_target_this_expr (NULL_TREE); switch (ccode) { case OMP_TEAMS: @@ -40822,15 +40821,9 @@ cp_parser_omp_target (cp_parser *parser, cp_token *pragma_tok, cclauses[C_OMP_CLAUSE_SPLIT_TARGET] = tc; } } - tree stmt = make_node (OMP_TARGET); - TREE_TYPE (stmt) = void_type_node; - OMP_TARGET_CLAUSES (stmt) = cclauses[C_OMP_CLAUSE_SPLIT_TARGET]; - OMP_TARGET_BODY (stmt) = body; - OMP_TARGET_COMBINED (stmt) = 1; - SET_EXPR_LOCATION (stmt, pragma_tok->location); - add_stmt (stmt); - pc = &OMP_TARGET_CLAUSES (stmt); - goto check_clauses; + finish_omp_target (pragma_tok->location, + cclauses[C_OMP_CLAUSE_SPLIT_TARGET], body, true); + return true; } else if (!flag_openmp) /* flag_openmp_simd */ { @@ -40867,49 +40860,14 @@ cp_parser_omp_target (cp_parser *parser, cp_token *pragma_tok, return false; } - stmt = make_node (OMP_TARGET); - TREE_TYPE (stmt) = void_type_node; - - 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); + tree clauses = cp_parser_omp_all_clauses (parser, OMP_TARGET_CLAUSE_MASK, + "#pragma omp target", pragma_tok); + c_omp_adjust_map_clauses (clauses, true); keep_next_level (true); - OMP_TARGET_BODY (stmt) = cp_parser_omp_structured_block (parser, if_p); + set_omp_target_this_expr (NULL_TREE); + tree body = cp_parser_omp_structured_block (parser, if_p); - SET_EXPR_LOCATION (stmt, pragma_tok->location); - add_stmt (stmt); - -check_clauses: - while (*pc) - { - if (OMP_CLAUSE_CODE (*pc) == OMP_CLAUSE_MAP) - switch (OMP_CLAUSE_MAP_KIND (*pc)) - { - case GOMP_MAP_TO: - case GOMP_MAP_ALWAYS_TO: - case GOMP_MAP_FROM: - case GOMP_MAP_ALWAYS_FROM: - case GOMP_MAP_TOFROM: - case GOMP_MAP_ALWAYS_TOFROM: - case GOMP_MAP_ALLOC: - 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), - "%<#pragma omp target%> with map-type other " - "than %, %, % or % " - "on % clause"); - *pc = OMP_CLAUSE_CHAIN (*pc); - continue; - } - pc = &OMP_CLAUSE_CHAIN (*pc); - } + finish_omp_target (pragma_tok->location, clauses, body, false); return true; } diff --git a/gcc/cp/semantics.c b/gcc/cp/semantics.c index 30d94aae960f..9eda407bbf17 100644 --- a/gcc/cp/semantics.c +++ b/gcc/cp/semantics.c @@ -61,6 +61,10 @@ static hash_map *omp_private_member_map; static vec omp_private_member_vec; static bool omp_private_member_ignore_next; +/* Used for OpenMP target region 'this' references. */ +static tree omp_target_this_expr = NULL_TREE; + +static hash_map omp_target_ptr_members_accessed; /* Deferred Access Checking Overview --------------------------------- @@ -1892,6 +1896,7 @@ tree finish_non_static_data_member (tree decl, tree object, tree qualifying_scope) { gcc_assert (TREE_CODE (decl) == FIELD_DECL); + tree orig_object = object; bool try_omp_private = !object && omp_private_member_map; tree ret; @@ -1930,6 +1935,14 @@ finish_non_static_data_member (tree decl, tree object, tree qualifying_scope) return error_mark_node; } + if (orig_object == NULL_TREE) + { + tree this_expr = TREE_OPERAND (object, 0); + + /* Acknowledge to OpenMP target that 'this' was referenced. */ + set_omp_target_this_expr (this_expr); + } + if (current_class_ptr) TREE_USED (current_class_ptr) = 1; if (processing_template_decl) @@ -1990,6 +2003,14 @@ finish_non_static_data_member (tree decl, tree object, tree qualifying_scope) if (v) ret = convert_from_reference (*v); } + else if (omp_target_this_expr + && TREE_TYPE (ret) + && POINTER_TYPE_P (TREE_TYPE (ret))) + { + if (omp_target_ptr_members_accessed.get (decl) == NULL) + omp_target_ptr_members_accessed.put (decl, ret); + } + return ret; } @@ -2748,8 +2769,15 @@ finish_this_expr (void) } if (result) - /* The keyword 'this' is a prvalue expression. */ - return rvalue (result); + { + /* The keyword 'this' is a prvalue expression. */ + result = rvalue (result); + + /* Acknowledge to OpenMP target that 'this' was referenced. */ + set_omp_target_this_expr (result); + + return result; + } tree fn = current_nonlambda_function (); if (fn && DECL_STATIC_FUNCTION_P (fn)) @@ -4751,12 +4779,12 @@ handle_omp_array_sections_1 (tree c, tree t, vec &types, { if (error_operand_p (t)) return error_mark_node; - if (REFERENCE_REF_P (t) - && TREE_CODE (TREE_OPERAND (t, 0)) == COMPONENT_REF) - t = TREE_OPERAND (t, 0); if ((ort == C_ORT_ACC || ort == C_ORT_OMP) && TREE_CODE (t) == FIELD_DECL) t = finish_non_static_data_member (t, NULL_TREE, NULL_TREE); + if (REFERENCE_REF_P (t) + && TREE_CODE (TREE_OPERAND (t, 0)) == COMPONENT_REF) + t = TREE_OPERAND (t, 0); ret = t; if (TREE_CODE (t) == COMPONENT_REF && (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP @@ -5386,6 +5414,8 @@ handle_omp_array_sections (tree c, enum c_omp_region_type ort) } tree c2 = build_omp_clause (OMP_CLAUSE_LOCATION (c), OMP_CLAUSE_MAP); + if (TREE_CODE (t) == FIELD_DECL) + t = finish_non_static_data_member (t, NULL_TREE, NULL_TREE); 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) @@ -6377,6 +6407,7 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort) bool order_seen = false; bool schedule_seen = false; bool oacc_async = false; + bool indir_component_ref_p = false; tree last_iterators = NULL_TREE; bool last_iterators_remove = false; /* 1 if normal/task reduction has been seen, -1 if inscan reduction @@ -7476,10 +7507,15 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort) t = TREE_OPERAND (t, 0); OMP_CLAUSE_DECL (c) = t; } + indir_component_ref_p = false; 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); + { + t = TREE_OPERAND (TREE_OPERAND (t, 0), 0); + indir_component_ref_p = true; + STRIP_NOPS (t); + } if (TREE_CODE (t) == COMPONENT_REF && ((ort & C_ORT_OMP_DECLARE_SIMD) == C_ORT_OMP || ort == C_ORT_ACC) @@ -7577,6 +7613,7 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort) && (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_MAP || (OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_FIRSTPRIVATE_POINTER)) + && !indir_component_ref_p && !cxx_mark_addressable (t)) remove = true; else if (!(OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP @@ -8629,6 +8666,256 @@ finish_omp_construct (enum tree_code code, tree body, tree clauses) return add_stmt (stmt); } +void +set_omp_target_this_expr (tree this_val) +{ + omp_target_this_expr = this_val; + + if (omp_target_this_expr == NULL_TREE) + omp_target_ptr_members_accessed.empty (); +} + +tree +finish_omp_target (location_t loc, tree clauses, tree body, bool combined_p) +{ + tree last_inserted_clause = NULL_TREE; + + if (omp_target_this_expr) + { + /* See if explicit user-specified map(this[:]) clause already exists. + If not, we create an implicit map(tofrom:this[:1]) clause. */ + tree *explicit_this_deref_map = NULL; + for (tree *c = &clauses; *c; c = &OMP_CLAUSE_CHAIN (*c)) + if (OMP_CLAUSE_CODE (*c) == OMP_CLAUSE_MAP + && TREE_CODE (OMP_CLAUSE_DECL (*c)) == INDIRECT_REF + && operand_equal_p (TREE_OPERAND (OMP_CLAUSE_DECL (*c), 0), + omp_target_this_expr)) + { + explicit_this_deref_map = c; + break; + } + + if (DECL_LAMBDA_FUNCTION_P (current_function_decl)) + { + /* For lambda functions, we need to first create a copy of the + __closure object. */ + tree closure = DECL_ARGUMENTS (current_function_decl); + tree c = build_omp_clause (loc, OMP_CLAUSE_MAP); + OMP_CLAUSE_SET_MAP_KIND (c, GOMP_MAP_TO); + OMP_CLAUSE_DECL (c) = build_simple_mem_ref (closure); + OMP_CLAUSE_SIZE (c) + = TYPE_SIZE_UNIT (TREE_TYPE (TREE_TYPE (closure))); + + tree c2 = build_omp_clause (loc, OMP_CLAUSE_MAP); + OMP_CLAUSE_SET_MAP_KIND (c2, GOMP_MAP_FIRSTPRIVATE_POINTER); + OMP_CLAUSE_DECL (c2) = closure; + OMP_CLAUSE_SIZE (c2) = size_zero_node; + OMP_CLAUSE_CHAIN (c2) = clauses; + OMP_CLAUSE_CHAIN (c) = c2; + last_inserted_clause = c2; + clauses = c; + + STRIP_NOPS (omp_target_this_expr); + gcc_assert (DECL_HAS_VALUE_EXPR_P (omp_target_this_expr)); + omp_target_this_expr = DECL_VALUE_EXPR (omp_target_this_expr); + + if (explicit_this_deref_map) + { + /* Transform *this into *__closure->this in maps. */ + tree this_map = *explicit_this_deref_map; + OMP_CLAUSE_DECL (this_map) + = build_simple_mem_ref (omp_target_this_expr); + tree nc = OMP_CLAUSE_CHAIN (this_map); + gcc_assert (OMP_CLAUSE_CODE (nc) == OMP_CLAUSE_MAP + && (OMP_CLAUSE_MAP_KIND (nc) + == GOMP_MAP_FIRSTPRIVATE_POINTER)); + OMP_CLAUSE_DECL (nc) = omp_target_this_expr; + OMP_CLAUSE_SET_MAP_KIND (nc, GOMP_MAP_ALWAYS_POINTER); + + /* Move map(*__closure->this) map(always_pointer:__closure->this) + sequence to right after __closure map. */ + *explicit_this_deref_map = OMP_CLAUSE_CHAIN (nc); + OMP_CLAUSE_CHAIN (nc) = OMP_CLAUSE_CHAIN (c2); + OMP_CLAUSE_CHAIN (c2) = this_map; + last_inserted_clause = nc; + } + else + { + tree c3 = build_omp_clause (loc, OMP_CLAUSE_MAP); + OMP_CLAUSE_SET_MAP_KIND (c3, GOMP_MAP_TOFROM); + OMP_CLAUSE_DECL (c3) + = build_simple_mem_ref (omp_target_this_expr); + OMP_CLAUSE_SIZE (c3) + = TYPE_SIZE_UNIT (TREE_TYPE (TREE_TYPE (omp_target_this_expr))); + + tree c4 = build_omp_clause (loc, OMP_CLAUSE_MAP); + OMP_CLAUSE_SET_MAP_KIND (c4, GOMP_MAP_ALWAYS_POINTER); + + OMP_CLAUSE_DECL (c4) = omp_target_this_expr; + OMP_CLAUSE_SIZE (c4) = size_zero_node; + + OMP_CLAUSE_CHAIN (c3) = c4; + OMP_CLAUSE_CHAIN (c4) = OMP_CLAUSE_CHAIN (c2); + OMP_CLAUSE_CHAIN (c2) = c3; + last_inserted_clause = c4; + } + } + else + { + /* For the non-lambda case, we only need to create map(this[:1]) when + it's not present, no transforming needed. */ + if (!explicit_this_deref_map) + { + tree c = build_omp_clause (loc, OMP_CLAUSE_MAP); + OMP_CLAUSE_SET_MAP_KIND (c, GOMP_MAP_TOFROM); + OMP_CLAUSE_DECL (c) = build_simple_mem_ref (omp_target_this_expr); + OMP_CLAUSE_SIZE (c) + = TYPE_SIZE_UNIT (TREE_TYPE (TREE_TYPE (omp_target_this_expr))); + + tree c2 = build_omp_clause (loc, OMP_CLAUSE_MAP); + OMP_CLAUSE_SET_MAP_KIND (c2, GOMP_MAP_FIRSTPRIVATE_POINTER); + STRIP_NOPS (omp_target_this_expr); + OMP_CLAUSE_DECL (c2) = omp_target_this_expr; + OMP_CLAUSE_SIZE (c2) = size_zero_node; + OMP_CLAUSE_CHAIN (c2) = clauses; + OMP_CLAUSE_CHAIN (c) = c2; + clauses = c; + last_inserted_clause = c2; + } + } + omp_target_this_expr = NULL_TREE; + } + + if (last_inserted_clause && !omp_target_ptr_members_accessed.is_empty ()) + for (hash_map::iterator i + = omp_target_ptr_members_accessed.begin (); + i != omp_target_ptr_members_accessed.end (); ++i) + { + /* For each referenced member that is of pointer or reference-to-pointer + type, create the equivalent of map(alloc:this->ptr[:0]). */ + tree field_decl = (*i).first; + tree ptr_member = (*i).second; + + for (tree nc = OMP_CLAUSE_CHAIN (last_inserted_clause); + nc != NULL_TREE; nc = OMP_CLAUSE_CHAIN (nc)) + { + /* If map(this->ptr[:N] already exists, avoid creating another + such map. */ + tree decl = OMP_CLAUSE_DECL (nc); + if ((TREE_CODE (decl) == INDIRECT_REF + || TREE_CODE (decl) == MEM_REF) + && operand_equal_p (TREE_OPERAND (decl, 0), + ptr_member)) + goto next_ptr_member; + } + + if (!cxx_mark_addressable (ptr_member)) + gcc_unreachable (); + + if (TREE_CODE (TREE_TYPE (field_decl)) == REFERENCE_TYPE) + { + /* For reference to pointers, we need to map the referenced pointer + first for things to be correct. */ + tree ptr_member_type = TREE_TYPE (ptr_member); + + /* Map pointer target as zero-length array section. */ + tree c = build_omp_clause (loc, OMP_CLAUSE_MAP); + OMP_CLAUSE_SET_MAP_KIND (c, GOMP_MAP_ALLOC); + OMP_CLAUSE_DECL (c) + = build1 (INDIRECT_REF, TREE_TYPE (ptr_member_type), ptr_member); + OMP_CLAUSE_SIZE (c) = size_zero_node; + OMP_CLAUSE_MAP_MAYBE_ZERO_LENGTH_ARRAY_SECTION (c) = 1; + + /* Map pointer to zero-length array section. */ + tree c2 = build_omp_clause (loc, OMP_CLAUSE_MAP); + OMP_CLAUSE_SET_MAP_KIND + (c2, GOMP_MAP_POINTER_TO_ZERO_LENGTH_ARRAY_SECTION); + OMP_CLAUSE_DECL (c2) = ptr_member; + OMP_CLAUSE_SIZE (c2) = size_zero_node; + + /* Attach reference-to-pointer field to pointer. */ + tree c3 = build_omp_clause (loc, OMP_CLAUSE_MAP); + OMP_CLAUSE_SET_MAP_KIND (c3, GOMP_MAP_ATTACH); + OMP_CLAUSE_DECL (c3) = TREE_OPERAND (ptr_member, 0); + OMP_CLAUSE_SIZE (c3) = size_zero_node; + + OMP_CLAUSE_CHAIN (c) = c2; + OMP_CLAUSE_CHAIN (c2) = c3; + OMP_CLAUSE_CHAIN (c3) = OMP_CLAUSE_CHAIN (last_inserted_clause); + + OMP_CLAUSE_CHAIN (last_inserted_clause) = c; + last_inserted_clause = c3; + } + else if (TREE_CODE (TREE_TYPE (field_decl)) == POINTER_TYPE) + { + /* Map pointer target as zero-length array section. */ + tree c = build_omp_clause (loc, OMP_CLAUSE_MAP); + OMP_CLAUSE_SET_MAP_KIND (c, GOMP_MAP_ALLOC); + OMP_CLAUSE_DECL (c) + = build2 (MEM_REF, char_type_node, ptr_member, + build_int_cst (build_pointer_type (char_type_node), 0)); + OMP_CLAUSE_SIZE (c) = size_zero_node; + OMP_CLAUSE_MAP_MAYBE_ZERO_LENGTH_ARRAY_SECTION (c) = 1; + + /* Attach zero-length array section to pointer. */ + tree c2 = build_omp_clause (loc, OMP_CLAUSE_MAP); + OMP_CLAUSE_SET_MAP_KIND + (c2, GOMP_MAP_ATTACH_ZERO_LENGTH_ARRAY_SECTION); + OMP_CLAUSE_DECL (c2) = ptr_member; + OMP_CLAUSE_SIZE (c2) = size_zero_node; + + OMP_CLAUSE_CHAIN (c) = c2; + OMP_CLAUSE_CHAIN (c2) = OMP_CLAUSE_CHAIN (last_inserted_clause); + OMP_CLAUSE_CHAIN (last_inserted_clause) = c; + last_inserted_clause = c2; + } + else + gcc_unreachable (); + + next_ptr_member: + ; + } + + tree stmt = make_node (OMP_TARGET); + TREE_TYPE (stmt) = void_type_node; + OMP_TARGET_CLAUSES (stmt) = clauses; + OMP_TARGET_BODY (stmt) = body; + OMP_TARGET_COMBINED (stmt) = combined_p; + SET_EXPR_LOCATION (stmt, loc); + + tree c = clauses; + while (c) + { + if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP) + switch (OMP_CLAUSE_MAP_KIND (c)) + { + case GOMP_MAP_TO: + case GOMP_MAP_ALWAYS_TO: + case GOMP_MAP_FROM: + case GOMP_MAP_ALWAYS_FROM: + case GOMP_MAP_TOFROM: + case GOMP_MAP_ALWAYS_TOFROM: + case GOMP_MAP_ALLOC: + case GOMP_MAP_FIRSTPRIVATE_POINTER: + case GOMP_MAP_FIRSTPRIVATE_REFERENCE: + case GOMP_MAP_ALWAYS_POINTER: + case GOMP_MAP_ATTACH_DETACH: + case GOMP_MAP_ATTACH: + case GOMP_MAP_ATTACH_ZERO_LENGTH_ARRAY_SECTION: + case GOMP_MAP_POINTER_TO_ZERO_LENGTH_ARRAY_SECTION: + break; + default: + error_at (OMP_CLAUSE_LOCATION (c), + "%<#pragma omp target%> with map-type other " + "than %, %, % or % " + "on % clause"); + break; + } + c = OMP_CLAUSE_CHAIN (c); + } + return add_stmt (stmt); +} + tree finish_omp_parallel (tree clauses, tree body) { diff --git a/gcc/omp-low.c b/gcc/omp-low.c index 7888006495e4..c3972383bcf6 100644 --- a/gcc/omp-low.c +++ b/gcc/omp-low.c @@ -12283,6 +12283,8 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx) case GOMP_MAP_ALWAYS_POINTER: case GOMP_MAP_ATTACH: case GOMP_MAP_DETACH: + case GOMP_MAP_ATTACH_ZERO_LENGTH_ARRAY_SECTION: + case GOMP_MAP_POINTER_TO_ZERO_LENGTH_ARRAY_SECTION: break; case GOMP_MAP_IF_PRESENT: case GOMP_MAP_FORCE_ALLOC: diff --git a/gcc/testsuite/g++.dg/gomp/target-this-1.C b/gcc/testsuite/g++.dg/gomp/target-this-1.C new file mode 100644 index 000000000000..de93a3e5e579 --- /dev/null +++ b/gcc/testsuite/g++.dg/gomp/target-this-1.C @@ -0,0 +1,33 @@ +// { dg-do compile } +// { dg-additional-options "-fdump-tree-gimple" } +extern "C" void abort (); + +struct S +{ + int a, b, c, d; + + int sum (void) + { + int val = 0; + val += a + b + this->c + this->d; + return val; + } + + int sum_offload (void) + { + int val = 0; + #pragma omp target map(val) + val += a + b + this->c + this->d; + return val; + } +}; + +int main (void) +{ + S s = { 1, 2, 3, 4 }; + if (s.sum () != s.sum_offload ()) + abort (); + return 0; +} + +/* { dg-final { scan-tree-dump {map\(tofrom:\*this \[len: [0-9]+\]\) map\(firstprivate:this \[pointer assign, bias: 0\]\)} "gimple" } } */ diff --git a/gcc/testsuite/g++.dg/gomp/target-this-2.C b/gcc/testsuite/g++.dg/gomp/target-this-2.C new file mode 100644 index 000000000000..a5e832130fba --- /dev/null +++ b/gcc/testsuite/g++.dg/gomp/target-this-2.C @@ -0,0 +1,49 @@ +// We use 'auto' without a function return type, so specify dialect here +// { dg-do compile } +// { dg-additional-options "-std=c++14 -fdump-tree-gimple" } + +extern "C" void abort (); + +struct T +{ + int x, y; + + auto sum_func (int n) + { + auto fn = [=](int m) -> int + { + int v; + v = (x + y) * n + m; + return v; + }; + return fn; + } + + auto sum_func_offload (int n) + { + auto fn = [=](int m) -> int + { + int v; + #pragma omp target map(from:v) + v = (x + y) * n + m; + return v; + }; + return fn; + } + +}; + +int main (void) +{ + T a = { 1, 2 }; + + auto s1 = a.sum_func (3); + auto s2 = a.sum_func_offload (3); + + if (s1 (1) != s2 (1)) + abort (); + + return 0; +} + +/* { dg-final { scan-tree-dump {map\(to:\*__closure \[len: [0-9]+\]\) map\(firstprivate:__closure \[pointer assign, bias: 0\]\) map\(tofrom:\*_[0-9]+ \[len: [0-9]+\]\) map\(always_pointer:__closure->__this \[pointer assign, bias: 0\]\)} "gimple" } } */ diff --git a/gcc/testsuite/g++.dg/gomp/target-this-3.C b/gcc/testsuite/g++.dg/gomp/target-this-3.C new file mode 100644 index 000000000000..208ea079b95f --- /dev/null +++ b/gcc/testsuite/g++.dg/gomp/target-this-3.C @@ -0,0 +1,105 @@ +// { dg-do compile } +// { dg-additional-options "-fdump-tree-gimple" } +#include +#include +extern "C" void abort (); + +struct S +{ + int * ptr; + int ptr_len; + + int *&refptr; + int refptr_len; + + bool set_ptr (int n) + { + bool mapped; + #pragma omp target map(from:mapped) + { + if (ptr != NULL) + for (int i = 0; i < ptr_len; i++) + ptr[i] = n; + mapped = (ptr != NULL); + } + return mapped; + } + + bool set_refptr (int n) + { + bool mapped; + #pragma omp target map(from:mapped) + { + if (refptr != NULL) + for (int i = 0; i < refptr_len; i++) + refptr[i] = n; + mapped = (refptr != NULL); + } + return mapped; + } +}; + +int main (void) +{ + #define N 10 + int *ptr1 = new int[N]; + int *ptr2 = new int[N]; + + memset (ptr1, 0, sizeof (int) * N); + memset (ptr2, 0, sizeof (int) * N); + + S s = { ptr1, N, ptr2, N }; + + bool mapped; + int val = 123; + + mapped = s.set_ptr (val); + if (mapped) + abort (); + if (s.ptr != ptr1) + abort (); + for (int i = 0; i < N; i++) + if (ptr1[i] != 0) + abort (); + + mapped = s.set_refptr (val); + if (mapped) + abort (); + if (s.refptr != ptr2) + abort (); + for (int i = 0; i < N; i++) + if (ptr2[i] != 0) + abort (); + + #pragma omp target data map(ptr1[:N]) + mapped = s.set_ptr (val); + + if (!mapped) + abort (); + if (s.set_refptr (0)) + abort (); + if (s.ptr != ptr1 || s.refptr != ptr2) + abort (); + for (int i = 0; i < N; i++) + if (ptr1[i] != val) + abort (); + + #pragma omp target data map(ptr2[:N]) + mapped = s.set_refptr (val); + + if (!mapped) + abort (); + if (s.set_ptr (0)) + abort (); + if (s.ptr != ptr1 || s.refptr != ptr2) + abort (); + for (int i = 0; i < N; i++) + if (ptr2[i] != val) + abort (); + + return 0; +} + +/* { dg-final { scan-tree-dump {#pragma omp target num_teams.* map\(tofrom:\*this \[len: [0-9]+\]\) map\(firstprivate:this \[pointer assign, bias: 0\]\) map\(alloc:\*this->refptr \[pointer assign, zero-length array section, bias: 0\]\) map\(attach:this->refptr \[bias: 0\]\) map\(from:mapped \[len: [0-9]+\]\) map\(alloc:\*_[0-9+] \[len: 0\]\) firstprivate\(n\)} "gimple" } } */ + +/* { dg-final { scan-tree-dump {#pragma omp target num_teams.* map\(tofrom:\*this \[len: [0-9]+\]\) map\(firstprivate:this \[pointer assign, bias: 0\]\) map\(attach_zero_length_array_section:this->ptr \[bias: 0\]\) map\(from:mapped \[len: [0-9]+\]\) map\(alloc:MEM.* \[len: 0\]\) firstprivate\(n\)} "gimple" } } */ diff --git a/gcc/testsuite/g++.dg/gomp/target-this-4.C b/gcc/testsuite/g++.dg/gomp/target-this-4.C new file mode 100644 index 000000000000..f42cf3845417 --- /dev/null +++ b/gcc/testsuite/g++.dg/gomp/target-this-4.C @@ -0,0 +1,107 @@ +// We use 'auto' without a function return type, so specify dialect here +// { dg-additional-options "-std=c++14 -fdump-tree-gimple" } +#include +#include + +struct T +{ + int *ptr; + int ptr_len; + + int *&refptr; + int refptr_len; + + auto set_ptr_func (int n) + { + auto fn = [=](void) -> bool + { + bool mapped; + #pragma omp target map(from:mapped) + { + if (ptr) + for (int i = 0; i < ptr_len; i++) + ptr[i] = n; + mapped = (ptr != NULL); + } + return mapped; + }; + return fn; + } + + auto set_refptr_func (int n) + { + auto fn = [=](void) -> bool + { + bool mapped; + #pragma omp target map(from:mapped) + { + if (refptr) + for (int i = 0; i < refptr_len; i++) + refptr[i] = n; + mapped = (refptr != NULL); + } + return mapped; + }; + return fn; + } +}; + +int main (void) +{ + #define N 10 + int *ptr1 = new int[N]; + int *ptr2 = new int[N]; + + memset (ptr1, 0, sizeof (int) * N); + memset (ptr2, 0, sizeof (int) * N); + + T a = { ptr1, N, ptr2, N }; + + auto p1 = a.set_ptr_func (1); + auto r2 = a.set_refptr_func (2); + + if (p1 ()) + abort (); + if (r2 ()) + abort (); + + if (a.ptr != ptr1) + abort (); + if (a.refptr != ptr2) + abort (); + + for (int i = 0; i < N; i++) + if (ptr1[i] != 0) + abort (); + + for (int i = 0; i < N; i++) + if (ptr2[i] != 0) + abort (); + + #pragma omp target data map(ptr1[:N], ptr2[:N]) + { + if (!p1 ()) + abort (); + if (!r2 ()) + abort (); + } + + if (a.ptr != ptr1) + abort (); + if (a.refptr != ptr2) + abort (); + + for (int i = 0; i < N; i++) + if (ptr1[i] != 1) + abort (); + + for (int i = 0; i < N; i++) + if (ptr2[i] != 2) + abort (); + + return 0; +} + +/* { dg-final { scan-tree-dump {#pragma omp target num_teams.* map\(to:\*__closure \[len: [0-9]+\]\) map\(firstprivate:__closure \[pointer assign, bias: 0\]\) map\(tofrom:\*_[0-9]+ \[len: [0-9]+\]\) map\(always_pointer:__closure->__this \[pointer assign, bias: 0\]\) map\(attach_zero_length_array_section:_[0-9]+->ptr \[bias: 0\]\) map\(from:mapped \[len: 1\]\) map\(alloc:MEM.* \[len: 0\]\) firstprivate\(n\) map\(alloc:MEM.* \[len: 0\]\) map\(firstprivate:this \[pointer assign, bias: 0\]\)} "gimple" } } */ + +/* { dg-final { scan-tree-dump {#pragma omp target num_teams.* map\(to:\*__closure \[len: [0-9]+\]\) map\(firstprivate:__closure \[pointer assign, bias: 0\]\) map\(tofrom:\*_[0-9]+ \[len: [0-9]+\]\) map\(always_pointer:__closure->__this \[pointer assign, bias: 0\]\) map\(alloc:\*_[0-9]+->refptr \[pointer assign, zero-length array section, bias: 0\]\) map\(attach:_3->refptr \[bias: 0\]\) map\(from:mapped \[len: [0-9]+\]\) map\(alloc:\*_[0-9]+ \[len: 0\]\) firstprivate\(n\) map\(alloc:MEM.* \[len: 0\]\) map\(firstprivate:this \[pointer assign, bias: 0\]\)} "gimple" } } */ diff --git a/gcc/tree-pretty-print.c b/gcc/tree-pretty-print.c index 3fae181fefbb..c8ccb43b072b 100644 --- a/gcc/tree-pretty-print.c +++ b/gcc/tree-pretty-print.c @@ -785,6 +785,7 @@ dump_omp_clause (pretty_printer *pp, tree clause, int spc, dump_flags_t flags) { case GOMP_MAP_ALLOC: case GOMP_MAP_POINTER: + case GOMP_MAP_POINTER_TO_ZERO_LENGTH_ARRAY_SECTION: pp_string (pp, "alloc"); break; case GOMP_MAP_IF_PRESENT: @@ -896,6 +897,9 @@ dump_omp_clause (pretty_printer *pp, tree clause, int spc, dump_flags_t flags) case GOMP_MAP_NONCONTIG_ARRAY_FORCE_PRESENT: pp_string (pp, "force_present,noncontig_array"); break; + case GOMP_MAP_ATTACH_ZERO_LENGTH_ARRAY_SECTION: + pp_string (pp, "attach_zero_length_array_section"); + break; default: gcc_unreachable (); } @@ -921,6 +925,9 @@ dump_omp_clause (pretty_printer *pp, tree clause, int spc, dump_flags_t flags) case GOMP_MAP_ALWAYS_POINTER: pp_string (pp, " [pointer assign, bias: "); break; + case GOMP_MAP_POINTER_TO_ZERO_LENGTH_ARRAY_SECTION: + pp_string (pp, " [pointer assign, zero-length array section, bias: "); + break; case GOMP_MAP_TO_PSET: pp_string (pp, " [pointer set, len: "); break; @@ -928,6 +935,7 @@ dump_omp_clause (pretty_printer *pp, tree clause, int spc, dump_flags_t flags) case GOMP_MAP_DETACH: case GOMP_MAP_FORCE_DETACH: case GOMP_MAP_ATTACH_DETACH: + case GOMP_MAP_ATTACH_ZERO_LENGTH_ARRAY_SECTION: pp_string (pp, " [bias: "); break; default: diff --git a/include/gomp-constants.h b/include/gomp-constants.h index 45d553c47c20..d3d4514c84b9 100644 --- a/include/gomp-constants.h +++ b/include/gomp-constants.h @@ -133,6 +133,11 @@ enum gomp_map_kind No refcount is bumped by this, and the store is done unconditionally. */ GOMP_MAP_ALWAYS_POINTER = (GOMP_MAP_FLAG_SPECIAL_2 | GOMP_MAP_FLAG_SPECIAL | 1), + /* Like GOMP_MAP_POINTER, but allow zero-length array section, i.e. set to + NULL if target is not mapped. */ + GOMP_MAP_POINTER_TO_ZERO_LENGTH_ARRAY_SECTION + = (GOMP_MAP_FLAG_SPECIAL_2 + | GOMP_MAP_FLAG_SPECIAL | 2), /* Forced deallocation of zero length array section. */ GOMP_MAP_DELETE_ZERO_LEN_ARRAY_SECTION = (GOMP_MAP_FLAG_SPECIAL_2 @@ -178,6 +183,12 @@ enum gomp_map_kind GOMP_MAP_NONCONTIG_ARRAY_FORCE_PRESENT = (GOMP_MAP_NONCONTIG_ARRAY | GOMP_MAP_FORCE_PRESENT), + /* Like GOMP_MAP_ATTACH, but allow attaching to zero-length array sections + (i.e. set to NULL when array section is not mapped) Currently only used + by OpenMP. */ + GOMP_MAP_ATTACH_ZERO_LENGTH_ARRAY_SECTION + = (GOMP_MAP_DEEP_COPY | 2), + /* Internal to GCC, not used in libgomp. */ /* Do not map, but pointer assign a pointer instead. */ GOMP_MAP_FIRSTPRIVATE_POINTER = (GOMP_MAP_LAST | 1), @@ -201,7 +212,8 @@ enum gomp_map_kind ((X) == GOMP_MAP_ALWAYS_POINTER) #define GOMP_MAP_POINTER_P(X) \ - ((X) == GOMP_MAP_POINTER) + ((X) == GOMP_MAP_POINTER \ + || (X) == GOMP_MAP_POINTER_TO_ZERO_LENGTH_ARRAY_SECTION) #define GOMP_MAP_ALWAYS_TO_P(X) \ (((X) == GOMP_MAP_ALWAYS_TO) || ((X) == GOMP_MAP_ALWAYS_TOFROM)) diff --git a/libgomp/libgomp.h b/libgomp/libgomp.h index 6377d2917650..096029fb8b07 100644 --- a/libgomp/libgomp.h +++ b/libgomp/libgomp.h @@ -1228,7 +1228,7 @@ extern uintptr_t gomp_map_val (struct target_mem_desc *, void **, size_t); extern void gomp_attach_pointer (struct gomp_device_descr *, struct goacc_asyncqueue *, splay_tree, splay_tree_key, uintptr_t, size_t, - struct gomp_coalesce_buf *); + struct gomp_coalesce_buf *, bool); extern void gomp_detach_pointer (struct gomp_device_descr *, struct goacc_asyncqueue *, splay_tree_key, uintptr_t, bool, struct gomp_coalesce_buf *); diff --git a/libgomp/oacc-mem.c b/libgomp/oacc-mem.c index 82da2196c94a..6bad8a5ffd87 100644 --- a/libgomp/oacc-mem.c +++ b/libgomp/oacc-mem.c @@ -966,7 +966,7 @@ acc_attach_async (void **hostaddr, int async) } gomp_attach_pointer (acc_dev, aq, &acc_dev->mem_map, n, (uintptr_t) hostaddr, - 0, NULL); + 0, NULL, false); gomp_mutex_unlock (&acc_dev->lock); } @@ -1199,7 +1199,7 @@ goacc_enter_data_internal (struct gomp_device_descr *acc_dev, size_t mapnum, if ((kinds[i] & 0xff) == GOMP_MAP_ATTACH) { gomp_attach_pointer (acc_dev, aq, &acc_dev->mem_map, n, - (uintptr_t) h, s, NULL); + (uintptr_t) h, s, NULL, false); /* OpenACC 'attach'/'detach' doesn't affect structured/dynamic reference counts ('n->refcount', 'n->dynamic_refcount'). */ } @@ -1217,7 +1217,8 @@ goacc_enter_data_internal (struct gomp_device_descr *acc_dev, size_t mapnum, splay_tree_key m = lookup_host (acc_dev, hostaddrs[j], sizeof (void *)); gomp_attach_pointer (acc_dev, aq, &acc_dev->mem_map, m, - (uintptr_t) hostaddrs[j], sizes[j], NULL); + (uintptr_t) hostaddrs[j], sizes[j], NULL, + false); } bool processed = false; diff --git a/libgomp/target.c b/libgomp/target.c index a51fba65cc46..0c9fe58e14d5 100644 --- a/libgomp/target.c +++ b/libgomp/target.c @@ -501,7 +501,8 @@ gomp_map_vars_existing (struct gomp_device_descr *devicep, struct gomp_coalesce_buf *cbuf, htab_t *refcount_set) { - assert (kind != GOMP_MAP_ATTACH); + assert (kind != GOMP_MAP_ATTACH + || kind != GOMP_MAP_ATTACH_ZERO_LENGTH_ARRAY_SECTION); tgt_var->key = oldn; tgt_var->copy_from = GOMP_MAP_COPY_FROM_P (kind); @@ -541,7 +542,8 @@ get_kind (bool short_mapkind, void *kinds, int idx) static void gomp_map_pointer (struct target_mem_desc *tgt, struct goacc_asyncqueue *aq, uintptr_t host_ptr, uintptr_t target_offset, uintptr_t bias, - struct gomp_coalesce_buf *cbuf) + struct gomp_coalesce_buf *cbuf, + bool allow_zero_length_array_sections) { struct gomp_device_descr *devicep = tgt->device_descr; struct splay_tree_s *mem_map = &devicep->mem_map; @@ -563,16 +565,24 @@ gomp_map_pointer (struct target_mem_desc *tgt, struct goacc_asyncqueue *aq, splay_tree_key n = gomp_map_lookup (mem_map, &cur_node); if (n == NULL) { - gomp_mutex_unlock (&devicep->lock); - gomp_fatal ("Pointer target of array section wasn't mapped"); - } - cur_node.host_start -= n->host_start; - cur_node.tgt_offset - = n->tgt->tgt_start + n->tgt_offset + cur_node.host_start; - /* At this point tgt_offset is target address of the - array section. Now subtract bias to get what we want - to initialize the pointer with. */ - cur_node.tgt_offset -= bias; + if (allow_zero_length_array_sections) + cur_node.tgt_offset = 0; + else + { + gomp_mutex_unlock (&devicep->lock); + gomp_fatal ("Pointer target of array section wasn't mapped"); + } + } + else + { + cur_node.host_start -= n->host_start; + cur_node.tgt_offset + = n->tgt->tgt_start + n->tgt_offset + cur_node.host_start; + /* At this point tgt_offset is target address of the + array section. Now subtract bias to get what we want + to initialize the pointer with. */ + cur_node.tgt_offset -= bias; + } gomp_copy_host2dev (devicep, aq, (void *) (tgt->tgt_start + target_offset), (void *) &cur_node.tgt_offset, sizeof (void *), true, cbuf); @@ -644,7 +654,8 @@ attribute_hidden void gomp_attach_pointer (struct gomp_device_descr *devicep, struct goacc_asyncqueue *aq, splay_tree mem_map, splay_tree_key n, uintptr_t attach_to, size_t bias, - struct gomp_coalesce_buf *cbufp) + struct gomp_coalesce_buf *cbufp, + bool allow_zero_length_array_sections) { struct splay_tree_key_s s; size_t size, idx; @@ -696,11 +707,21 @@ gomp_attach_pointer (struct gomp_device_descr *devicep, if (!tn) { - gomp_mutex_unlock (&devicep->lock); - gomp_fatal ("pointer target not mapped for attach"); + if (allow_zero_length_array_sections) + { + /* When allowing attachment to zero-length array sections, we + allow attaching to NULL pointers when the target region is not + mapped. */ + data = 0; + } + else + { + gomp_mutex_unlock (&devicep->lock); + gomp_fatal ("pointer target not mapped for attach"); + } } - - data = tn->tgt->tgt_start + tn->tgt_offset + target - tn->host_start; + else + data = tn->tgt->tgt_start + tn->tgt_offset + target - tn->host_start; gomp_debug (1, "%s: attaching host %p, target %p (struct base %p) to %p\n", @@ -960,7 +981,9 @@ gomp_map_vars_internal (struct gomp_device_descr *devicep, has_firstprivate = true; continue; } - else if ((kind & typemask) == GOMP_MAP_ATTACH) + else if ((kind & typemask) == GOMP_MAP_ATTACH + || ((kind & typemask) + == GOMP_MAP_ATTACH_ZERO_LENGTH_ARRAY_SECTION)) { tgt->list[i].key = NULL; has_firstprivate = true; @@ -1268,7 +1291,7 @@ gomp_map_vars_internal (struct gomp_device_descr *devicep, (uintptr_t) *(void **) hostaddrs[j], k->tgt_offset + ((uintptr_t) hostaddrs[j] - k->host_start), - sizes[j], cbufp); + sizes[j], cbufp, false); } } i = j - 1; @@ -1395,6 +1418,7 @@ gomp_map_vars_internal (struct gomp_device_descr *devicep, ++i; continue; case GOMP_MAP_ATTACH: + case GOMP_MAP_ATTACH_ZERO_LENGTH_ARRAY_SECTION: { cur_node.host_start = (uintptr_t) hostaddrs[i]; cur_node.host_end = cur_node.host_start + sizeof (void *); @@ -1411,9 +1435,12 @@ gomp_map_vars_internal (struct gomp_device_descr *devicep, structured/dynamic reference counts ('n->refcount', 'n->dynamic_refcount'). */ + bool zlas + = ((kind & typemask) + == GOMP_MAP_ATTACH_ZERO_LENGTH_ARRAY_SECTION); gomp_attach_pointer (devicep, aq, mem_map, n, (uintptr_t) hostaddrs[i], sizes[i], - cbufp); + cbufp, zlas); } else if ((pragma_kind & GOMP_MAP_VARS_OPENACC) != 0) { @@ -1529,9 +1556,12 @@ gomp_map_vars_internal (struct gomp_device_descr *devicep, cbufp); break; case GOMP_MAP_POINTER: - gomp_map_pointer (tgt, aq, - (uintptr_t) *(void **) k->host_start, - k->tgt_offset, sizes[i], cbufp); + case GOMP_MAP_POINTER_TO_ZERO_LENGTH_ARRAY_SECTION: + gomp_map_pointer + (tgt, aq, (uintptr_t) *(void **) k->host_start, + k->tgt_offset, sizes[i], cbufp, + ((kind & typemask) + == GOMP_MAP_POINTER_TO_ZERO_LENGTH_ARRAY_SECTION)); break; case GOMP_MAP_TO_PSET: gomp_copy_host2dev (devicep, aq, @@ -1573,7 +1603,7 @@ gomp_map_vars_internal (struct gomp_device_descr *devicep, k->tgt_offset + ((uintptr_t) hostaddrs[j] - k->host_start), - sizes[j], cbufp); + sizes[j], cbufp, false); } } i = j - 1; diff --git a/libgomp/testsuite/libgomp.c++/target-this-1.C b/libgomp/testsuite/libgomp.c++/target-this-1.C new file mode 100644 index 000000000000..a591ea4c5643 --- /dev/null +++ b/libgomp/testsuite/libgomp.c++/target-this-1.C @@ -0,0 +1,29 @@ +extern "C" void abort (); + +struct S +{ + int a, b, c, d; + + int sum (void) + { + int val = 0; + val += a + b + this->c + this->d; + return val; + } + + int sum_offload (void) + { + int val = 0; + #pragma omp target map(val) + val += a + b + this->c + this->d; + return val; + } +}; + +int main (void) +{ + S s = { 1, 2, 3, 4 }; + if (s.sum () != s.sum_offload ()) + abort (); + return 0; +} diff --git a/libgomp/testsuite/libgomp.c++/target-this-2.C b/libgomp/testsuite/libgomp.c++/target-this-2.C new file mode 100644 index 000000000000..8119be8c2c5d --- /dev/null +++ b/libgomp/testsuite/libgomp.c++/target-this-2.C @@ -0,0 +1,47 @@ + +// We use 'auto' without a function return type, so specify dialect here +// { dg-additional-options "-std=c++14" } + +extern "C" void abort (); + +struct T +{ + int x, y; + + auto sum_func (int n) + { + auto fn = [=](int m) -> int + { + int v; + v = (x + y) * n + m; + return v; + }; + return fn; + } + + auto sum_func_offload (int n) + { + auto fn = [=](int m) -> int + { + int v; + #pragma omp target map(from:v) + v = (x + y) * n + m; + return v; + }; + return fn; + } + +}; + +int main (void) +{ + T a = { 1, 2 }; + + auto s1 = a.sum_func (3); + auto s2 = a.sum_func_offload (3); + + if (s1 (1) != s2 (1)) + abort (); + + return 0; +} diff --git a/libgomp/testsuite/libgomp.c++/target-this-3.C b/libgomp/testsuite/libgomp.c++/target-this-3.C new file mode 100644 index 000000000000..e15f69a16233 --- /dev/null +++ b/libgomp/testsuite/libgomp.c++/target-this-3.C @@ -0,0 +1,99 @@ +#include +#include +extern "C" void abort (); + +struct S +{ + int * ptr; + int ptr_len; + + int *&refptr; + int refptr_len; + + bool set_ptr (int n) + { + bool mapped; + #pragma omp target map(from:mapped) + { + if (ptr != NULL) + for (int i = 0; i < ptr_len; i++) + ptr[i] = n; + mapped = (ptr != NULL); + } + return mapped; + } + + bool set_refptr (int n) + { + bool mapped; + #pragma omp target map(from:mapped) + { + if (refptr != NULL) + for (int i = 0; i < refptr_len; i++) + refptr[i] = n; + mapped = (refptr != NULL); + } + return mapped; + } +}; + +int main (void) +{ + #define N 10 + int *ptr1 = new int[N]; + int *ptr2 = new int[N]; + + memset (ptr1, 0, sizeof (int) * N); + memset (ptr2, 0, sizeof (int) * N); + + S s = { ptr1, N, ptr2, N }; + + bool mapped; + int val = 123; + + mapped = s.set_ptr (val); + if (mapped) + abort (); + if (s.ptr != ptr1) + abort (); + for (int i = 0; i < N; i++) + if (ptr1[i] != 0) + abort (); + + mapped = s.set_refptr (val); + if (mapped) + abort (); + if (s.refptr != ptr2) + abort (); + for (int i = 0; i < N; i++) + if (ptr2[i] != 0) + abort (); + + #pragma omp target data map(ptr1[:N]) + mapped = s.set_ptr (val); + + if (!mapped) + abort (); + if (s.set_refptr (0)) + abort (); + if (s.ptr != ptr1 || s.refptr != ptr2) + abort (); + for (int i = 0; i < N; i++) + if (ptr1[i] != val) + abort (); + + #pragma omp target data map(ptr2[:N]) + mapped = s.set_refptr (val); + + if (!mapped) + abort (); + if (s.set_ptr (0)) + abort (); + if (s.ptr != ptr1 || s.refptr != ptr2) + abort (); + for (int i = 0; i < N; i++) + if (ptr2[i] != val) + abort (); + + return 0; +} diff --git a/libgomp/testsuite/libgomp.c++/target-this-4.C b/libgomp/testsuite/libgomp.c++/target-this-4.C new file mode 100644 index 000000000000..9f53677a2400 --- /dev/null +++ b/libgomp/testsuite/libgomp.c++/target-this-4.C @@ -0,0 +1,104 @@ + +// We use 'auto' without a function return type, so specify dialect here +// { dg-additional-options "-std=c++14" } +#include +#include + +struct T +{ + int *ptr; + int ptr_len; + + int *&refptr; + int refptr_len; + + auto set_ptr_func (int n) + { + auto fn = [=](void) -> bool + { + bool mapped; + #pragma omp target map(from:mapped) + { + if (ptr) + for (int i = 0; i < ptr_len; i++) + ptr[i] = n; + mapped = (ptr != NULL); + } + return mapped; + }; + return fn; + } + + auto set_refptr_func (int n) + { + auto fn = [=](void) -> bool + { + bool mapped; + #pragma omp target map(from:mapped) + { + if (refptr) + for (int i = 0; i < refptr_len; i++) + refptr[i] = n; + mapped = (refptr != NULL); + } + return mapped; + }; + return fn; + } +}; + +int main (void) +{ + #define N 10 + int *ptr1 = new int[N]; + int *ptr2 = new int[N]; + + memset (ptr1, 0, sizeof (int) * N); + memset (ptr2, 0, sizeof (int) * N); + + T a = { ptr1, N, ptr2, N }; + + auto p1 = a.set_ptr_func (1); + auto r2 = a.set_refptr_func (2); + + if (p1 ()) + abort (); + if (r2 ()) + abort (); + + if (a.ptr != ptr1) + abort (); + if (a.refptr != ptr2) + abort (); + + for (int i = 0; i < N; i++) + if (ptr1[i] != 0) + abort (); + + for (int i = 0; i < N; i++) + if (ptr2[i] != 0) + abort (); + + #pragma omp target data map(ptr1[:N], ptr2[:N]) + { + if (!p1 ()) + abort (); + if (!r2 ()) + abort (); + } + + if (a.ptr != ptr1) + abort (); + if (a.refptr != ptr2) + abort (); + + for (int i = 0; i < N; i++) + if (ptr1[i] != 1) + abort (); + + for (int i = 0; i < N; i++) + if (ptr2[i] != 2) + abort (); + + return 0; +}