From mboxrd@z Thu Jan 1 00:00:00 1970 Return-Path: Received: by sourceware.org (Postfix, from userid 1709) id 3F0383898516; Mon, 31 May 2021 09:02:53 +0000 (GMT) DKIM-Filter: OpenDKIM Filter v2.11.0 sourceware.org 3F0383898516 Content-Type: text/plain; charset="us-ascii" MIME-Version: 1.0 Content-Transfer-Encoding: 7bit From: Chung-Lin Tang To: gcc-cvs@gcc.gnu.org Subject: [gcc/devel/omp/gcc-11] Revert "OpenMP 5.0: map this[:1] in C++ non-static member functions (PR 92120)" X-Act-Checkin: gcc X-Git-Author: Chung-Lin Tang X-Git-Refname: refs/heads/devel/omp/gcc-11 X-Git-Oldrev: de5cd4140c24292898bffd6145594f0b4170f28a X-Git-Newrev: a81101b9f200ff8dd614deea39f50cb03991e462 Message-Id: <20210531090253.3F0383898516@sourceware.org> Date: Mon, 31 May 2021 09:02:53 +0000 (GMT) X-BeenThere: gcc-cvs@gcc.gnu.org X-Mailman-Version: 2.1.29 Precedence: list List-Id: Gcc-cvs mailing list List-Unsubscribe: , List-Archive: List-Help: List-Subscribe: , X-List-Received-Date: Mon, 31 May 2021 09:02:53 -0000 https://gcc.gnu.org/g:a81101b9f200ff8dd614deea39f50cb03991e462 commit a81101b9f200ff8dd614deea39f50cb03991e462 Author: Chung-Lin Tang Date: Wed May 26 19:12:05 2021 +0800 Revert "OpenMP 5.0: map this[:1] in C++ non-static member functions (PR 92120)" This reverts commit d511585fcfa94fdd8a13a82c027af403749eb4d1. Diff: --- gcc/cp/cp-tree.h | 2 - gcc/cp/lambda.c | 3 - gcc/cp/parser.c | 62 +++++- gcc/cp/semantics.c | 299 +------------------------- gcc/omp-low.c | 2 - gcc/testsuite/g++.dg/gomp/target-this-1.C | 33 --- gcc/testsuite/g++.dg/gomp/target-this-2.C | 49 ----- gcc/testsuite/g++.dg/gomp/target-this-3.C | 105 --------- gcc/testsuite/g++.dg/gomp/target-this-4.C | 107 --------- gcc/tree-pretty-print.c | 8 - include/gomp-constants.h | 14 +- libgomp/libgomp.h | 2 +- libgomp/oacc-mem.c | 7 +- libgomp/target.c | 78 +++---- libgomp/testsuite/libgomp.c++/target-this-1.C | 29 --- libgomp/testsuite/libgomp.c++/target-this-2.C | 47 ---- libgomp/testsuite/libgomp.c++/target-this-3.C | 99 --------- libgomp/testsuite/libgomp.c++/target-this-4.C | 104 --------- 18 files changed, 87 insertions(+), 963 deletions(-) diff --git a/gcc/cp/cp-tree.h b/gcc/cp/cp-tree.h index 4ee5e0fd324..17ba7cbd222 100644 --- a/gcc/cp/cp-tree.h +++ b/gcc/cp/cp-tree.h @@ -7566,8 +7566,6 @@ extern tree start_lambda_function (tree fn, tree lambda_expr); extern void finish_lambda_function (tree body); extern bool regenerated_lambda_fn_p (tree); extern tree most_general_lambda (tree); -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 23c1682893c..16e2b4c18b4 100644 --- a/gcc/cp/lambda.c +++ b/gcc/cp/lambda.c @@ -845,9 +845,6 @@ 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 0927f29bfe7..901402bccb1 100644 --- a/gcc/cp/parser.c +++ b/gcc/cp/parser.c @@ -42092,6 +42092,8 @@ 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); @@ -42144,7 +42146,6 @@ 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: @@ -42196,9 +42197,15 @@ cp_parser_omp_target (cp_parser *parser, cp_token *pragma_tok, cclauses[C_OMP_CLAUSE_SPLIT_TARGET] = tc; } } - finish_omp_target (pragma_tok->location, - cclauses[C_OMP_CLAUSE_SPLIT_TARGET], body, true); - return true; + 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; } else if (!flag_openmp) /* flag_openmp_simd */ { @@ -42235,14 +42242,49 @@ cp_parser_omp_target (cp_parser *parser, cp_token *pragma_tok, return false; } - tree clauses = cp_parser_omp_all_clauses (parser, OMP_TARGET_CLAUSE_MASK, - "#pragma omp target", pragma_tok); - c_omp_adjust_map_clauses (clauses, true); + 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); keep_next_level (true); - set_omp_target_this_expr (NULL_TREE); - tree body = cp_parser_omp_structured_block (parser, if_p); + OMP_TARGET_BODY (stmt) = cp_parser_omp_structured_block (parser, if_p); - finish_omp_target (pragma_tok->location, clauses, body, false); + 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); + } return true; } diff --git a/gcc/cp/semantics.c b/gcc/cp/semantics.c index cc049f5e8de..80e486a090d 100644 --- a/gcc/cp/semantics.c +++ b/gcc/cp/semantics.c @@ -61,10 +61,6 @@ 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 --------------------------------- @@ -2080,7 +2076,6 @@ 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; @@ -2119,14 +2114,6 @@ 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) @@ -2187,14 +2174,6 @@ 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; } @@ -2931,15 +2910,8 @@ finish_this_expr (void) } if (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; - } + /* The keyword 'this' is a prvalue expression. */ + return rvalue (result); tree fn = current_nonlambda_function (); if (fn && DECL_STATIC_FUNCTION_P (fn)) @@ -4961,12 +4933,12 @@ handle_omp_array_sections_1 (tree c, tree t, vec &types, { if (error_operand_p (t)) return error_mark_node; - 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); + 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); ret = t; if (TREE_CODE (t) == COMPONENT_REF && (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP @@ -5604,8 +5576,6 @@ 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) @@ -6597,7 +6567,6 @@ 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 @@ -7836,15 +7805,10 @@ 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); - indir_component_ref_p = true; - STRIP_NOPS (t); - } + t = TREE_OPERAND (TREE_OPERAND (t, 0), 0); if (TREE_CODE (t) == COMPONENT_REF && ((ort & C_ORT_OMP_DECLARE_SIMD) == C_ORT_OMP || ort == C_ORT_ACC) @@ -7942,7 +7906,6 @@ 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 @@ -9062,256 +9025,6 @@ 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 a16603ab441..3f7403dd206 100644 --- a/gcc/omp-low.c +++ b/gcc/omp-low.c @@ -12565,8 +12565,6 @@ 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 deleted file mode 100644 index de93a3e5e57..00000000000 --- a/gcc/testsuite/g++.dg/gomp/target-this-1.C +++ /dev/null @@ -1,33 +0,0 @@ -// { 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 deleted file mode 100644 index a5e832130fb..00000000000 --- a/gcc/testsuite/g++.dg/gomp/target-this-2.C +++ /dev/null @@ -1,49 +0,0 @@ -// 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 deleted file mode 100644 index 208ea079b95..00000000000 --- a/gcc/testsuite/g++.dg/gomp/target-this-3.C +++ /dev/null @@ -1,105 +0,0 @@ -// { 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 deleted file mode 100644 index f42cf384541..00000000000 --- a/gcc/testsuite/g++.dg/gomp/target-this-4.C +++ /dev/null @@ -1,107 +0,0 @@ -// 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 7d386fca8ee..9c21179c9d5 100644 --- a/gcc/tree-pretty-print.c +++ b/gcc/tree-pretty-print.c @@ -835,7 +835,6 @@ 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: @@ -947,9 +946,6 @@ 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 (); } @@ -975,9 +971,6 @@ 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; @@ -985,7 +978,6 @@ 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 d9b80ac9732..43e4949da3d 100644 --- a/include/gomp-constants.h +++ b/include/gomp-constants.h @@ -133,11 +133,6 @@ 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 @@ -183,12 +178,6 @@ 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), @@ -212,8 +201,7 @@ enum gomp_map_kind ((X) == GOMP_MAP_ALWAYS_POINTER) #define GOMP_MAP_POINTER_P(X) \ - ((X) == GOMP_MAP_POINTER \ - || (X) == GOMP_MAP_POINTER_TO_ZERO_LENGTH_ARRAY_SECTION) + ((X) == GOMP_MAP_POINTER) #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 0c8085e350c..62b5ae9de8f 100644 --- a/libgomp/libgomp.h +++ b/libgomp/libgomp.h @@ -1233,7 +1233,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 *, bool); + struct gomp_coalesce_buf *); 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 fe403f3ecc1..9498ede3132 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, false); + 0, NULL); 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, false); + (uintptr_t) h, s, NULL); /* OpenACC 'attach'/'detach' doesn't affect structured/dynamic reference counts ('n->refcount', 'n->dynamic_refcount'). */ } @@ -1217,8 +1217,7 @@ 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, - false); + (uintptr_t) hostaddrs[j], sizes[j], NULL); } bool processed = false; diff --git a/libgomp/target.c b/libgomp/target.c index 2633ab1e43f..332b0e2cc8f 100644 --- a/libgomp/target.c +++ b/libgomp/target.c @@ -514,8 +514,7 @@ gomp_map_vars_existing (struct gomp_device_descr *devicep, struct gomp_coalesce_buf *cbuf, htab_t *refcount_set) { - assert (kind != GOMP_MAP_ATTACH - || kind != GOMP_MAP_ATTACH_ZERO_LENGTH_ARRAY_SECTION); + assert (kind != GOMP_MAP_ATTACH); tgt_var->key = oldn; tgt_var->copy_from = GOMP_MAP_COPY_FROM_P (kind); @@ -555,8 +554,7 @@ 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, - bool allow_zero_length_array_sections) + struct gomp_coalesce_buf *cbuf) { struct gomp_device_descr *devicep = tgt->device_descr; struct splay_tree_s *mem_map = &devicep->mem_map; @@ -578,24 +576,16 @@ 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) { - 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_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; gomp_copy_host2dev (devicep, aq, (void *) (tgt->tgt_start + target_offset), (void *) &cur_node.tgt_offset, sizeof (void *), true, cbuf); @@ -667,8 +657,7 @@ 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, - bool allow_zero_length_array_sections) + struct gomp_coalesce_buf *cbufp) { struct splay_tree_key_s s; size_t size, idx; @@ -720,21 +709,11 @@ gomp_attach_pointer (struct gomp_device_descr *devicep, if (!tn) { - 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"); - } + gomp_mutex_unlock (&devicep->lock); + gomp_fatal ("pointer target not mapped for attach"); } - else - data = tn->tgt->tgt_start + tn->tgt_offset + target - tn->host_start; + + 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", @@ -994,9 +973,7 @@ gomp_map_vars_internal (struct gomp_device_descr *devicep, has_firstprivate = true; continue; } - else if ((kind & typemask) == GOMP_MAP_ATTACH - || ((kind & typemask) - == GOMP_MAP_ATTACH_ZERO_LENGTH_ARRAY_SECTION)) + else if ((kind & typemask) == GOMP_MAP_ATTACH) { tgt->list[i].key = NULL; has_firstprivate = true; @@ -1304,7 +1281,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, false); + sizes[j], cbufp); } } i = j - 1; @@ -1431,7 +1408,6 @@ 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 *); @@ -1448,12 +1424,9 @@ 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, zlas); + cbufp); } else if ((pragma_kind & GOMP_MAP_VARS_OPENACC) != 0) { @@ -1569,12 +1542,9 @@ gomp_map_vars_internal (struct gomp_device_descr *devicep, cbufp); break; case GOMP_MAP_POINTER: - 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)); + gomp_map_pointer (tgt, aq, + (uintptr_t) *(void **) k->host_start, + k->tgt_offset, sizes[i], cbufp); break; case GOMP_MAP_TO_PSET: gomp_copy_host2dev (devicep, aq, @@ -1616,7 +1586,7 @@ gomp_map_vars_internal (struct gomp_device_descr *devicep, k->tgt_offset + ((uintptr_t) hostaddrs[j] - k->host_start), - sizes[j], cbufp, false); + sizes[j], cbufp); } } i = j - 1; diff --git a/libgomp/testsuite/libgomp.c++/target-this-1.C b/libgomp/testsuite/libgomp.c++/target-this-1.C deleted file mode 100644 index a591ea4c564..00000000000 --- a/libgomp/testsuite/libgomp.c++/target-this-1.C +++ /dev/null @@ -1,29 +0,0 @@ -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 deleted file mode 100644 index 8119be8c2c5..00000000000 --- a/libgomp/testsuite/libgomp.c++/target-this-2.C +++ /dev/null @@ -1,47 +0,0 @@ - -// 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 deleted file mode 100644 index e15f69a1623..00000000000 --- a/libgomp/testsuite/libgomp.c++/target-this-3.C +++ /dev/null @@ -1,99 +0,0 @@ -#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 deleted file mode 100644 index 9f53677a240..00000000000 --- a/libgomp/testsuite/libgomp.c++/target-this-4.C +++ /dev/null @@ -1,104 +0,0 @@ - -// 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; -}