From mboxrd@z Thu Jan 1 00:00:00 1970 Return-Path: Received: from esa4.mentor.iphmx.com (esa4.mentor.iphmx.com [68.232.137.252]) by sourceware.org (Postfix) with ESMTPS id 223D7385702D for ; Wed, 11 Aug 2021 16:58:53 +0000 (GMT) DMARC-Filter: OpenDMARC Filter v1.4.1 sourceware.org 223D7385702D 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: gadJB3vgsovLlJmKY/ofaYqQdxt1AfabpYw35wExC8MVPoCMf4B5azsR54x/nBUxvv9C+dtl+5 f8l+7vXBp0KZIidraXQmMu6zJKTkTYB/R1+v8kAwWyLiYH12PYDZ8YpBWOhO9jtHsfyzSyRI/H VWMpirFtTXFwZ1Ri892dD21KZLzOXSRP++PXnl3/fLrXXT92ENlWuRcIRkfe73qDCDgOgq6dCq 6gqN0gz6fVgRmj1A+q+skMKzYAWyzbij5q+nFMeaEkq4vUvZQOzX7aBRuyxMnfioanS8A4CDWZ NcgwXg7cAnCOD+LEmhOUJluL X-IronPort-AV: E=Sophos;i="5.84,313,1620720000"; d="scan'208";a="64742311" Received: from orw-gwy-02-in.mentorg.com ([192.94.38.167]) by esa4.mentor.iphmx.com with ESMTP; 11 Aug 2021 08:58:50 -0800 IronPort-SDR: SdR2OGPq3P+wLbfc546KEiOCsnVtw/oj9UgOQJk7baT5pqCnMOD/L4F07COcMehtnmCBBNWWMr qTmjJIrZHkP7+kmJ3z0Tn6n8KJGmTA1mHytCPsPqpRDOPNI2DotRTwnqgyTz0XMNqxsvM8KjlL g49TkecWzvJ6EVAJpJJcqQzKVMtC+UPSlp6nPmsu7yXJg3G9ETK1hCMVl2Qtn3+P8gWVqAINOR qK2L6sB7ezhwFa3L+sP0tV6jG2M2Z/UqPK0L/NVvL0QJsx8bNQ2ncunpPBF1tXTah9wgXoDTED +2s= From: Julian Brown To: CC: Chung-Lin Tang , Jakub Jelinek Subject: [PATCH 3/8] Remove array section base-pointer mapping semantics, and other front-end adjustments (mainline trunk) Date: Wed, 11 Aug 2021 09:58:26 -0700 Message-ID: <8b40bf41dc37fc1243cbb5a299e1f8afccbf2f0a.1628697740.git.julian@codesourcery.com> X-Mailer: git-send-email 2.29.2 In-Reply-To: References: MIME-Version: 1.0 Content-Transfer-Encoding: 8bit Content-Type: text/plain X-Originating-IP: [137.202.0.90] X-ClientProxiedBy: SVR-IES-MBX-04.mgc.mentorg.com (139.181.222.4) To SVR-IES-MBX-04.mgc.mentorg.com (139.181.222.4) X-Spam-Status: No, score=-11.8 required=5.0 tests=BAYES_00, GIT_PATCH_0, HEADER_FROM_DIFFERENT_DOMAINS, KAM_DMARC_STATUS, 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: , X-List-Received-Date: Wed, 11 Aug 2021 16:58:58 -0000 From: Chung-Lin Tang This is a version of a patch by Chung-Lin, merged to current mainline. Any errors introduced are my own! It was previously posted here: https://gcc.gnu.org/pipermail/gcc-patches/2021-May/571195.html Chung-Lin's description from the previous submission follows (edited a little for formatting). This is a version of this patch: https://gcc.gnu.org/pipermail/gcc-patches/2021-May/570075.html for mainline trunk. 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: 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: (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: 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. Tested on x86_64-linux with nvptx offloading with no regressions. This patch was merged and tested atop of the prior submitted patches: (a) https://gcc.gnu.org/pipermail/gcc-patches/2021-May/570886.html "[PATCH, OpenMP 5.0] Improve OpenMP target support for C++ (includes PR92120 v3)" (b) https://gcc.gnu.org/pipermail/gcc-patches/2021-May/570365.html "[PATCH, OpenMP 5.0] Implement relaxation of implicit map vs. existing device mappings (for mainline trunk)" so you might queued this one later than those for review. Thanks, Chung-Lin 2021-05-25 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. * g++.dg/gomp/target-lambda-1.C: Likewise. * g++.dg/gomp/target-this-3.C: Likewise. * g++.dg/gomp/target-this-4.C: Likewise. * 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. -------------- next part -------------- --- gcc/c/c-parser.c | 54 ++++- gcc/c/c-typeck.c | 88 +++++-- gcc/cp/parser.c | 56 ++++- gcc/cp/semantics.c | 143 ++++++++++-- gcc/fortran/trans-openmp.c | 3 + gcc/gimplify.c | 220 +++++++++++++++--- .../goacc/deep-copy-arrayofstruct.c | 5 +- .../c-c++-common/gomp/target-enter-data-1.c | 24 ++ .../c-c++-common/gomp/target-implicit-map-2.c | 52 +++++ libgomp/target.c | 106 +++++++-- libgomp/testsuite/libgomp.c++/target-11.C | 14 +- libgomp/testsuite/libgomp.c++/target-12.C | 2 +- libgomp/testsuite/libgomp.c++/target-15.C | 20 +- libgomp/testsuite/libgomp.c++/target-16.C | 20 +- libgomp/testsuite/libgomp.c++/target-17.C | 20 +- libgomp/testsuite/libgomp.c++/target-21.C | 8 +- libgomp/testsuite/libgomp.c++/target-23.C | 4 +- .../target-implicit-map-2.c | 46 ++++ libgomp/testsuite/libgomp.c/target-23.c | 2 +- libgomp/testsuite/libgomp.c/target-29.c | 20 +- 20 files changed, 754 insertions(+), 153 deletions(-) create mode 100644 gcc/testsuite/c-c++-common/gomp/target-enter-data-1.c create mode 100644 gcc/testsuite/c-c++-common/gomp/target-implicit-map-2.c create mode 100644 libgomp/testsuite/libgomp.c-c++-common/target-implicit-map-2.c diff --git a/gcc/c/c-parser.c b/gcc/c/c-parser.c index e8aaec75677..785e2ebaeb5 100644 --- a/gcc/c/c-parser.c +++ b/gcc/c/c-parser.c @@ -12930,6 +12930,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, @@ -12942,6 +12951,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) { @@ -13061,6 +13071,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))) @@ -13088,9 +13099,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)) @@ -13101,9 +13116,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 `:'. */ @@ -13132,8 +13151,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) @@ -16046,7 +16092,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: @@ -16055,7 +16101,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 773cd2f8703..0e756815383 100644 --- a/gcc/c/c-typeck.c +++ b/gcc/c/c-typeck.c @@ -13090,6 +13090,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 @@ -13111,10 +13123,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) { @@ -13403,15 +13419,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 @@ -14761,13 +14787,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)) @@ -14778,6 +14811,7 @@ c_finish_omp_clauses (tree clauses, enum c_omp_region_type ort) remove = true; break; } + if (bitmap_bit_p (&map_field_head, DECL_UID (t))) break; if (bitmap_bit_p (&map_head, DECL_UID (t))) @@ -14834,14 +14868,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_) { @@ -14877,7 +14929,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), @@ -14886,6 +14938,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; @@ -14957,7 +15018,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 d90408aa3a1..88d45fe3fda 100644 --- a/gcc/cp/parser.c +++ b/gcc/cp/parser.c @@ -35805,11 +35805,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) @@ -35890,6 +35901,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))) @@ -35913,14 +35925,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)); @@ -35929,7 +35946,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 `:'. */ @@ -35942,6 +35962,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)) { @@ -35960,8 +35982,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; @@ -39223,11 +39269,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 6778efae606..fb99f0e360f 100644 --- a/gcc/cp/semantics.c +++ b/gcc/cp/semantics.c @@ -4995,6 +4995,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 @@ -5019,10 +5031,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)) @@ -5321,15 +5337,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 @@ -5599,16 +5625,35 @@ 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_ACC && 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_ACC && 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 @@ -5632,8 +5677,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)))) { @@ -7741,15 +7788,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)) @@ -7760,6 +7814,7 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort) remove = true; break; } + if (bitmap_bit_p (&map_field_head, DECL_UID (t))) break; if (bitmap_bit_p (&map_head, DECL_UID (t))) @@ -7820,15 +7875,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_) @@ -7863,6 +7936,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; @@ -7959,7 +8050,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), @@ -8006,8 +8098,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: @@ -8036,7 +8133,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 ace4faf038a..9dc2b6fc6a5 100644 --- a/gcc/fortran/trans-openmp.c +++ b/gcc/fortran/trans-openmp.c @@ -2458,6 +2458,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 327bb815323..fb35d240b34 100644 --- a/gcc/gimplify.c +++ b/gcc/gimplify.c @@ -8481,7 +8481,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; @@ -8528,10 +8528,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; @@ -8540,6 +8541,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) @@ -8553,12 +8555,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 @@ -8638,21 +8650,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 @@ -8744,6 +8842,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; @@ -9213,6 +9312,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 @@ -9231,6 +9331,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); } @@ -9277,7 +9378,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 (); @@ -9321,7 +9424,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 @@ -9356,7 +9460,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) @@ -9369,13 +9481,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) @@ -9407,6 +9521,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) @@ -9446,9 +9586,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 @@ -9533,9 +9678,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) @@ -9623,16 +9770,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); } @@ -10404,6 +10556,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 410a0ffb7fa..adb415c1f6f 100644 --- a/libgomp/target.c +++ b/libgomp/target.c @@ -534,11 +534,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, 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 *), 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, cbuf); } gomp_increment_refcount (oldn, refcount_set); @@ -1955,16 +1974,45 @@ gomp_update (struct gomp_device_descr *devicep, size_t mapnum, void **hostaddrs, } - 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 (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 *), 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, - NULL); - if (GOMP_MAP_COPY_FROM_P (kind & typemask)) - gomp_copy_dev2host (devicep, NULL, hostaddr, devaddr, size); + if (GOMP_MAP_COPY_TO_P (kind & typemask)) + gomp_copy_host2dev (devicep, NULL, devaddr, hostaddr, size, + NULL); + if (GOMP_MAP_COPY_FROM_P (kind & typemask)) + gomp_copy_dev2host (devicep, NULL, hostaddr, devaddr, size); + } } } gomp_mutex_unlock (&devicep->lock); @@ -2824,11 +2872,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) -- 2.29.2