From patchwork Fri Nov 19 13:54:12 2021 Content-Type: text/plain; charset="utf-8" MIME-Version: 1.0 Content-Transfer-Encoding: 8bit X-Patchwork-Submitter: Chung-Lin Tang X-Patchwork-Id: 47930 Return-Path: X-Original-To: patchwork@sourceware.org Delivered-To: patchwork@sourceware.org Received: from server2.sourceware.org (localhost [IPv6:::1]) by sourceware.org (Postfix) with ESMTP id 4C21B385C8B2 for ; Fri, 19 Nov 2021 13:55:02 +0000 (GMT) X-Original-To: gcc-patches@gcc.gnu.org Delivered-To: gcc-patches@gcc.gnu.org Received: from esa1.mentor.iphmx.com (esa1.mentor.iphmx.com [68.232.129.153]) by sourceware.org (Postfix) with ESMTPS id 1F03E385AC2B; Fri, 19 Nov 2021 13:54:24 +0000 (GMT) DMARC-Filter: OpenDMARC Filter v1.4.1 sourceware.org 1F03E385AC2B Authentication-Results: sourceware.org; dmarc=none (p=none dis=none) header.from=codesourcery.com Authentication-Results: sourceware.org; spf=pass smtp.mailfrom=mentor.com IronPort-SDR: RIhiN4UISEZ00H+vJJJO1eaHjAQTJESkRVDG2m1MEHSKo9uCERS/y12pqbBA30mHnhEke92MOe J/gJl0cXhwgwCTjRpytnTgJRKI380o6z/HFe3GA52K9/HELR4vRN5Lil7EkBa+zLt/Cq/3DhiT vVX2Q5RVtfVjAIKFzffQD98/WGi+YR7WSe0JfSajPyn6/T5LdLUgih5pq+zcv3sBpu3WXTqODC 26m7TeRWK3xEeHCWnIf980MImhk9IQ2rPuRu7pCQIDvAh3/xdVwPW+UawiAPvLudesSVbDgAN0 GDq2/JqZYc02Fzs2wV2RKoZF X-IronPort-AV: E=Sophos;i="5.87,247,1631606400"; d="scan'208";a="71185807" Received: from orw-gwy-01-in.mentorg.com ([192.94.38.165]) by esa1.mentor.iphmx.com with ESMTP; 19 Nov 2021 05:54:23 -0800 IronPort-SDR: dPYBYmOTu+YHHB1NGAjrprxz3UAQ23iZbbtt/Aopljpzc854wcttV+sCX7AJCY14Vp957MjTdc +b5mDsodZcIBkaDH3WTO/nREN46Tu+4Gxa+UeLKA+y83YMZslxL5X6tyeeOBYdfeKkHawlI2+c vb9iDHszAAAAeqUy2tcbR20wVhULp7mMHTdnogggma5XApDNY5GhrP2j6Hl3gpuTG/W5LoODeR hsiT1CkGT6gej/4BP/xDuf6z5Cv0uO6lTUUqU5hoh5RVGpJrOi/6HYcfKQDVSOhfYatiZM9YLp 7ZQ= Message-ID: Date: Fri, 19 Nov 2021 21:54:12 +0800 MIME-Version: 1.0 User-Agent: Mozilla/5.0 (Macintosh; Intel Mac OS X 10.13; rv:91.0) Gecko/20100101 Thunderbird/91.3.1 From: Chung-Lin Tang Subject: [PATCH, v2, OpenMP 5.0] Remove array section base-pointer mapping semantics, and other front-end adjustments (mainline trunk) To: Jakub Jelinek , gcc-patches , Tobias Burnus , Fortran List , Catherine Moore , Thomas Schwinge , Julian Brown References: <309eb444-6b9b-18b3-90a7-ffd54d0d0335@codesourcery.com> Content-Language: en-US In-Reply-To: <309eb444-6b9b-18b3-90a7-ffd54d0d0335@codesourcery.com> X-ClientProxiedBy: svr-orw-mbx-01.mgc.mentorg.com (147.34.90.201) To svr-orw-mbx-02.mgc.mentorg.com (147.34.90.202) X-Spam-Status: No, score=-8.9 required=5.0 tests=BAYES_00, BODY_8BITS, GIT_PATCH_0, HEADER_FROM_DIFFERENT_DOMAINS, KAM_DMARC_STATUS, KAM_MANYTO, KAM_SHORT, SPF_HELO_PASS, SPF_PASS, TXREP autolearn=ham autolearn_force=no version=3.4.4 X-Spam-Checker-Version: SpamAssassin 3.4.4 (2020-01-24) on server2.sourceware.org X-BeenThere: gcc-patches@gcc.gnu.org X-Mailman-Version: 2.1.29 Precedence: list List-Id: Gcc-patches mailing list List-Unsubscribe: , List-Archive: List-Post: List-Help: List-Subscribe: , Errors-To: gcc-patches-bounces+patchwork=sourceware.org@gcc.gnu.org Sender: "Gcc-patches" Hi Jakub, attached is a rebased version of this "OpenMP fixes/adjustments" patch. This version removes some of the (ort == C_ORT_OMP || ort == C_ORT_ACC) stuff that's not needed in handle_omp_array_sections_1 and [c_]finish_omp_clauses. Note that this is meant to be patched atop of the recent also posted C++ PR92120 v5 patch: https://gcc.gnu.org/pipermail/gcc-patches/2021-November/584602.html Again, tested without regressions (together with the PR92120 patch), awaiting review. Thanks, Chung-Lin (ChangeLog updated below) On 2021/5/25 9:36 PM, Chung-Lin Tang wrote: > > This patch largely implements three pieces of functionality: > > (1) Per discussion and clarification on the omp-lang mailing list, > standards conforming behavior for mapping array sections should *NOT* also map the base-pointer, > i.e for this code: > >     struct S { int *ptr; ... }; >     struct S s; >     #pragma omp target enter data map(to: s.ptr[:100]) > > Currently we generate after gimplify: > #pragma omp target enter data map(struct:s [len: 1]) map(alloc:s.ptr [len: 8]) \ >                                map(to:*_1 [len: 400]) map(attach:s.ptr [bias: 0]) > > which is deemed incorrect. After this patch, the gimplify results are now adjusted to: > #pragma omp target enter data map(to:*_1 [len: 400]) map(attach:s.ptr [bias: 0]) > (the attach operation is still generated, and if s.ptr is already mapped prior, attachment will happen) > > The correct way of achieving the base-pointer-also-mapped behavior would be to use: > #pragma omp target enter data map(to: s.ptr, s.ptr[:100]) > > This adjustment in behavior required a number of small adjustments here and there in gimplify, including > to accomodate map sequences for C++ references. > > There is also a small Fortran front-end patch involved (hence CCing Tobias and fortran@). > The new gimplify processing changed behavior in handling GOMP_MAP_ALWAYS_POINTER maps such that > the libgomp.fortran/struct-elem-map-1.f90 regressed. It appeared that the Fortran FE was generating > a GOMP_MAP_ALWAYS_POINTER for array types, which didn't seem quite correct, and the pre-patch behavior > was removing this map anyways. I have a small change in trans-openmp.c:gfc_trans_omp_array_section > to not generate the map in this case, and so far no bad test results. > > (2) The second part (though kind of related to the first above) are fixes in libgomp/target.c > to not overwrite attached pointers when handling device<->host copies, mainly for the "always" case. > This behavior is also noted in the 5.0 spec, but not yet properly coded before. > > (3) The third is a set of changes to the C/C++ front-ends to extend the allowed component access syntax > in map clauses. This is actually mainly an effort to allow SPEC HPC to compile, so despite in the long > term the entire map clause syntax parsing is probably going to be revamped, we're still adding this in > for now. These changes are enabled for both OpenACC and OpenMP. 2021-11-19 Chung-Lin Tang gcc/c/ChangeLog: * c-parser.c (struct omp_dim): New struct type for use inside c_parser_omp_variable_list. (c_parser_omp_variable_list): Allow multiple levels of array and component accesses in array section base-pointer expression. (c_parser_omp_clause_to): Set 'allow_deref' to true in call to c_parser_omp_var_list_parens. (c_parser_omp_clause_from): Likewise. * c-typeck.c (handle_omp_array_sections_1): Extend allowed range of base-pointer expressions involving INDIRECT/MEM/ARRAY_REF and POINTER_PLUS_EXPR. (c_finish_omp_clauses): Extend allowed ranged of expressions involving INDIRECT/MEM/ARRAY_REF and POINTER_PLUS_EXPR. gcc/cp/ChangeLog: * parser.c (struct omp_dim): New struct type for use inside cp_parser_omp_var_list_no_open. (cp_parser_omp_var_list_no_open): Allow multiple levels of array and component accesses in array section base-pointer expression. (cp_parser_omp_all_clauses): Set 'allow_deref' to true in call to cp_parser_omp_var_list for to/from clauses. * semantics.c (handle_omp_array_sections_1): Extend allowed range of base-pointer expressions involving INDIRECT/MEM/ARRAY_REF and POINTER_PLUS_EXPR. (handle_omp_array_sections): Adjust pointer map generation of references. (finish_omp_clauses): Extend allowed ranged of expressions involving INDIRECT/MEM/ARRAY_REF and POINTER_PLUS_EXPR. gcc/fortran/ChangeLog: * trans-openmp.c (gfc_trans_omp_array_section): Do not generate GOMP_MAP_ALWAYS_POINTER map for main array maps of ARRAY_TYPE type. gcc/ChangeLog: * gimplify.c (extract_base_bit_offset): Add 'tree *offsetp' parameter, accomodate case where 'offset' return of get_inner_reference is non-NULL. (is_or_contains_p): Further robustify conditions. (omp_target_reorder_clauses): In alloc/to/from sorting phase, also move following GOMP_MAP_ALWAYS_POINTER maps along. Add new sorting phase where we make sure pointers with an attach/detach map are ordered correctly. (gimplify_scan_omp_clauses): Add modifications to avoid creating GOMP_MAP_STRUCT and associated alloc map for attach/detach maps. gcc/testsuite/ChangeLog: * c-c++-common/goacc/deep-copy-arrayofstruct.c: Adjust testcase. * c-c++-common/gomp/target-enter-data-1.c: New testcase. * c-c++-common/gomp/target-implicit-map-2.c: New testcase. libgomp/ChangeLog: * target.c (gomp_map_vars_existing): Make sure attached pointer is not overwritten during cross-host/device copying. (gomp_update): Likewise. (gomp_exit_data): Likewise. * testsuite/libgomp.c++/target-11.C: Adjust testcase. * testsuite/libgomp.c++/target-12.C: Likewise. * testsuite/libgomp.c++/target-15.C: Likewise. * testsuite/libgomp.c++/target-16.C: Likewise. * testsuite/libgomp.c++/target-17.C: Likewise. * testsuite/libgomp.c++/target-21.C: Likewise. * testsuite/libgomp.c++/target-23.C: Likewise. * testsuite/libgomp.c/target-23.c: Likewise. * testsuite/libgomp.c/target-29.c: Likewise. * testsuite/libgomp.c-c++-common/target-implicit-map-2.c: New testcase. diff --git a/gcc/c/c-parser.c b/gcc/c/c-parser.c index acbf20dcb58..1f27f35fdb9 100644 --- a/gcc/c/c-parser.c +++ b/gcc/c/c-parser.c @@ -12970,6 +12970,15 @@ c_parser_oacc_wait_list (c_parser *parser, location_t clause_loc, tree list) The optional ALLOW_DEREF argument is true if list items can use the deref (->) operator. */ +struct omp_dim +{ + tree low_bound, length; + location_t loc; + bool no_colon; + omp_dim (tree lb, tree len, location_t lo, bool nc) + : low_bound (lb), length (len), loc (lo), no_colon (nc) {} +}; + static tree c_parser_omp_variable_list (c_parser *parser, location_t clause_loc, @@ -12982,6 +12991,7 @@ c_parser_omp_variable_list (c_parser *parser, while (1) { + auto_vec dims; bool array_section_p = false; if (kind == OMP_CLAUSE_DEPEND || kind == OMP_CLAUSE_AFFINITY) { @@ -13101,6 +13111,7 @@ c_parser_omp_variable_list (c_parser *parser, case OMP_CLAUSE_MAP: case OMP_CLAUSE_FROM: case OMP_CLAUSE_TO: + start_component_ref: while (c_parser_next_token_is (parser, CPP_DOT) || (allow_deref && c_parser_next_token_is (parser, CPP_DEREF))) @@ -13128,9 +13139,13 @@ c_parser_omp_variable_list (c_parser *parser, case OMP_CLAUSE_REDUCTION: case OMP_CLAUSE_IN_REDUCTION: case OMP_CLAUSE_TASK_REDUCTION: + array_section_p = false; + dims.truncate (0); while (c_parser_next_token_is (parser, CPP_OPEN_SQUARE)) { + location_t loc = UNKNOWN_LOCATION; tree low_bound = NULL_TREE, length = NULL_TREE; + bool no_colon = false; c_parser_consume_token (parser); if (!c_parser_next_token_is (parser, CPP_COLON)) @@ -13141,9 +13156,13 @@ c_parser_omp_variable_list (c_parser *parser, expr = convert_lvalue_to_rvalue (expr_loc, expr, false, true); low_bound = expr.value; + loc = expr_loc; } if (c_parser_next_token_is (parser, CPP_CLOSE_SQUARE)) - length = integer_one_node; + { + length = integer_one_node; + no_colon = true; + } else { /* Look for `:'. */ @@ -13172,8 +13191,35 @@ c_parser_omp_variable_list (c_parser *parser, break; } - t = tree_cons (low_bound, length, t); + dims.safe_push (omp_dim (low_bound, length, loc, no_colon)); + } + + if (t != error_mark_node) + { + if ((kind == OMP_CLAUSE_MAP + || kind == OMP_CLAUSE_FROM + || kind == OMP_CLAUSE_TO) + && !array_section_p + && (c_parser_next_token_is (parser, CPP_DOT) + || (allow_deref + && c_parser_next_token_is (parser, + CPP_DEREF)))) + { + for (unsigned i = 0; i < dims.length (); i++) + { + gcc_assert (dims[i].length == integer_one_node); + t = build_array_ref (dims[i].loc, + t, dims[i].low_bound); + } + goto start_component_ref; + } + else + { + for (unsigned i = 0; i < dims.length (); i++) + t = tree_cons (dims[i].low_bound, dims[i].length, t); + } } + if ((kind == OMP_CLAUSE_DEPEND || kind == OMP_CLAUSE_AFFINITY) && t != error_mark_node && parser->tokens_avail != 2) @@ -16420,7 +16466,7 @@ c_parser_omp_clause_device_type (c_parser *parser, tree list) static tree c_parser_omp_clause_to (c_parser *parser, tree list) { - return c_parser_omp_var_list_parens (parser, OMP_CLAUSE_TO, list); + return c_parser_omp_var_list_parens (parser, OMP_CLAUSE_TO, list, true); } /* OpenMP 4.0: @@ -16429,7 +16475,7 @@ c_parser_omp_clause_to (c_parser *parser, tree list) static tree c_parser_omp_clause_from (c_parser *parser, tree list) { - return c_parser_omp_var_list_parens (parser, OMP_CLAUSE_FROM, list); + return c_parser_omp_var_list_parens (parser, OMP_CLAUSE_FROM, list, true); } /* OpenMP 4.0: diff --git a/gcc/c/c-typeck.c b/gcc/c/c-typeck.c index c0ebb319aff..ee6362d4274 100644 --- a/gcc/c/c-typeck.c +++ b/gcc/c/c-typeck.c @@ -13217,6 +13217,18 @@ handle_omp_array_sections_1 (tree c, tree t, vec &types, t, omp_clause_code_name[OMP_CLAUSE_CODE (c)]); return error_mark_node; } + while (TREE_CODE (t) == INDIRECT_REF) + { + t = TREE_OPERAND (t, 0); + STRIP_NOPS (t); + if (TREE_CODE (t) == POINTER_PLUS_EXPR) + t = TREE_OPERAND (t, 0); + } + while (TREE_CODE (t) == COMPOUND_EXPR) + { + t = TREE_OPERAND (t, 1); + STRIP_NOPS (t); + } if (TREE_CODE (t) == COMPONENT_REF && (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP || OMP_CLAUSE_CODE (c) == OMP_CLAUSE_TO @@ -13238,10 +13250,14 @@ handle_omp_array_sections_1 (tree c, tree t, vec &types, return error_mark_node; } t = TREE_OPERAND (t, 0); - if (TREE_CODE (t) == MEM_REF) + while (TREE_CODE (t) == MEM_REF + || TREE_CODE (t) == INDIRECT_REF + || TREE_CODE (t) == ARRAY_REF) { t = TREE_OPERAND (t, 0); STRIP_NOPS (t); + if (TREE_CODE (t) == POINTER_PLUS_EXPR) + t = TREE_OPERAND (t, 0); } if (ort == C_ORT_ACC && TREE_CODE (t) == MEM_REF) { @@ -13530,15 +13546,25 @@ handle_omp_array_sections_1 (tree c, tree t, vec &types, return error_mark_node; } /* If there is a pointer type anywhere but in the very first - array-section-subscript, the array section can't be contiguous. */ + array-section-subscript, the array section could be non-contiguous. */ if (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_DEPEND && OMP_CLAUSE_CODE (c) != OMP_CLAUSE_AFFINITY && TREE_CODE (TREE_CHAIN (t)) == TREE_LIST) { - error_at (OMP_CLAUSE_LOCATION (c), - "array section is not contiguous in %qs clause", - omp_clause_code_name[OMP_CLAUSE_CODE (c)]); - return error_mark_node; + /* If any prior dimension has a non-one length, then deem this + array section as non-contiguous. */ + for (tree d = TREE_CHAIN (t); TREE_CODE (d) == TREE_LIST; + d = TREE_CHAIN (d)) + { + tree d_length = TREE_VALUE (d); + if (d_length == NULL_TREE || !integer_onep (d_length)) + { + error_at (OMP_CLAUSE_LOCATION (c), + "array section is not contiguous in %qs clause", + omp_clause_code_name[OMP_CLAUSE_CODE (c)]); + return error_mark_node; + } + } } } else @@ -14888,13 +14914,20 @@ c_finish_omp_clauses (tree clauses, enum c_omp_region_type ort) if (TREE_CODE (t) == COMPONENT_REF && TREE_CODE (TREE_TYPE (t)) == ARRAY_TYPE) { - while (TREE_CODE (t) == COMPONENT_REF) - t = TREE_OPERAND (t, 0); - if (TREE_CODE (t) == MEM_REF) + do { t = TREE_OPERAND (t, 0); - STRIP_NOPS (t); + if (TREE_CODE (t) == MEM_REF + || TREE_CODE (t) == INDIRECT_REF) + { + t = TREE_OPERAND (t, 0); + STRIP_NOPS (t); + if (TREE_CODE (t) == POINTER_PLUS_EXPR) + t = TREE_OPERAND (t, 0); + } } + while (TREE_CODE (t) == COMPONENT_REF); + if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP && OMP_CLAUSE_MAP_IMPLICIT (c) && (bitmap_bit_p (&map_head, DECL_UID (t)) @@ -14961,14 +14994,32 @@ c_finish_omp_clauses (tree clauses, enum c_omp_region_type ort) bias) to zero here, so it is not set erroneously to the pointer size later on in gimplify.c. */ OMP_CLAUSE_SIZE (c) = size_zero_node; + while (TREE_CODE (t) == INDIRECT_REF + || TREE_CODE (t) == ARRAY_REF) + { + t = TREE_OPERAND (t, 0); + STRIP_NOPS (t); + if (TREE_CODE (t) == POINTER_PLUS_EXPR) + t = TREE_OPERAND (t, 0); + } + while (TREE_CODE (t) == COMPOUND_EXPR) + { + t = TREE_OPERAND (t, 1); + STRIP_NOPS (t); + } indir_component_ref_p = false; if (TREE_CODE (t) == COMPONENT_REF - && TREE_CODE (TREE_OPERAND (t, 0)) == MEM_REF) + && (TREE_CODE (TREE_OPERAND (t, 0)) == MEM_REF + || TREE_CODE (TREE_OPERAND (t, 0)) == INDIRECT_REF + || TREE_CODE (TREE_OPERAND (t, 0)) == ARRAY_REF)) { t = TREE_OPERAND (TREE_OPERAND (t, 0), 0); indir_component_ref_p = true; STRIP_NOPS (t); + if (TREE_CODE (t) == POINTER_PLUS_EXPR) + t = TREE_OPERAND (t, 0); } + if (TREE_CODE (t) == COMPONENT_REF && OMP_CLAUSE_CODE (c) != OMP_CLAUSE__CACHE_) { @@ -15004,7 +15055,7 @@ c_finish_omp_clauses (tree clauses, enum c_omp_region_type ort) break; } t = TREE_OPERAND (t, 0); - if (ort == C_ORT_ACC && TREE_CODE (t) == MEM_REF) + if (TREE_CODE (t) == MEM_REF) { if (maybe_ne (mem_ref_offset (t), 0)) error_at (OMP_CLAUSE_LOCATION (c), @@ -15013,6 +15064,15 @@ c_finish_omp_clauses (tree clauses, enum c_omp_region_type ort) else t = TREE_OPERAND (t, 0); } + while (TREE_CODE (t) == MEM_REF + || TREE_CODE (t) == INDIRECT_REF + || TREE_CODE (t) == ARRAY_REF) + { + t = TREE_OPERAND (t, 0); + STRIP_NOPS (t); + if (TREE_CODE (t) == POINTER_PLUS_EXPR) + t = TREE_OPERAND (t, 0); + } } if (remove) break; @@ -15084,7 +15144,8 @@ c_finish_omp_clauses (tree clauses, enum c_omp_region_type ort) "%qD appears more than once in data clauses", t); remove = true; } - else if (bitmap_bit_p (&map_head, DECL_UID (t))) + else if (bitmap_bit_p (&map_head, DECL_UID (t)) + && !bitmap_bit_p (&map_field_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 8c27ea12013..fceea16a976 100644 --- a/gcc/cp/parser.c +++ b/gcc/cp/parser.c @@ -36323,11 +36323,22 @@ check_no_duplicate_clause (tree clauses, enum omp_clause_code code, The optional ALLOW_DEREF argument is true if list items can use the deref (->) operator. */ +struct omp_dim +{ + tree low_bound, length; + location_t loc; + bool no_colon; + omp_dim (tree lb, tree len, location_t lo, bool nc) + : low_bound (lb), length (len), loc (lo), no_colon (nc) {} +}; + static tree cp_parser_omp_var_list_no_open (cp_parser *parser, enum omp_clause_code kind, tree list, bool *colon, bool allow_deref = false) { + auto_vec dims; + bool array_section_p; cp_token *token; bool saved_colon_corrects_to_scope_p = parser->colon_corrects_to_scope_p; if (colon) @@ -36408,6 +36419,7 @@ cp_parser_omp_var_list_no_open (cp_parser *parser, enum omp_clause_code kind, case OMP_CLAUSE_MAP: case OMP_CLAUSE_FROM: case OMP_CLAUSE_TO: + start_component_ref: while (cp_lexer_next_token_is (parser->lexer, CPP_DOT) || (allow_deref && cp_lexer_next_token_is (parser->lexer, CPP_DEREF))) @@ -36431,14 +36443,19 @@ cp_parser_omp_var_list_no_open (cp_parser *parser, enum omp_clause_code kind, case OMP_CLAUSE_REDUCTION: case OMP_CLAUSE_IN_REDUCTION: case OMP_CLAUSE_TASK_REDUCTION: + array_section_p = false; + dims.truncate (0); while (cp_lexer_next_token_is (parser->lexer, CPP_OPEN_SQUARE)) { + location_t loc = UNKNOWN_LOCATION; tree low_bound = NULL_TREE, length = NULL_TREE; + bool no_colon = false; parser->colon_corrects_to_scope_p = false; cp_lexer_consume_token (parser->lexer); if (!cp_lexer_next_token_is (parser->lexer, CPP_COLON)) { + loc = cp_lexer_peek_token (parser->lexer)->location; low_bound = cp_parser_expression (parser); /* Later handling is not prepared to see through these. */ gcc_checking_assert (!location_wrapper_p (low_bound)); @@ -36447,7 +36464,10 @@ cp_parser_omp_var_list_no_open (cp_parser *parser, enum omp_clause_code kind, parser->colon_corrects_to_scope_p = saved_colon_corrects_to_scope_p; if (cp_lexer_next_token_is (parser->lexer, CPP_CLOSE_SQUARE)) - length = integer_one_node; + { + length = integer_one_node; + no_colon = true; + } else { /* Look for `:'. */ @@ -36460,6 +36480,8 @@ cp_parser_omp_var_list_no_open (cp_parser *parser, enum omp_clause_code kind, } if (kind == OMP_CLAUSE_DEPEND || kind == OMP_CLAUSE_AFFINITY) cp_parser_commit_to_tentative_parse (parser); + else + array_section_p = true; if (!cp_lexer_next_token_is (parser->lexer, CPP_CLOSE_SQUARE)) { @@ -36478,8 +36500,32 @@ cp_parser_omp_var_list_no_open (cp_parser *parser, enum omp_clause_code kind, goto skip_comma; } - decl = tree_cons (low_bound, length, decl); + dims.safe_push (omp_dim (low_bound, length, loc, no_colon)); + } + + if ((kind == OMP_CLAUSE_MAP + || kind == OMP_CLAUSE_FROM + || kind == OMP_CLAUSE_TO) + && !array_section_p + && (cp_lexer_next_token_is (parser->lexer, CPP_DOT) + || (allow_deref + && cp_lexer_next_token_is (parser->lexer, + CPP_DEREF)))) + { + for (unsigned i = 0; i < dims.length (); i++) + { + gcc_assert (dims[i].length == integer_one_node); + decl = build_array_ref (dims[i].loc, + decl, dims[i].low_bound); + } + goto start_component_ref; } + else + { + for (unsigned i = 0; i < dims.length (); i++) + decl = tree_cons (dims[i].low_bound, dims[i].length, decl); + } + break; default: break; @@ -39981,11 +40027,13 @@ cp_parser_omp_all_clauses (cp_parser *parser, omp_clause_mask mask, clauses = cp_parser_omp_var_list (parser, OMP_CLAUSE_TO_DECLARE, clauses); else - clauses = cp_parser_omp_var_list (parser, OMP_CLAUSE_TO, clauses); + clauses = cp_parser_omp_var_list (parser, OMP_CLAUSE_TO, clauses, + true); c_name = "to"; break; case PRAGMA_OMP_CLAUSE_FROM: - clauses = cp_parser_omp_var_list (parser, OMP_CLAUSE_FROM, clauses); + clauses = cp_parser_omp_var_list (parser, OMP_CLAUSE_FROM, clauses, + true); c_name = "from"; break; case PRAGMA_OMP_CLAUSE_UNIFORM: diff --git a/gcc/cp/semantics.c b/gcc/cp/semantics.c index c64b45c0cee..898fcea6bb6 100644 --- a/gcc/cp/semantics.c +++ b/gcc/cp/semantics.c @@ -5030,6 +5030,18 @@ handle_omp_array_sections_1 (tree c, tree t, vec &types, && TREE_CODE (TREE_OPERAND (t, 0)) == COMPONENT_REF) t = TREE_OPERAND (t, 0); ret = t; + while (TREE_CODE (t) == INDIRECT_REF) + { + t = TREE_OPERAND (t, 0); + STRIP_NOPS (t); + if (TREE_CODE (t) == POINTER_PLUS_EXPR) + t = TREE_OPERAND (t, 0); + } + while (TREE_CODE (t) == COMPOUND_EXPR) + { + t = TREE_OPERAND (t, 1); + STRIP_NOPS (t); + } if (TREE_CODE (t) == COMPONENT_REF && (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP || OMP_CLAUSE_CODE (c) == OMP_CLAUSE_TO @@ -5054,10 +5066,14 @@ handle_omp_array_sections_1 (tree c, tree t, vec &types, return error_mark_node; } t = TREE_OPERAND (t, 0); - if (TREE_CODE (t) == INDIRECT_REF) + while (TREE_CODE (t) == MEM_REF + || TREE_CODE (t) == INDIRECT_REF + || TREE_CODE (t) == ARRAY_REF) { t = TREE_OPERAND (t, 0); STRIP_NOPS (t); + if (TREE_CODE (t) == POINTER_PLUS_EXPR) + t = TREE_OPERAND (t, 0); } } if (REFERENCE_REF_P (t)) @@ -5341,15 +5357,25 @@ handle_omp_array_sections_1 (tree c, tree t, vec &types, return error_mark_node; } /* If there is a pointer type anywhere but in the very first - array-section-subscript, the array section can't be contiguous. */ + array-section-subscript, the array section could be non-contiguous. */ if (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_AFFINITY && OMP_CLAUSE_CODE (c) != OMP_CLAUSE_DEPEND && TREE_CODE (TREE_CHAIN (t)) == TREE_LIST) { - error_at (OMP_CLAUSE_LOCATION (c), - "array section is not contiguous in %qs clause", - omp_clause_code_name[OMP_CLAUSE_CODE (c)]); - return error_mark_node; + /* If any prior dimension has a non-one length, then deem this + array section as non-contiguous. */ + for (tree d = TREE_CHAIN (t); TREE_CODE (d) == TREE_LIST; + d = TREE_CHAIN (d)) + { + tree d_length = TREE_VALUE (d); + if (d_length == NULL_TREE || !integer_onep (d_length)) + { + error_at (OMP_CLAUSE_LOCATION (c), + "array section is not contiguous in %qs clause", + omp_clause_code_name[OMP_CLAUSE_CODE (c)]); + return error_mark_node; + } + } } } else @@ -5619,16 +5645,37 @@ handle_omp_array_sections (tree c, enum c_omp_region_type ort) default: break; } + bool reference_always_pointer = true; tree c2 = build_omp_clause (OMP_CLAUSE_LOCATION (c), OMP_CLAUSE_MAP); if (TREE_CODE (t) == COMPONENT_REF) - OMP_CLAUSE_SET_MAP_KIND (c2, GOMP_MAP_ATTACH_DETACH); + { + OMP_CLAUSE_SET_MAP_KIND (c2, GOMP_MAP_ATTACH_DETACH); + + if ((ort & C_ORT_OMP_DECLARE_SIMD) == C_ORT_OMP + && TYPE_REF_P (TREE_TYPE (t))) + { + if (TREE_CODE (TREE_TYPE (TREE_TYPE (t))) == ARRAY_TYPE) + OMP_CLAUSE_SET_MAP_KIND (c2, GOMP_MAP_ALWAYS_POINTER); + else + t = convert_from_reference (t); + + reference_always_pointer = false; + } + } else if (REFERENCE_REF_P (t) && TREE_CODE (TREE_OPERAND (t, 0)) == COMPONENT_REF) { - t = TREE_OPERAND (t, 0); - gomp_map_kind k = (ort == C_ORT_ACC) ? GOMP_MAP_ATTACH_DETACH - : GOMP_MAP_ALWAYS_POINTER; + gomp_map_kind k; + if ((ort & C_ORT_OMP_DECLARE_SIMD) == C_ORT_OMP + && TREE_CODE (TREE_TYPE (t)) == POINTER_TYPE) + k = GOMP_MAP_ATTACH_DETACH; + else + { + t = TREE_OPERAND (t, 0); + k = (ort == C_ORT_ACC + ? GOMP_MAP_ATTACH_DETACH : GOMP_MAP_ALWAYS_POINTER); + } OMP_CLAUSE_SET_MAP_KIND (c2, k); } else @@ -5652,8 +5699,10 @@ handle_omp_array_sections (tree c, enum c_omp_region_type ort) OMP_CLAUSE_SIZE (c2) = t; OMP_CLAUSE_CHAIN (c2) = OMP_CLAUSE_CHAIN (c); OMP_CLAUSE_CHAIN (c) = c2; + ptr = OMP_CLAUSE_DECL (c2); - if (OMP_CLAUSE_MAP_KIND (c2) != GOMP_MAP_FIRSTPRIVATE_POINTER + if (reference_always_pointer + && OMP_CLAUSE_MAP_KIND (c2) != GOMP_MAP_FIRSTPRIVATE_POINTER && TYPE_REF_P (TREE_TYPE (ptr)) && INDIRECT_TYPE_P (TREE_TYPE (TREE_TYPE (ptr)))) { @@ -7855,15 +7904,22 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort) if (TREE_CODE (t) == COMPONENT_REF && TREE_CODE (TREE_TYPE (t)) == ARRAY_TYPE) { - while (TREE_CODE (t) == COMPONENT_REF) - t = TREE_OPERAND (t, 0); - if (REFERENCE_REF_P (t)) - t = TREE_OPERAND (t, 0); - if (TREE_CODE (t) == INDIRECT_REF) + do { t = TREE_OPERAND (t, 0); - STRIP_NOPS (t); + if (REFERENCE_REF_P (t)) + t = TREE_OPERAND (t, 0); + if (TREE_CODE (t) == MEM_REF + || TREE_CODE (t) == INDIRECT_REF) + { + t = TREE_OPERAND (t, 0); + STRIP_NOPS (t); + if (TREE_CODE (t) == POINTER_PLUS_EXPR) + t = TREE_OPERAND (t, 0); + } } + while (TREE_CODE (t) == COMPONENT_REF); + if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP && OMP_CLAUSE_MAP_IMPLICIT (c) && (bitmap_bit_p (&map_head, DECL_UID (t)) @@ -7934,15 +7990,33 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort) && TREE_CODE (TREE_OPERAND (t, 0)) == COMPONENT_REF) { t = TREE_OPERAND (t, 0); - OMP_CLAUSE_DECL (c) = t; + if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP + && OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_ATTACH_DETACH) + OMP_CLAUSE_DECL (c) = t; + } + while (TREE_CODE (t) == INDIRECT_REF + || TREE_CODE (t) == ARRAY_REF) + { + t = TREE_OPERAND (t, 0); + STRIP_NOPS (t); + if (TREE_CODE (t) == POINTER_PLUS_EXPR) + t = TREE_OPERAND (t, 0); + } + while (TREE_CODE (t) == COMPOUND_EXPR) + { + t = TREE_OPERAND (t, 1); + STRIP_NOPS (t); } indir_component_ref_p = false; if (TREE_CODE (t) == COMPONENT_REF - && TREE_CODE (TREE_OPERAND (t, 0)) == INDIRECT_REF) + && (TREE_CODE (TREE_OPERAND (t, 0)) == INDIRECT_REF + || TREE_CODE (TREE_OPERAND (t, 0)) == ARRAY_REF)) { t = TREE_OPERAND (TREE_OPERAND (t, 0), 0); indir_component_ref_p = true; STRIP_NOPS (t); + if (TREE_CODE (t) == POINTER_PLUS_EXPR) + t = TREE_OPERAND (t, 0); } if (TREE_CODE (t) == COMPONENT_REF && OMP_CLAUSE_CODE (c) != OMP_CLAUSE__CACHE_) @@ -7977,6 +8051,24 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort) break; } t = TREE_OPERAND (t, 0); + if (TREE_CODE (t) == MEM_REF) + { + if (maybe_ne (mem_ref_offset (t), 0)) + error_at (OMP_CLAUSE_LOCATION (c), + "cannot dereference %qE in %qs clause", t, + omp_clause_code_name[OMP_CLAUSE_CODE (c)]); + else + t = TREE_OPERAND (t, 0); + } + while (TREE_CODE (t) == MEM_REF + || TREE_CODE (t) == INDIRECT_REF + || TREE_CODE (t) == ARRAY_REF) + { + t = TREE_OPERAND (t, 0); + STRIP_NOPS (t); + if (TREE_CODE (t) == POINTER_PLUS_EXPR) + t = TREE_OPERAND (t, 0); + } } if (remove) break; @@ -8074,7 +8166,8 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort) "%qD appears more than once in data clauses", t); remove = true; } - else if (bitmap_bit_p (&map_head, DECL_UID (t))) + else if (bitmap_bit_p (&map_head, DECL_UID (t)) + && !bitmap_bit_p (&map_field_head, DECL_UID (t))) { if (ort == C_ORT_ACC) error_at (OMP_CLAUSE_LOCATION (c), @@ -8121,8 +8214,13 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort) else { bitmap_set_bit (&map_head, DECL_UID (t)); - if (t != OMP_CLAUSE_DECL (c) - && TREE_CODE (OMP_CLAUSE_DECL (c)) == COMPONENT_REF) + + tree decl = OMP_CLAUSE_DECL (c); + if (t != decl + && (TREE_CODE (decl) == COMPONENT_REF + || (INDIRECT_REF_P (decl) + && TREE_CODE (TREE_OPERAND (decl, 0)) == COMPONENT_REF + && TYPE_REF_P (TREE_TYPE (TREE_OPERAND (decl, 0)))))) bitmap_set_bit (&map_field_head, DECL_UID (t)); } handle_map_references: @@ -8151,7 +8249,7 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort) tree c2 = build_omp_clause (OMP_CLAUSE_LOCATION (c), OMP_CLAUSE_MAP); if (TREE_CODE (t) == COMPONENT_REF) - OMP_CLAUSE_SET_MAP_KIND (c2, GOMP_MAP_ATTACH_DETACH); + OMP_CLAUSE_SET_MAP_KIND (c2, GOMP_MAP_ALWAYS_POINTER); else OMP_CLAUSE_SET_MAP_KIND (c2, GOMP_MAP_FIRSTPRIVATE_REFERENCE); diff --git a/gcc/fortran/trans-openmp.c b/gcc/fortran/trans-openmp.c index b86c7cf9833..a762d715f4c 100644 --- a/gcc/fortran/trans-openmp.c +++ b/gcc/fortran/trans-openmp.c @@ -2460,6 +2460,9 @@ gfc_trans_omp_array_section (stmtblock_t *block, gfc_omp_namelist *n, TREE_TYPE (TREE_TYPE (decl)), decl, offset, NULL_TREE, NULL_TREE); OMP_CLAUSE_DECL (node) = offset; + + if (ptr_kind == GOMP_MAP_ALWAYS_POINTER) + return; } else { diff --git a/gcc/gimplify.c b/gcc/gimplify.c index ed46fe3c461..9d37766f853 100644 --- a/gcc/gimplify.c +++ b/gcc/gimplify.c @@ -8666,7 +8666,7 @@ insert_struct_comp_map (enum tree_code code, tree c, tree struct_node, static tree extract_base_bit_offset (tree base, tree *base_ref, poly_int64 *bitposp, - poly_offset_int *poffsetp) + poly_offset_int *poffsetp, tree *offsetp) { tree offset; poly_int64 bitsize, bitpos; @@ -8713,10 +8713,11 @@ extract_base_bit_offset (tree base, tree *base_ref, poly_int64 *bitposp, && TREE_CODE (TREE_TYPE (TREE_OPERAND (base, 0))) == REFERENCE_TYPE) base = TREE_OPERAND (base, 0); - gcc_assert (offset == NULL_TREE || poly_int_tree_p (offset)); - - if (offset) - poffset = wi::to_poly_offset (offset); + if (offset && poly_int_tree_p (offset)) + { + poffset = wi::to_poly_offset (offset); + offset = NULL_TREE; + } else poffset = 0; @@ -8725,6 +8726,7 @@ extract_base_bit_offset (tree base, tree *base_ref, poly_int64 *bitposp, *bitposp = bitpos; *poffsetp = poffset; + *offsetp = offset; /* Set *BASE_REF if BASE was a dereferenced reference variable. */ if (base_ref && orig_base != base) @@ -8738,12 +8740,22 @@ extract_base_bit_offset (tree base, tree *base_ref, poly_int64 *bitposp, 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; + if ((TREE_CODE (expr) == INDIRECT_REF && TREE_CODE (base_ptr) == MEM_REF) + || (TREE_CODE (expr) == MEM_REF && TREE_CODE (base_ptr) == INDIRECT_REF)) + return operand_equal_p (TREE_OPERAND (expr, 0), + TREE_OPERAND (base_ptr, 0)); + while (!operand_equal_p (expr, base_ptr)) + { + if (TREE_CODE (base_ptr) == COMPOUND_EXPR) + base_ptr = TREE_OPERAND (base_ptr, 1); + if (TREE_CODE (base_ptr) == COMPONENT_REF + || TREE_CODE (base_ptr) == POINTER_PLUS_EXPR + || TREE_CODE (base_ptr) == SAVE_EXPR) + base_ptr = TREE_OPERAND (base_ptr, 0); + else + break; + } + return operand_equal_p (expr, base_ptr); } /* Implement OpenMP 5.x map ordering rules for target directives. There are @@ -8823,21 +8835,107 @@ omp_target_reorder_clauses (tree *list_p) 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; + if (atf[j]) + { + tree *cp2 = atf[j]; + tree decl2 = OMP_CLAUSE_DECL (*cp2); + + 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; + + if (*cp2 != NULL_TREE + && OMP_CLAUSE_CODE (*cp2) == OMP_CLAUSE_MAP + && OMP_CLAUSE_MAP_KIND (*cp2) == GOMP_MAP_ALWAYS_POINTER) + { + tree c2 = *cp2; + *cp2 = OMP_CLAUSE_CHAIN (c2); + OMP_CLAUSE_CHAIN (c2) = OMP_CLAUSE_CHAIN (c); + OMP_CLAUSE_CHAIN (c) = c2; + } + + atf[j] = NULL; } - } + } } } + + /* For attach_detach map clauses, if there is another map that maps the + attached/detached pointer, make sure that map is ordered before the + attach_detach. */ + atf.truncate (0); + 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 clauses, and + always_pointer/attach_detach clauses. */ + 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 + || k == GOMP_MAP_ATTACH_DETACH + || k == GOMP_MAP_ALWAYS_POINTER) + atf.safe_push (cp); + } + + for (unsigned int i = 0; i < atf.length (); i++) + if (atf[i]) + { + tree *cp = atf[i]; + tree ptr = OMP_CLAUSE_DECL (*cp); + STRIP_TYPE_NOPS (ptr); + if (OMP_CLAUSE_MAP_KIND (*cp) == GOMP_MAP_ATTACH_DETACH) + for (unsigned int j = i + 1; j < atf.length (); j++) + { + tree *cp2 = atf[j]; + tree decl2 = OMP_CLAUSE_DECL (*cp2); + if (OMP_CLAUSE_MAP_KIND (*cp2) != GOMP_MAP_ATTACH_DETACH + && OMP_CLAUSE_MAP_KIND (*cp2) != GOMP_MAP_ALWAYS_POINTER + && is_or_contains_p (decl2, ptr)) + { + /* Move *cp2 to before *cp. */ + tree c = *cp2; + *cp2 = OMP_CLAUSE_CHAIN (c); + OMP_CLAUSE_CHAIN (c) = *cp; + *cp = c; + atf[j] = NULL; + + /* If decl2 is of the form '*decl2_opnd0', and followed by an + ALWAYS_POINTER or ATTACH_DETACH of 'decl2_opnd0', move the + pointer operation along with *cp2. This can happen for C++ + reference sequences. */ + if (j + 1 < atf.length () + && (TREE_CODE (decl2) == INDIRECT_REF + || TREE_CODE (decl2) == MEM_REF)) + { + tree *cp3 = atf[j + 1]; + tree decl3 = OMP_CLAUSE_DECL (*cp3); + tree decl2_opnd0 = TREE_OPERAND (decl2, 0); + if ((OMP_CLAUSE_MAP_KIND (*cp3) == GOMP_MAP_ALWAYS_POINTER + || OMP_CLAUSE_MAP_KIND (*cp3) == GOMP_MAP_ATTACH_DETACH) + && operand_equal_p (decl3, decl2_opnd0)) + { + /* Also move *cp3 to before *cp. */ + c = *cp3; + *cp2 = OMP_CLAUSE_CHAIN (c); + OMP_CLAUSE_CHAIN (c) = *cp; + *cp = c; + atf[j + 1] = NULL; + j += 1; + } + } + } + } + } } /* DECL is supposed to have lastprivate semantics in the outer contexts @@ -8929,6 +9027,7 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p, struct gimplify_omp_ctx *ctx, *outer_ctx; tree c; hash_map *struct_map_to_clause = NULL; + hash_map *struct_seen_clause = NULL; hash_set *struct_deref_set = NULL; tree *prev_list_p = NULL, *orig_list_p = list_p; int handled_depend_iterators = -1; @@ -9404,6 +9503,7 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p, } bool indir_p = false; bool component_ref_p = false; + tree indir_base = NULL_TREE; tree orig_decl = decl; tree decl_ref = NULL_TREE; if ((region_type & (ORT_ACC | ORT_TARGET | ORT_TARGET_DATA)) != 0 @@ -9422,6 +9522,7 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p, == POINTER_TYPE)) { indir_p = true; + indir_base = decl; decl = TREE_OPERAND (decl, 0); STRIP_NOPS (decl); } @@ -9468,7 +9569,9 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p, != GOMP_MAP_POINTER) || OMP_CLAUSE_DECL (next_clause) != decl) && (!struct_deref_set - || !struct_deref_set->contains (decl))) + || !struct_deref_set->contains (decl)) + && (!struct_map_to_clause + || !struct_map_to_clause->get (indir_base))) { if (!struct_deref_set) struct_deref_set = new hash_set (); @@ -9512,7 +9615,8 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p, if ((DECL_P (decl) || (component_ref_p && (INDIRECT_REF_P (decl) - || TREE_CODE (decl) == MEM_REF))) + || TREE_CODE (decl) == MEM_REF + || TREE_CODE (decl) == ARRAY_REF))) && 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 @@ -9547,7 +9651,15 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p, remove = true; break; } - if (OMP_CLAUSE_CHAIN (*prev_list_p) != c) + + /* The below prev_list_p based error recovery code is + currently no longer valid for OpenMP. */ + if (code != OMP_TARGET + && code != OMP_TARGET_DATA + && code != OMP_TARGET_UPDATE + && code != OMP_TARGET_ENTER_DATA + && code != OMP_TARGET_EXIT_DATA + && OMP_CLAUSE_CHAIN (*prev_list_p) != c) { tree ch = OMP_CLAUSE_CHAIN (*prev_list_p); if (ch == NULL_TREE || OMP_CLAUSE_CHAIN (ch) != c) @@ -9560,13 +9672,15 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p, poly_offset_int offset1; poly_int64 bitpos1; + tree tree_offset1; tree base_ref; tree base = extract_base_bit_offset (OMP_CLAUSE_DECL (c), &base_ref, - &bitpos1, &offset1); + &bitpos1, &offset1, + &tree_offset1); - gcc_assert (base == decl); + bool do_map_struct = (base == decl && !tree_offset1); splay_tree_node n = (DECL_P (decl) @@ -9598,6 +9712,32 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p, OMP_CLAUSE_SET_MAP_KIND (c, k); has_attachments = true; } + + /* We currently don't handle non-constant offset accesses wrt to + GOMP_MAP_STRUCT elements. */ + if (!do_map_struct) + goto skip_map_struct; + + /* Nor for attach_detach for OpenMP. */ + if ((code == OMP_TARGET + || code == OMP_TARGET_DATA + || code == OMP_TARGET_UPDATE + || code == OMP_TARGET_ENTER_DATA + || code == OMP_TARGET_EXIT_DATA) + && attach_detach) + { + if (DECL_P (decl)) + { + if (struct_seen_clause == NULL) + struct_seen_clause + = new hash_map; + if (!struct_seen_clause->get (decl)) + struct_seen_clause->put (decl, list_p); + } + + goto skip_map_struct; + } + if ((DECL_P (decl) && (n == NULL || (n->value & GOVD_MAP) == 0)) || (!DECL_P (decl) @@ -9637,9 +9777,14 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p, struct_map_to_clause->put (decl, l); if (ptr || attach_detach) { - insert_struct_comp_map (code, c, l, *prev_list_p, + tree **sc = (struct_seen_clause + ? struct_seen_clause->get (decl) + : NULL); + tree *insert_node_pos = sc ? *sc : prev_list_p; + + insert_struct_comp_map (code, c, l, *insert_node_pos, NULL); - *prev_list_p = l; + *insert_node_pos = l; prev_list_p = NULL; } else @@ -9725,9 +9870,11 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p, tree sc_decl = OMP_CLAUSE_DECL (*sc); poly_offset_int offsetn; poly_int64 bitposn; + tree tree_offsetn; tree base = extract_base_bit_offset (sc_decl, NULL, - &bitposn, &offsetn); + &bitposn, &offsetn, + &tree_offsetn); if (base != decl) break; if (scp) @@ -9815,16 +9962,21 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p, continue; } } + skip_map_struct: + ; } else if ((code == OACC_ENTER_DATA || code == OACC_EXIT_DATA || code == OACC_DATA || code == OACC_PARALLEL || code == OACC_KERNELS - || code == OACC_SERIAL) + || code == OACC_SERIAL + || code == OMP_TARGET_ENTER_DATA + || code == OMP_TARGET_EXIT_DATA) && OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ATTACH_DETACH) { - gomp_map_kind k = (code == OACC_EXIT_DATA + gomp_map_kind k = ((code == OACC_EXIT_DATA + || code == OMP_TARGET_EXIT_DATA) ? GOMP_MAP_DETACH : GOMP_MAP_ATTACH); OMP_CLAUSE_SET_MAP_KIND (c, k); } @@ -10656,6 +10808,8 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p, ctx->clauses = *orig_list_p; gimplify_omp_ctxp = ctx; + if (struct_seen_clause) + delete struct_seen_clause; if (struct_map_to_clause) delete struct_map_to_clause; if (struct_deref_set) diff --git a/gcc/testsuite/c-c++-common/goacc/deep-copy-arrayofstruct.c b/gcc/testsuite/c-c++-common/goacc/deep-copy-arrayofstruct.c index d411bcfa8e7..4247607b61c 100644 --- a/gcc/testsuite/c-c++-common/goacc/deep-copy-arrayofstruct.c +++ b/gcc/testsuite/c-c++-common/goacc/deep-copy-arrayofstruct.c @@ -37,13 +37,12 @@ int main(int argc, char* argv[]) { int j, k; for (k = 0; k < S; k++) -#pragma acc parallel loop copy(m[k].a[0:N]) /* { dg-error "expected .\\\). before .\\\.. token" } */ +#pragma acc parallel loop copy(m[k].a[0:N]) for (j = 0; j < N; j++) m[k].a[j]++; for (k = 0; k < S; k++) -#pragma acc parallel loop copy(m[k].b[0:N], m[k].c[5:N-10]) /* { dg-error "expected .\\\). before .\\\.. token" } */ - /* { dg-error ".m. appears more than once in data clauses" "" { target c++ } .-1 } */ +#pragma acc parallel loop copy(m[k].b[0:N], m[k].c[5:N-10]) for (j = 0; j < N; j++) { m[k].b[j]++; diff --git a/gcc/testsuite/c-c++-common/gomp/target-enter-data-1.c b/gcc/testsuite/c-c++-common/gomp/target-enter-data-1.c new file mode 100644 index 00000000000..ce766d29e2d --- /dev/null +++ b/gcc/testsuite/c-c++-common/gomp/target-enter-data-1.c @@ -0,0 +1,24 @@ +/* { dg-do compile } */ +/* { dg-additional-options "-fopenmp -fdump-tree-gimple" } */ + +struct bar +{ + int num_vectors; + double *vectors; +}; + +struct foo +{ + int num_vectors; + struct bar *bars; + double **vectors; +}; + +void func (struct foo *f, int n, int m) +{ + #pragma omp target enter data map (to: f->vectors[m][:n]) + #pragma omp target enter data map (to: f->bars[n].vectors[:m]) + #pragma omp target enter data map (to: f->bars[n].vectors[:f->bars[n].num_vectors]) +} + +/* { dg-final { scan-tree-dump-times "map\\(to:\\*_\[0-9\]+ \\\[len: _\[0-9\]+\\\]\\) map\\(attach:\[^-\]+->vectors \\\[bias: \[^\]\]+\\\]\\)" 3 "gimple" } } */ diff --git a/gcc/testsuite/c-c++-common/gomp/target-implicit-map-2.c b/gcc/testsuite/c-c++-common/gomp/target-implicit-map-2.c new file mode 100644 index 00000000000..3aa1a8fc55e --- /dev/null +++ b/gcc/testsuite/c-c++-common/gomp/target-implicit-map-2.c @@ -0,0 +1,52 @@ +/* { dg-do compile } */ +/* { dg-additional-options "-fdump-tree-gimple" } */ +#include + +#define N 10 + +struct S +{ + int a, b; + int *ptr; + int c, d; +}; + +int +main (void) +{ + struct S a; + a.ptr = (int *) malloc (sizeof (int) * N); + + for (int i = 0; i < N; i++) + a.ptr[i] = 0; + + #pragma omp target enter data map(to: a.ptr, a.ptr[:N]) + + #pragma omp target + for (int i = 0; i < N; i++) + a.ptr[i] += 1; + + #pragma omp target update from(a.ptr[:N]) + + for (int i = 0; i < N; i++) + if (a.ptr[i] != 1) + abort (); + + #pragma omp target map(a.ptr[:N]) + for (int i = 0; i < N; i++) + a.ptr[i] += 1; + + #pragma omp target update from(a.ptr[:N]) + + for (int i = 0; i < N; i++) + if (a.ptr[i] != 2) + abort (); + + #pragma omp target exit data map(from:a.ptr, a.ptr[:N]) + + return 0; +} + +/* { dg-final { scan-tree-dump {#pragma omp target num_teams.* map\(tofrom:a \[len: [0-9]+\]\[implicit\]\)} "gimple" } } */ + +/* { dg-final { scan-tree-dump {#pragma omp target num_teams.* map\(tofrom:a \[len: [0-9]+\]\[implicit\]\) map\(tofrom:\*_[0-9]+ \[len: [0-9]+\]\) map\(attach:a\.ptr \[bias: 0\]\)} "gimple" } } */ diff --git a/libgomp/target.c b/libgomp/target.c index bb31b1991d1..917da807d9c 100644 --- a/libgomp/target.c +++ b/libgomp/target.c @@ -581,11 +581,30 @@ gomp_map_vars_existing (struct gomp_device_descr *devicep, address/length adjustment is a TODO. */ assert (!implicit_subset); - gomp_copy_host2dev (devicep, aq, - (void *) (oldn->tgt->tgt_start + oldn->tgt_offset - + newn->host_start - oldn->host_start), - (void *) newn->host_start, - newn->host_end - newn->host_start, false, cbuf); + if (oldn->aux && oldn->aux->attach_count) + { + /* We have to be careful not to overwrite still attached pointers + during the copyback to host. */ + uintptr_t addr = newn->host_start; + while (addr < newn->host_end) + { + size_t i = (addr - oldn->host_start) / sizeof (void *); + if (oldn->aux->attach_count[i] == 0) + gomp_copy_host2dev (devicep, aq, + (void *) (oldn->tgt->tgt_start + + oldn->tgt_offset + + addr - oldn->host_start), + (void *) addr, + sizeof (void *), false, cbuf); + addr += sizeof (void *); + } + } + else + gomp_copy_host2dev (devicep, aq, + (void *) (oldn->tgt->tgt_start + oldn->tgt_offset + + newn->host_start - oldn->host_start), + (void *) newn->host_start, + newn->host_end - newn->host_start, false, cbuf); } gomp_increment_refcount (oldn, refcount_set); @@ -2011,17 +2030,45 @@ gomp_update (struct gomp_device_descr *devicep, size_t mapnum, void **hostaddrs, (void *) n->host_end); } - - void *hostaddr = (void *) cur_node.host_start; - void *devaddr = (void *) (n->tgt->tgt_start + n->tgt_offset - + cur_node.host_start - n->host_start); - size_t size = cur_node.host_end - cur_node.host_start; - - if (GOMP_MAP_COPY_TO_P (kind & typemask)) - gomp_copy_host2dev (devicep, NULL, devaddr, hostaddr, size, - false, NULL); - if (GOMP_MAP_COPY_FROM_P (kind & typemask)) - gomp_copy_dev2host (devicep, NULL, hostaddr, devaddr, size); + if (n->aux && n->aux->attach_count) + { + uintptr_t addr = cur_node.host_start; + while (addr < cur_node.host_end) + { + /* We have to be careful not to overwrite still attached + pointers during host<->device updates. */ + size_t i = (addr - cur_node.host_start) / sizeof (void *); + if (n->aux->attach_count[i] == 0) + { + void *devaddr = (void *) (n->tgt->tgt_start + + n->tgt_offset + + addr - n->host_start); + if (GOMP_MAP_COPY_TO_P (kind & typemask)) + gomp_copy_host2dev (devicep, NULL, + devaddr, (void *) addr, + sizeof (void *), false, NULL); + if (GOMP_MAP_COPY_FROM_P (kind & typemask)) + gomp_copy_dev2host (devicep, NULL, + (void *) addr, devaddr, + sizeof (void *)); + } + addr += sizeof (void *); + } + } + else + { + void *hostaddr = (void *) cur_node.host_start; + void *devaddr = (void *) (n->tgt->tgt_start + n->tgt_offset + + cur_node.host_start + - n->host_start); + size_t size = cur_node.host_end - cur_node.host_start; + + if (GOMP_MAP_COPY_TO_P (kind & typemask)) + gomp_copy_host2dev (devicep, NULL, devaddr, hostaddr, size, + false, NULL); + if (GOMP_MAP_COPY_FROM_P (kind & typemask)) + gomp_copy_dev2host (devicep, NULL, hostaddr, devaddr, size); + } } } gomp_mutex_unlock (&devicep->lock); @@ -2915,11 +2962,31 @@ gomp_exit_data (struct gomp_device_descr *devicep, size_t mapnum, if ((kind == GOMP_MAP_FROM && do_copy) || kind == GOMP_MAP_ALWAYS_FROM) - gomp_copy_dev2host (devicep, NULL, (void *) cur_node.host_start, - (void *) (k->tgt->tgt_start + k->tgt_offset - + cur_node.host_start - - k->host_start), - cur_node.host_end - cur_node.host_start); + { + if (k->aux && k->aux->attach_count) + { + /* We have to be careful not to overwrite still attached + pointers during the copyback to host. */ + uintptr_t addr = k->host_start; + while (addr < k->host_end) + { + size_t i = (addr - k->host_start) / sizeof (void *); + if (k->aux->attach_count[i] == 0) + gomp_copy_dev2host (devicep, NULL, (void *) addr, + (void *) (k->tgt->tgt_start + + k->tgt_offset + + addr - k->host_start), + sizeof (void *)); + addr += sizeof (void *); + } + } + else + gomp_copy_dev2host (devicep, NULL, (void *) cur_node.host_start, + (void *) (k->tgt->tgt_start + k->tgt_offset + + cur_node.host_start + - k->host_start), + cur_node.host_end - cur_node.host_start); + } /* Structure elements lists are removed altogether at once, which may cause immediate deallocation of the target_mem_desc, causing diff --git a/libgomp/testsuite/libgomp.c++/target-11.C b/libgomp/testsuite/libgomp.c++/target-11.C index fe99603351d..87c2980b4b5 100644 --- a/libgomp/testsuite/libgomp.c++/target-11.C +++ b/libgomp/testsuite/libgomp.c++/target-11.C @@ -23,9 +23,11 @@ foo () e = c + 18; D s = { a, b + 2, { 0, a + 16, 0, d }, c + 3, e }; int err = 0; - #pragma omp target map (to:s.v.b[0:z + 7], s.template u[z + 1:z + 4]) \ - map (tofrom:s.s[3:3], s. template v. template d[z + 1:z + 3]) \ - map (from: s.w[z:4], s.x[1:3], err) private (i) + #pragma omp target map (to: s.v.b, s.v.b[0:z + 7]) \ + map (s.template u, s.template u[z + 1:z + 4]) \ + map (tofrom: s.s, s.s[3:3]) \ + map (tofrom: s. template v. template d[z + 1:z + 3])\ + map (from: s.w, s.w[z:4], s.x, s.x[1:3], err) private (i) { err = 0; for (i = 0; i < 7; i++) @@ -80,9 +82,9 @@ main () e = c + 18; S s = { a, b + 2, { 0, a + 16, 0, d }, c + 3, e }; int err = 0; - #pragma omp target map (to:s.v.b[0:z + 7], s.u[z + 1:z + 4]) \ - map (tofrom:s.s[3:3], s.v.d[z + 1:z + 3]) \ - map (from: s.w[z:4], s.x[1:3], err) private (i) + #pragma omp target map (to: s.v.b, s.v.b[0:z + 7], s.u, s.u[z + 1:z + 4]) \ + map (tofrom: s.s, s.s[3:3], s.v.d[z + 1:z + 3]) \ + map (from: s.w, s.w[z:4], s.x, s.x[1:3], err) private (i) { err = 0; for (i = 0; i < 7; i++) diff --git a/libgomp/testsuite/libgomp.c++/target-12.C b/libgomp/testsuite/libgomp.c++/target-12.C index 3b4ed57df68..480e479c262 100644 --- a/libgomp/testsuite/libgomp.c++/target-12.C +++ b/libgomp/testsuite/libgomp.c++/target-12.C @@ -53,7 +53,7 @@ main () int u[10] = { 0, 1, 2, 3, 4, 5, 6, 7, 8, 9 }, err = 0; S s = { 9, u + 3, { 10, 11, 12, 13, 14 } }; int *v = u + 4; - #pragma omp target enter data map (to: s.s, s.u[0:5]) map (alloc: s.v[1:3]) + #pragma omp target enter data map (to: s.s, s.u, s.u[0:5]) map (alloc: s.v[1:3]) s.s++; u[3]++; s.v[1]++; diff --git a/libgomp/testsuite/libgomp.c++/target-15.C b/libgomp/testsuite/libgomp.c++/target-15.C index 4b320c31229..53626b2547e 100644 --- a/libgomp/testsuite/libgomp.c++/target-15.C +++ b/libgomp/testsuite/libgomp.c++/target-15.C @@ -14,7 +14,7 @@ foo (S s) d = id; int err; - #pragma omp target map(tofrom: s.a, s.b, s.c[1:2], s.d[-2:3], s.e, s.f, s.g[1:2], s.h[2:3]) map(to: sep) map(from: err) + #pragma omp target map(tofrom: s.a, s.b, s.c[1:2], s.d, s.d[-2:3], s.e, s.f, s.g[1:2], s.h, s.h[2:3]) map(to: sep) map(from: err) { err = s.a != 11 || s.b[0] != 12 || s.b[1] != 13; err |= s.c[1] != 15 || s.c[2] != 16 || s.d[-2] != 18 || s.d[-1] != 19 || s.d[0] != 20; @@ -48,7 +48,7 @@ foo (S s) || omp_target_is_present (&s.h, d) || omp_target_is_present (&s.h[2], d))) abort (); - #pragma omp target data map(alloc: s.a, s.b, s.c[1:2], s.d[-2:3], s.e, s.f, s.g[1:2], s.h[2:3]) + #pragma omp target data map(alloc: s.a, s.b, s.c[1:2], s.d, s.d[-2:3], s.e, s.f, s.g[1:2], s.h, s.h[2:3]) { if (!omp_target_is_present (&s.a, d) || !omp_target_is_present (s.b, d) @@ -61,8 +61,8 @@ foo (S s) || !omp_target_is_present (&s.h, d) || !omp_target_is_present (&s.h[2], d)) abort (); - #pragma omp target update to(s.a, s.b, s.c[1:2], s.d[-2:3], s.e, s.f, s.g[1:2], s.h[2:3]) - #pragma omp target map(alloc: s.a, s.b, s.c[1:2], s.d[-2:3], s.e, s.f, s.g[1:2], s.h[2:3]) map(from: err) + #pragma omp target update to(s.a, s.b, s.c[1:2], s.d, s.d[-2:3], s.e, s.f, s.g[1:2], s.h, s.h[2:3]) + #pragma omp target map(alloc: s.a, s.b, s.c[1:2], s.d, s.d[-2:3], s.e, s.f, s.g[1:2], s.h, s.h[2:3]) map(from: err) { err = s.a != 50 || s.b[0] != 49 || s.b[1] != 48; err |= s.c[1] != 47 || s.c[2] != 46 || s.d[-2] != 45 || s.d[-1] != 44 || s.d[0] != 43; @@ -73,7 +73,7 @@ foo (S s) s.e = 25; s.f[0] = 26; s.f[1] = 27; s.g[1] = 28; s.g[2] = 29; s.h[2] = 30; s.h[3] = 31; s.h[4] = 32; } - #pragma omp target update from(s.a, s.b, s.c[1:2], s.d[-2:3], s.e, s.f, s.g[1:2], s.h[2:3]) + #pragma omp target update from(s.a, s.b, s.c[1:2], s.d, s.d[-2:3], s.e, s.f, s.g[1:2], s.h, s.h[2:3]) } if (sep && (omp_target_is_present (&s.a, d) @@ -97,7 +97,7 @@ foo (S s) s.c[1] = 36; s.c[2] = 37; s.d[-2] = 38; s.d[-1] = 39; s.d[0] = 40; s.e = 41; s.f[0] = 42; s.f[1] = 43; s.g[1] = 44; s.g[2] = 45; s.h[2] = 46; s.h[3] = 47; s.h[4] = 48; - #pragma omp target enter data map(alloc: s.a, s.b, s.c[1:2], s.d[-2:3], s.e, s.f, s.g[1:2], s.h[2:3]) + #pragma omp target enter data map(alloc: s.a, s.b, s.c[1:2], s.d, s.d[-2:3], s.e, s.f, s.g[1:2], s.h, s.h[2:3]) if (!omp_target_is_present (&s.a, d) || !omp_target_is_present (s.b, d) || !omp_target_is_present (&s.c[1], d) @@ -109,8 +109,8 @@ foo (S s) || !omp_target_is_present (&s.h, d) || !omp_target_is_present (&s.h[2], d)) abort (); - #pragma omp target enter data map(always, to: s.a, s.b, s.c[1:2], s.d[-2:3], s.e, s.f, s.g[1:2], s.h[2:3]) - #pragma omp target map(alloc: s.a, s.b, s.c[1:2], s.d[-2:3], s.e, s.f, s.g[1:2], s.h[2:3]) map(from: err) + #pragma omp target enter data map(always, to: s.a, s.b, s.c[1:2], s.d, s.d[-2:3], s.e, s.f, s.g[1:2], s.h, s.h[2:3]) + #pragma omp target map(alloc: s.a, s.b, s.c[1:2], s.d, s.d[-2:3], s.e, s.f, s.g[1:2], s.h, s.h[2:3]) map(from: err) { err = s.a != 33 || s.b[0] != 34 || s.b[1] != 35; err |= s.c[1] != 36 || s.c[2] != 37 || s.d[-2] != 38 || s.d[-1] != 39 || s.d[0] != 40; @@ -121,7 +121,7 @@ foo (S s) s.e = 31; s.f[0] = 40; s.f[1] = 39; s.g[1] = 38; s.g[2] = 37; s.h[2] = 36; s.h[3] = 35; s.h[4] = 34; } - #pragma omp target exit data map(always, from: s.a, s.b, s.c[1:2], s.d[-2:3], s.e, s.f, s.g[1:2], s.h[2:3]) + #pragma omp target exit data map(always, from: s.a, s.b, s.c[1:2], s.d, s.d[-2:3], s.e, s.f, s.g[1:2], s.h, s.h[2:3]) if (!omp_target_is_present (&s.a, d) || !omp_target_is_present (s.b, d) || !omp_target_is_present (&s.c[1], d) @@ -133,7 +133,7 @@ foo (S s) || !omp_target_is_present (&s.h, d) || !omp_target_is_present (&s.h[2], d)) abort (); - #pragma omp target exit data map(release: s.a, s.b, s.c[1:2], s.d[-2:3], s.e, s.f, s.g[1:2], s.h[2:3]) + #pragma omp target exit data map(release: s.a, s.b, s.c[1:2], s.d, s.d[-2:3], s.e, s.f, s.g[1:2], s.h, s.h[2:3]) if (sep && (omp_target_is_present (&s.a, d) || omp_target_is_present (s.b, d) diff --git a/libgomp/testsuite/libgomp.c++/target-16.C b/libgomp/testsuite/libgomp.c++/target-16.C index cd102d90594..b8be7cc922f 100644 --- a/libgomp/testsuite/libgomp.c++/target-16.C +++ b/libgomp/testsuite/libgomp.c++/target-16.C @@ -16,7 +16,7 @@ foo (S s) d = id; int err; - #pragma omp target map(tofrom: s.a, s.b, s.c[1:2], s.d[-2:3], s.e, s.f, s.g[1:2], s.h[2:3]) map(to: sep) map(from: err) + #pragma omp target map(tofrom: s.a, s.b, s.c[1:2], s.d, s.d[-2:3], s.e, s.f, s.g[1:2], s.h, s.h[2:3]) map(to: sep) map(from: err) { err = s.a != 11 || s.b[0] != 12 || s.b[1] != 13; err |= s.c[1] != 15 || s.c[2] != 16 || s.d[-2] != 18 || s.d[-1] != 19 || s.d[0] != 20; @@ -50,7 +50,7 @@ foo (S s) || omp_target_is_present (&s.h, d) || omp_target_is_present (&s.h[2], d))) abort (); - #pragma omp target data map(alloc: s.a, s.b, s.c[1:2], s.d[-2:3], s.e, s.f, s.g[1:2], s.h[2:3]) + #pragma omp target data map(alloc: s.a, s.b, s.c[1:2], s.d, s.d[-2:3], s.e, s.f, s.g[1:2], s.h, s.h[2:3]) { if (!omp_target_is_present (&s.a, d) || !omp_target_is_present (s.b, d) @@ -63,8 +63,8 @@ foo (S s) || !omp_target_is_present (&s.h, d) || !omp_target_is_present (&s.h[2], d)) abort (); - #pragma omp target update to(s.a, s.b, s.c[1:2], s.d[-2:3], s.e, s.f, s.g[1:2], s.h[2:3]) - #pragma omp target map(alloc: s.a, s.b, s.c[1:2], s.d[-2:3], s.e, s.f, s.g[1:2], s.h[2:3]) map(from: err) + #pragma omp target update to(s.a, s.b, s.c[1:2], s.d, s.d[-2:3], s.e, s.f, s.g[1:2], s.h, s.h[2:3]) + #pragma omp target map(alloc: s.a, s.b, s.c[1:2], s.d, s.d[-2:3], s.e, s.f, s.g[1:2], s.h, s.h[2:3]) map(from: err) { err = s.a != 50 || s.b[0] != 49 || s.b[1] != 48; err |= s.c[1] != 47 || s.c[2] != 46 || s.d[-2] != 45 || s.d[-1] != 44 || s.d[0] != 43; @@ -75,7 +75,7 @@ foo (S s) s.e = 25; s.f[0] = 26; s.f[1] = 27; s.g[1] = 28; s.g[2] = 29; s.h[2] = 30; s.h[3] = 31; s.h[4] = 32; } - #pragma omp target update from(s.a, s.b, s.c[1:2], s.d[-2:3], s.e, s.f, s.g[1:2], s.h[2:3]) + #pragma omp target update from(s.a, s.b, s.c[1:2], s.d, s.d[-2:3], s.e, s.f, s.g[1:2], s.h, s.h[2:3]) } if (sep && (omp_target_is_present (&s.a, d) @@ -99,7 +99,7 @@ foo (S s) s.c[1] = 36; s.c[2] = 37; s.d[-2] = 38; s.d[-1] = 39; s.d[0] = 40; s.e = 41; s.f[0] = 42; s.f[1] = 43; s.g[1] = 44; s.g[2] = 45; s.h[2] = 46; s.h[3] = 47; s.h[4] = 48; - #pragma omp target enter data map(alloc: s.a, s.b, s.c[1:2], s.d[-2:3], s.e, s.f, s.g[1:2], s.h[2:3]) + #pragma omp target enter data map(alloc: s.a, s.b, s.c[1:2], s.d, s.d[-2:3], s.e, s.f, s.g[1:2], s.h, s.h[2:3]) if (!omp_target_is_present (&s.a, d) || !omp_target_is_present (s.b, d) || !omp_target_is_present (&s.c[1], d) @@ -111,8 +111,8 @@ foo (S s) || !omp_target_is_present (&s.h, d) || !omp_target_is_present (&s.h[2], d)) abort (); - #pragma omp target enter data map(always, to: s.a, s.b, s.c[1:2], s.d[-2:3], s.e, s.f, s.g[1:2], s.h[2:3]) - #pragma omp target map(alloc: s.a, s.b, s.c[1:2], s.d[-2:3], s.e, s.f, s.g[1:2], s.h[2:3]) map(from: err) + #pragma omp target enter data map(always, to: s.a, s.b, s.c[1:2], s.d, s.d[-2:3], s.e, s.f, s.g[1:2], s.h, s.h[2:3]) + #pragma omp target map(alloc: s.a, s.b, s.c[1:2], s.d, s.d[-2:3], s.e, s.f, s.g[1:2], s.h, s.h[2:3]) map(from: err) { err = s.a != 33 || s.b[0] != 34 || s.b[1] != 35; err |= s.c[1] != 36 || s.c[2] != 37 || s.d[-2] != 38 || s.d[-1] != 39 || s.d[0] != 40; @@ -123,7 +123,7 @@ foo (S s) s.e = 31; s.f[0] = 40; s.f[1] = 39; s.g[1] = 38; s.g[2] = 37; s.h[2] = 36; s.h[3] = 35; s.h[4] = 34; } - #pragma omp target exit data map(always, from: s.a, s.b, s.c[1:2], s.d[-2:3], s.e, s.f, s.g[1:2], s.h[2:3]) + #pragma omp target exit data map(always, from: s.a, s.b, s.c[1:2], s.d, s.d[-2:3], s.e, s.f, s.g[1:2], s.h, s.h[2:3]) if (!omp_target_is_present (&s.a, d) || !omp_target_is_present (s.b, d) || !omp_target_is_present (&s.c[1], d) @@ -135,7 +135,7 @@ foo (S s) || !omp_target_is_present (&s.h, d) || !omp_target_is_present (&s.h[2], d)) abort (); - #pragma omp target exit data map(release: s.a, s.b, s.c[1:2], s.d[-2:3], s.e, s.f, s.g[1:2], s.h[2:3]) + #pragma omp target exit data map(release: s.a, s.b, s.c[1:2], s.d, s.d[-2:3], s.e, s.f, s.g[1:2], s.h, s.h[2:3]) if (sep && (omp_target_is_present (&s.a, d) || omp_target_is_present (s.b, d) diff --git a/libgomp/testsuite/libgomp.c++/target-17.C b/libgomp/testsuite/libgomp.c++/target-17.C index d81ff19a411..f97476aafc4 100644 --- a/libgomp/testsuite/libgomp.c++/target-17.C +++ b/libgomp/testsuite/libgomp.c++/target-17.C @@ -16,7 +16,7 @@ foo (S s) d = id; int err; - #pragma omp target map(tofrom: s.a, s.b, s.c[1:2], s.d[-2:3], s.e, s.f, s.g[1:2], s.h[2:3]) map(to: sep) map(from: err) + #pragma omp target map(tofrom: s.a, s.b, s.c[1:2], s.d, s.d[-2:3], s.e, s.f, s.g[1:2], s.h, s.h[2:3]) map(to: sep) map(from: err) { err = s.a != 11 || s.b[0] != 12 || s.b[1] != 13; err |= s.c[1] != 15 || s.c[2] != 16 || s.d[-2] != 18 || s.d[-1] != 19 || s.d[0] != 20; @@ -50,7 +50,7 @@ foo (S s) || omp_target_is_present (&s.h, d) || omp_target_is_present (&s.h[2], d))) abort (); - #pragma omp target data map(alloc: s.a, s.b, s.c[1:2], s.d[-2:3], s.e, s.f, s.g[1:2], s.h[2:3]) + #pragma omp target data map(alloc: s.a, s.b, s.c[1:2], s.d, s.d[-2:3], s.e, s.f, s.g[1:2], s.h, s.h[2:3]) { if (!omp_target_is_present (&s.a, d) || !omp_target_is_present (s.b, d) @@ -63,8 +63,8 @@ foo (S s) || !omp_target_is_present (&s.h, d) || !omp_target_is_present (&s.h[2], d)) abort (); - #pragma omp target update to(s.a, s.b, s.c[1:2], s.d[-2:3], s.e, s.f, s.g[1:2], s.h[2:3]) - #pragma omp target map(alloc: s.a, s.b, s.c[1:2], s.d[-2:3], s.e, s.f, s.g[1:2], s.h[2:3]) map(from: err) + #pragma omp target update to(s.a, s.b, s.c[1:2], s.d, s.d[-2:3], s.e, s.f, s.g[1:2], s.h, s.h[2:3]) + #pragma omp target map(alloc: s.a, s.b, s.c[1:2], s.d, s.d[-2:3], s.e, s.f, s.g[1:2], s.h, s.h[2:3]) map(from: err) { err = s.a != 50 || s.b[0] != 49 || s.b[1] != 48; err |= s.c[1] != 47 || s.c[2] != 46 || s.d[-2] != 45 || s.d[-1] != 44 || s.d[0] != 43; @@ -75,7 +75,7 @@ foo (S s) s.e = 25; s.f[0] = 26; s.f[1] = 27; s.g[1] = 28; s.g[2] = 29; s.h[2] = 30; s.h[3] = 31; s.h[4] = 32; } - #pragma omp target update from(s.a, s.b, s.c[1:2], s.d[-2:3], s.e, s.f, s.g[1:2], s.h[2:3]) + #pragma omp target update from(s.a, s.b, s.c[1:2], s.d, s.d[-2:3], s.e, s.f, s.g[1:2], s.h, s.h[2:3]) } if (sep && (omp_target_is_present (&s.a, d) @@ -99,7 +99,7 @@ foo (S s) s.c[1] = 36; s.c[2] = 37; s.d[-2] = 38; s.d[-1] = 39; s.d[0] = 40; s.e = 41; s.f[0] = 42; s.f[1] = 43; s.g[1] = 44; s.g[2] = 45; s.h[2] = 46; s.h[3] = 47; s.h[4] = 48; - #pragma omp target enter data map(alloc: s.a, s.b, s.c[1:2], s.d[-2:3], s.e, s.f, s.g[1:2], s.h[2:3]) + #pragma omp target enter data map(alloc: s.a, s.b, s.c[1:2], s.d, s.d[-2:3], s.e, s.f, s.g[1:2], s.h, s.h[2:3]) if (!omp_target_is_present (&s.a, d) || !omp_target_is_present (s.b, d) || !omp_target_is_present (&s.c[1], d) @@ -111,8 +111,8 @@ foo (S s) || !omp_target_is_present (&s.h, d) || !omp_target_is_present (&s.h[2], d)) abort (); - #pragma omp target enter data map(always, to: s.a, s.b, s.c[1:2], s.d[-2:3], s.e, s.f, s.g[1:2], s.h[2:3]) - #pragma omp target map(alloc: s.a, s.b, s.c[1:2], s.d[-2:3], s.e, s.f, s.g[1:2], s.h[2:3]) map(from: err) + #pragma omp target enter data map(always, to: s.a, s.b, s.c[1:2], s.d, s.d[-2:3], s.e, s.f, s.g[1:2], s.h, s.h[2:3]) + #pragma omp target map(alloc: s.a, s.b, s.c[1:2], s.d, s.d[-2:3], s.e, s.f, s.g[1:2], s.h, s.h[2:3]) map(from: err) { err = s.a != 33 || s.b[0] != 34 || s.b[1] != 35; err |= s.c[1] != 36 || s.c[2] != 37 || s.d[-2] != 38 || s.d[-1] != 39 || s.d[0] != 40; @@ -123,7 +123,7 @@ foo (S s) s.e = 31; s.f[0] = 40; s.f[1] = 39; s.g[1] = 38; s.g[2] = 37; s.h[2] = 36; s.h[3] = 35; s.h[4] = 34; } - #pragma omp target exit data map(always, from: s.a, s.b, s.c[1:2], s.d[-2:3], s.e, s.f, s.g[1:2], s.h[2:3]) + #pragma omp target exit data map(always, from: s.a, s.b, s.c[1:2], s.d, s.d[-2:3], s.e, s.f, s.g[1:2], s.h, s.h[2:3]) if (!omp_target_is_present (&s.a, d) || !omp_target_is_present (s.b, d) || !omp_target_is_present (&s.c[1], d) @@ -135,7 +135,7 @@ foo (S s) || !omp_target_is_present (&s.h, d) || !omp_target_is_present (&s.h[2], d)) abort (); - #pragma omp target exit data map(release: s.a, s.b, s.c[1:2], s.d[-2:3], s.e, s.f, s.g[1:2], s.h[2:3]) + #pragma omp target exit data map(release: s.a, s.b, s.c[1:2], s.d, s.d[-2:3], s.e, s.f, s.g[1:2], s.h, s.h[2:3]) if (sep && (omp_target_is_present (&s.a, d) || omp_target_is_present (s.b, d) diff --git a/libgomp/testsuite/libgomp.c++/target-21.C b/libgomp/testsuite/libgomp.c++/target-21.C index 21a2f299bbb..da17b5745de 100644 --- a/libgomp/testsuite/libgomp.c++/target-21.C +++ b/libgomp/testsuite/libgomp.c++/target-21.C @@ -7,7 +7,7 @@ void foo (S s) { int err; - #pragma omp target map (s.x[0:N], s.y[0:N]) map (s.t.t[16:3]) map (from: err) + #pragma omp target map (s.x[0:N], s.y, s.y[0:N]) map (s.t.t[16:3]) map (from: err) { err = s.x[2] != 28 || s.y[2] != 37 || s.t.t[17] != 81; s.x[2]++; @@ -38,7 +38,7 @@ void foo2 (S &s) { int err; - #pragma omp target map (s.x[N:10], s.y[N:10]) map (from: err) map (s.t.t[N+16:N+3]) + #pragma omp target map (s.x[N:10], s.y, s.y[N:10]) map (from: err) map (s.t.t[N+16:N+3]) { err = s.x[2] != 30 || s.y[2] != 38 || s.t.t[17] != 81; s.x[2]++; @@ -69,7 +69,7 @@ void foo3 (U s) { int err; - #pragma omp target map (s.x[0:10], s.y[0:10]) map (from: err) map (s.t.t[16:3]) + #pragma omp target map (s.x[0:10], s.y, s.y[0:10]) map (from: err) map (s.t.t[16:3]) { err = s.x[2] != 32 || s.y[2] != 39 || s.t.t[17] != 82; s.x[2]++; @@ -100,7 +100,7 @@ void foo4 (U &s) { int err; - #pragma omp target map (s.x[0:10], s.y[0:10]) map (from: err) map (s.t.t[16:3]) + #pragma omp target map (s.x[0:10], s.y, s.y[0:10]) map (from: err) map (s.t.t[16:3]) { err = s.x[2] != 34 || s.y[2] != 40 || s.t.t[17] != 82; s.x[2]++; diff --git a/libgomp/testsuite/libgomp.c++/target-23.C b/libgomp/testsuite/libgomp.c++/target-23.C index d4f9ff3e983..63d343624b0 100644 --- a/libgomp/testsuite/libgomp.c++/target-23.C +++ b/libgomp/testsuite/libgomp.c++/target-23.C @@ -16,13 +16,13 @@ main (void) s->data[i] = 0; #pragma omp target enter data map(to: s) - #pragma omp target enter data map(to: s->data[:SZ]) + #pragma omp target enter data map(to: s->data, s->data[:SZ]) #pragma omp target { for (int i = 0; i < SZ; i++) s->data[i] = i; } - #pragma omp target exit data map(from: s->data[:SZ]) + #pragma omp target exit data map(from: s->data, s->data[:SZ]) #pragma omp target exit data map(from: s) for (int i = 0; i < SZ; i++) diff --git a/libgomp/testsuite/libgomp.c-c++-common/target-implicit-map-2.c b/libgomp/testsuite/libgomp.c-c++-common/target-implicit-map-2.c new file mode 100644 index 00000000000..974a9786c3f --- /dev/null +++ b/libgomp/testsuite/libgomp.c-c++-common/target-implicit-map-2.c @@ -0,0 +1,46 @@ +#include + +#define N 10 + +struct S +{ + int a, b; + int *ptr; + int c, d; +}; + +int +main (void) +{ + struct S a; + a.ptr = (int *) malloc (sizeof (int) * N); + + for (int i = 0; i < N; i++) + a.ptr[i] = 0; + + #pragma omp target enter data map(to: a.ptr, a.ptr[:N]) + + #pragma omp target + for (int i = 0; i < N; i++) + a.ptr[i] += 1; + + #pragma omp target update from(a.ptr[:N]) + + for (int i = 0; i < N; i++) + if (a.ptr[i] != 1) + abort (); + + #pragma omp target map(a.ptr[:N]) + for (int i = 0; i < N; i++) + a.ptr[i] += 1; + + #pragma omp target update from(a.ptr[:N]) + + for (int i = 0; i < N; i++) + if (a.ptr[i] != 2) + abort (); + + #pragma omp target exit data map(from:a.ptr, a.ptr[:N]) + + return 0; +} diff --git a/libgomp/testsuite/libgomp.c/target-23.c b/libgomp/testsuite/libgomp.c/target-23.c index fb1532a07b2..d56b13acf82 100644 --- a/libgomp/testsuite/libgomp.c/target-23.c +++ b/libgomp/testsuite/libgomp.c/target-23.c @@ -8,7 +8,7 @@ main () int u[10] = { 0, 1, 2, 3, 4, 5, 6, 7, 8, 9 }, err = 0; struct S s = { 9, u + 3, { 10, 11, 12, 13, 14 } }; int *v = u + 4; - #pragma omp target enter data map (to: s.s, s.u[0:5]) map (alloc: s.v[1:3]) + #pragma omp target enter data map (to: s.s, s.u, s.u[0:5]) map (alloc: s.v[1:3]) s.s++; u[3]++; s.v[1]++; diff --git a/libgomp/testsuite/libgomp.c/target-29.c b/libgomp/testsuite/libgomp.c/target-29.c index e5095a1b6b8..4a286649811 100644 --- a/libgomp/testsuite/libgomp.c/target-29.c +++ b/libgomp/testsuite/libgomp.c/target-29.c @@ -14,7 +14,7 @@ foo (struct S s) d = id; int err; - #pragma omp target map(tofrom: s.a, s.b, s.c[1:2], s.d[-2:3]) map(to: sep) map(from: err) + #pragma omp target map(tofrom: s.a, s.b, s.c[1:2], s.d, s.d[-2:3]) map(to: sep) map(from: err) { err = s.a != 11 || s.b[0] != 12 || s.b[1] != 13; err |= s.c[1] != 15 || s.c[2] != 16 || s.d[-2] != 18 || s.d[-1] != 19 || s.d[0] != 20; @@ -35,7 +35,7 @@ foo (struct S s) || omp_target_is_present (s.d, d) || omp_target_is_present (&s.d[-2], d))) abort (); - #pragma omp target data map(alloc: s.a, s.b, s.c[1:2], s.d[-2:3]) + #pragma omp target data map(alloc: s.a, s.b, s.c[1:2], s.d, s.d[-2:3]) { if (!omp_target_is_present (&s.a, d) || !omp_target_is_present (s.b, d) @@ -43,15 +43,15 @@ foo (struct S s) || !omp_target_is_present (s.d, d) || !omp_target_is_present (&s.d[-2], d)) abort (); - #pragma omp target update to(s.a, s.b, s.c[1:2], s.d[-2:3]) - #pragma omp target map(alloc: s.a, s.b, s.c[1:2], s.d[-2:3]) map(from: err) + #pragma omp target update to(s.a, s.b, s.c[1:2], s.d, s.d[-2:3]) + #pragma omp target map(alloc: s.a, s.b, s.c[1:2], s.d, s.d[-2:3]) map(from: err) { err = s.a != 50 || s.b[0] != 49 || s.b[1] != 48; err |= s.c[1] != 47 || s.c[2] != 46 || s.d[-2] != 45 || s.d[-1] != 44 || s.d[0] != 43; s.a = 17; s.b[0] = 18; s.b[1] = 19; s.c[1] = 20; s.c[2] = 21; s.d[-2] = 22; s.d[-1] = 23; s.d[0] = 24; } - #pragma omp target update from(s.a, s.b, s.c[1:2], s.d[-2:3]) + #pragma omp target update from(s.a, s.b, s.c[1:2], s.d, s.d[-2:3]) } if (sep && (omp_target_is_present (&s.a, d) @@ -66,29 +66,29 @@ foo (struct S s) if (err) abort (); s.a = 33; s.b[0] = 34; s.b[1] = 35; s.c[1] = 36; s.c[2] = 37; s.d[-2] = 38; s.d[-1] = 39; s.d[0] = 40; - #pragma omp target enter data map(alloc: s.a, s.b, s.c[1:2], s.d[-2:3]) + #pragma omp target enter data map(alloc: s.a, s.b, s.c[1:2], s.d, s.d[-2:3]) if (!omp_target_is_present (&s.a, d) || !omp_target_is_present (s.b, d) || !omp_target_is_present (&s.c[1], d) || !omp_target_is_present (s.d, d) || !omp_target_is_present (&s.d[-2], d)) abort (); - #pragma omp target enter data map(always, to: s.a, s.b, s.c[1:2], s.d[-2:3]) - #pragma omp target map(alloc: s.a, s.b, s.c[1:2], s.d[-2:3]) map(from: err) + #pragma omp target enter data map(always, to: s.a, s.b, s.c[1:2], s.d, s.d[-2:3]) + #pragma omp target map(alloc: s.a, s.b, s.c[1:2], s.d, s.d[-2:3]) map(from: err) { err = s.a != 33 || s.b[0] != 34 || s.b[1] != 35; err |= s.c[1] != 36 || s.c[2] != 37 || s.d[-2] != 38 || s.d[-1] != 39 || s.d[0] != 40; s.a = 49; s.b[0] = 48; s.b[1] = 47; s.c[1] = 46; s.c[2] = 45; s.d[-2] = 44; s.d[-1] = 43; s.d[0] = 42; } - #pragma omp target exit data map(always, from: s.a, s.b, s.c[1:2], s.d[-2:3]) + #pragma omp target exit data map(always, from: s.a, s.b, s.c[1:2], s.d, s.d[-2:3]) if (!omp_target_is_present (&s.a, d) || !omp_target_is_present (s.b, d) || !omp_target_is_present (&s.c[1], d) || !omp_target_is_present (s.d, d) || !omp_target_is_present (&s.d[-2], d)) abort (); - #pragma omp target exit data map(release: s.a, s.b, s.c[1:2], s.d[-2:3]) + #pragma omp target exit data map(release: s.a, s.b, s.c[1:2], s.d, s.d[-2:3]) if (sep && (omp_target_is_present (&s.a, d) || omp_target_is_present (s.b, d)