From patchwork Fri Oct 1 17:09:03 2021 Content-Type: text/plain; charset="utf-8" MIME-Version: 1.0 Content-Transfer-Encoding: 7bit X-Patchwork-Submitter: Julian Brown X-Patchwork-Id: 45701 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 7BDA7385AC1E for ; Fri, 1 Oct 2021 17:11:09 +0000 (GMT) X-Original-To: gcc-patches@gcc.gnu.org Delivered-To: gcc-patches@gcc.gnu.org Received: from esa4.mentor.iphmx.com (esa4.mentor.iphmx.com [68.232.137.252]) by sourceware.org (Postfix) with ESMTPS id 787DF3857424 for ; Fri, 1 Oct 2021 17:09:20 +0000 (GMT) DMARC-Filter: OpenDMARC Filter v1.4.1 sourceware.org 787DF3857424 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: ZugmAeCSTmNXRU1bALExZGyR//gNYcW6s1wFEGpL23ELf5WE7dmYIDLfJsEVh159cpOBWIBj0L K19vEeGAfiy0M3mcjZK43IruXwBHWFynzzi5iOUx9jqOmVUa+NP9Pm8evY1KZP3gQjMvAi3gdr CKBcoPv1AMmbV1etW92DiVjbPASU0LIU9K9uAhJlpyZyCtVHCt5vfNFGVqwr9PMiFqNs+hQXed OsrVZP6reCIme/vjlC1bfx5YLim3360wk+QbGyjZBpabhHx72XMbnEIKxZksXIHGM372UlWgQZ ZxtvKYJ2VPS5i2+ZGuWcHE/E X-IronPort-AV: E=Sophos;i="5.85,339,1624348800"; d="scan'208";a="66726629" Received: from orw-gwy-01-in.mentorg.com ([192.94.38.165]) by esa4.mentor.iphmx.com with ESMTP; 01 Oct 2021 09:09:19 -0800 IronPort-SDR: csWoJ4bM65tkqBO7vRbtB6YqFL+DTiiK+MBulMxShmkYuoKvXixOB+vRqhX+ge3NHWDiTHjNnr 1eoPRiaKkru+RqWUfvPgcSKe8X+MHc9Erv7jktWRYx1Y/QNft573ugMp2bgHElDUyBN/3ln4Rn mxbWUevvNMZ1qW8inWM/0+k0h1FM3piUN+kiE8TJMvXCtWElH9u4MVpHmdduJ2RsamRjRrRCKo pth4C6kTLPd6bh7yf4Va90XHaiQdWRSgL2SOV3mUHr/Q/D2P7ldigCp5uE84b4dUhsUKCA9J6u a8U= From: Julian Brown To: Subject: [PATCH 05/11] OpenMP/OpenACC: Hoist struct sibling list handling in gimplification Date: Fri, 1 Oct 2021 10:09:03 -0700 Message-ID: <00411b16d398cb33a52f4357a527adf4bca4a94d.1633106213.git.julian@codesourcery.com> X-Mailer: git-send-email 2.29.2 In-Reply-To: References: MIME-Version: 1.0 X-Originating-IP: [137.202.0.90] X-ClientProxiedBy: svr-ies-mbx-06.mgc.mentorg.com (139.181.222.6) 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, 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: , Cc: Jakub Jelinek , Thomas Schwinge Errors-To: gcc-patches-bounces+patchwork=sourceware.org@gcc.gnu.org Sender: "Gcc-patches" This patch lifts struct sibling-list handling out of the main loop in gimplify_scan_omp_clauses. The reasons for this are several: first, it means that we can subject created sibling list groups to topological sorting (see previous patch) so base-pointer data dependencies are handled correctly. Secondly, it means that in the first pass gathering up sibling lists from parsed OpenMP/OpenACC clauses, we don't need to worry about gimplifying: that means we can see struct bases & components we need to sort sibling lists properly, even when we're using a non-DECL_P struct base. Gimplification proper still happens Thirdly, because we use more than one pass through the clause list and gather appropriate data, we can tell if we're mapping a whole struct in a different node, and avoid building struct sibling lists for that struct appropriately. Fourthly, we can re-use the node grouping functions from the previous patch, and thus mostly avoid the "prev_list_p" handling in gimplify_scan_omp_clauses that tracks the first node in such groups at present. Some redundant code has been removed and code paths for OpenACC/OpenMP are now shared where appropriate, though OpenACC doesn't do the topological sorting of nodes (yet?). OK for mainline? Thanks, Julian 2021-09-29 Julian Brown gcc/ * gimplify.c (gimplify_omp_var_data): Remove GOVD_MAP_HAS_ATTACHMENTS. (extract_base_bit_offset): Remove OFFSETP parameter. (strip_components_and_deref): Extend with POINTER_PLUS_EXPR and COMPOUND_EXPR handling. (aggregate_base_p): Remove. (omp_group_last, omp_group_base): Add GOMP_MAP_STRUCT handling. (build_struct_group): Remove CTX, DECL, PD, COMPONENT_REF_P, FLAGS, STRUCT_SEEN_CLAUSE, PRE_P, CONT parameters. Replace PREV_LIST_P and C parameters with GRP_START_P and GRP_END. Add INNER. Update calls to extract_base_bit_offset. Remove gimplification of clauses for OpenMP. Rework inner struct handling for OpenACC. Don't use context's variables splay tree. (omp_build_struct_sibling_lists): New function, extracted from gimplify_scan_omp_clauses and refactored. (gimplify_scan_omp_clauses): Call above function to handle struct sibling lists. Remove STRUCT_MAP_TO_CLAUSE, STRUCT_SEEN_CLAUSE, STRUCT_DEREF_SET. Rework flag handling, adding decl for struct variables. (gimplify_adjust_omp_clauses_1): Remove GOVD_MAP_HAS_ATTACHMENTS handling, unused now. gcc/testsuite/ * g++.dg/goacc/member-array-acc.C: Update expected output. * g++.dg/gomp/target-3.C: Likewise. * g++.dg/gomp/target-lambda-1.C: Likewise. * g++.dg/gomp/target-this-2.C: Likewise. * g++.dg/gomp/target-this-4.C: Likewise. --- gcc/gimplify.c | 943 ++++++++---------- gcc/testsuite/g++.dg/goacc/member-array-acc.C | 2 +- gcc/testsuite/g++.dg/gomp/target-3.C | 4 +- gcc/testsuite/g++.dg/gomp/target-lambda-1.C | 2 +- gcc/testsuite/g++.dg/gomp/target-this-2.C | 2 +- gcc/testsuite/g++.dg/gomp/target-this-4.C | 4 +- 6 files changed, 410 insertions(+), 547 deletions(-) diff --git a/gcc/gimplify.c b/gcc/gimplify.c index c10a3e8842a..31e2e4d9fe7 100644 --- a/gcc/gimplify.c +++ b/gcc/gimplify.c @@ -125,10 +125,6 @@ enum gimplify_omp_var_data /* Flag for GOVD_REDUCTION: inscan seen in {in,ex}clusive clause. */ GOVD_REDUCTION_INSCAN = 0x2000000, - /* Flag for GOVD_MAP: (struct) vars that have pointer attachments for - fields. */ - GOVD_MAP_HAS_ATTACHMENTS = 0x4000000, - /* Flag for GOVD_FIRSTPRIVATE: OMP_CLAUSE_FIRSTPRIVATE_IMPLICIT. */ GOVD_FIRSTPRIVATE_IMPLICIT = 0x8000000, @@ -8642,7 +8638,7 @@ build_struct_comp_nodes (enum tree_code code, tree grp_start, tree grp_end, static tree extract_base_bit_offset (tree base, poly_int64 *bitposp, - poly_offset_int *poffsetp, tree *offsetp) + poly_offset_int *poffsetp) { tree offset; poly_int64 bitsize, bitpos; @@ -8670,7 +8666,6 @@ extract_base_bit_offset (tree base, poly_int64 *bitposp, *bitposp = bitpos; *poffsetp = poffset; - *offsetp = offset; return base; } @@ -8683,8 +8678,15 @@ strip_components_and_deref (tree expr) while (TREE_CODE (expr) == COMPONENT_REF || TREE_CODE (expr) == INDIRECT_REF || (TREE_CODE (expr) == MEM_REF - && integer_zerop (TREE_OPERAND (expr, 1)))) - expr = TREE_OPERAND (expr, 0); + && integer_zerop (TREE_OPERAND (expr, 1))) + || TREE_CODE (expr) == POINTER_PLUS_EXPR + || TREE_CODE (expr) == COMPOUND_EXPR) + if (TREE_CODE (expr) == COMPOUND_EXPR) + expr = TREE_OPERAND (expr, 1); + else + expr = TREE_OPERAND (expr, 0); + + STRIP_NOPS (expr); return expr; } @@ -8700,34 +8702,6 @@ strip_indirections (tree expr) return expr; } -/* Return TRUE if EXPR is something we will use as the base of an aggregate - access, either: - - - a DECL_P. - - a struct component with no indirection ("a.b.c"). - - a struct component with indirection ("a->b->c"). -*/ - -static bool -aggregate_base_p (tree expr) -{ - while (TREE_CODE (expr) == COMPONENT_REF - && (DECL_P (TREE_OPERAND (expr, 0)) - || (TREE_CODE (TREE_OPERAND (expr, 0)) == COMPONENT_REF))) - expr = TREE_OPERAND (expr, 0); - - if (DECL_P (expr)) - return true; - - if (TREE_CODE (expr) == COMPONENT_REF - && (TREE_CODE (TREE_OPERAND (expr, 0)) == INDIRECT_REF - || (TREE_CODE (TREE_OPERAND (expr, 0)) == MEM_REF - && integer_zerop (TREE_OPERAND (TREE_OPERAND (expr, 0), 1))))) - return true; - - return false; -} - enum omp_tsort_mark { UNVISITED, TEMPORARY, @@ -8956,6 +8930,18 @@ omp_group_last (tree *start_p) || OMP_CLAUSE_MAP_KIND (nc) == GOMP_MAP_DETACH)) grp_last_p = &OMP_CLAUSE_CHAIN (c); break; + + case GOMP_MAP_STRUCT: + { + unsigned HOST_WIDE_INT num_mappings + = tree_to_uhwi (OMP_CLAUSE_SIZE (c)); + if (OMP_CLAUSE_MAP_KIND (nc) == GOMP_MAP_FIRSTPRIVATE_POINTER + || OMP_CLAUSE_MAP_KIND (nc) == GOMP_MAP_FIRSTPRIVATE_REFERENCE) + grp_last_p = &OMP_CLAUSE_CHAIN (*grp_last_p); + for (unsigned i = 0; i < num_mappings; i++) + grp_last_p = &OMP_CLAUSE_CHAIN (*grp_last_p); + } + break; } return grp_last_p; @@ -9089,6 +9075,21 @@ omp_group_base (omp_mapping_group *grp, unsigned int *chained, internal_error ("unexpected mapping node"); return error_mark_node; + case GOMP_MAP_STRUCT: + { + unsigned HOST_WIDE_INT num_mappings + = tree_to_uhwi (OMP_CLAUSE_SIZE (node)); + node = OMP_CLAUSE_CHAIN (node); + if (OMP_CLAUSE_MAP_KIND (node) == GOMP_MAP_FIRSTPRIVATE_POINTER + || OMP_CLAUSE_MAP_KIND (node) == GOMP_MAP_FIRSTPRIVATE_REFERENCE) + { + *firstprivate = OMP_CLAUSE_DECL (node); + node = OMP_CLAUSE_CHAIN (node); + } + *chained = num_mappings; + return node; + } + case GOMP_MAP_FORCE_DEVICEPTR: case GOMP_MAP_DEVICE_RESIDENT: case GOMP_MAP_LINK: @@ -9751,21 +9752,16 @@ move_concat_nodes_after (tree first_new, tree *last_new_tail, tree *first_ptr, next node. PREV_LIST_P and LIST_P may be modified by the function when a list rearrangement has taken place. */ -static tree -build_struct_group (struct gimplify_omp_ctx *ctx, - enum omp_region_type region_type, enum tree_code code, - tree decl, tree *pd, bool component_ref_p, - unsigned int *flags, tree c, +static tree * +build_struct_group (enum omp_region_type region_type, enum tree_code code, hash_map *&struct_map_to_clause, - hash_map *&struct_seen_clause, - tree *&prev_list_p, tree *&list_p, gimple_seq *pre_p, - bool *cont) + tree *grp_start_p, tree grp_end, tree *inner) { poly_offset_int coffset; poly_int64 cbitpos; - tree tree_coffset; - tree ocd = OMP_CLAUSE_DECL (c); + tree ocd = OMP_CLAUSE_DECL (grp_end); bool openmp = !(region_type & ORT_ACC); + tree *continue_at = NULL; while (TREE_CODE (ocd) == ARRAY_REF) ocd = TREE_OPERAND (ocd, 0); @@ -9773,90 +9769,31 @@ build_struct_group (struct gimplify_omp_ctx *ctx, if (TREE_CODE (ocd) == INDIRECT_REF) ocd = TREE_OPERAND (ocd, 0); - tree base = extract_base_bit_offset (ocd, &cbitpos, &coffset, &tree_coffset); - tree sbase; + tree base = extract_base_bit_offset (ocd, &cbitpos, &coffset); - if (openmp) - { - if (TREE_CODE (base) == INDIRECT_REF - && TREE_CODE (TREE_TYPE (TREE_OPERAND (base, 0))) == REFERENCE_TYPE) - sbase = strip_indirections (base); - else - sbase = base; - } - else - { - sbase = strip_indirections (base); + bool ptr = (OMP_CLAUSE_MAP_KIND (grp_end) == GOMP_MAP_ALWAYS_POINTER); + bool attach_detach = ((OMP_CLAUSE_MAP_KIND (grp_end) + == GOMP_MAP_ATTACH_DETACH) + || (OMP_CLAUSE_MAP_KIND (grp_end) + == GOMP_MAP_ATTACH_ZERO_LENGTH_ARRAY_SECTION)); + bool attach = (OMP_CLAUSE_MAP_KIND (grp_end) == GOMP_MAP_ATTACH + || OMP_CLAUSE_MAP_KIND (grp_end) == GOMP_MAP_DETACH); - STRIP_NOPS (sbase); - } - - bool do_map_struct = (sbase == decl && !tree_coffset); - - /* Here, DECL is usually a DECL_P, unless we have chained indirect member - accesses, e.g. mystruct->a->b. In that case it'll be the "mystruct->a" - part. */ - splay_tree_node n - = (DECL_P (decl) - ? splay_tree_lookup (ctx->variables, (splay_tree_key) decl) - : NULL); - bool ptr = (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ALWAYS_POINTER); - bool attach_detach = (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ATTACH_DETACH); - bool attach = (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ATTACH - || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_DETACH); - bool has_attachments = false; - - /* For OpenACC, pointers in structs should trigger an attach action. */ - if (attach_detach - && ((region_type & (ORT_ACC | ORT_TARGET | ORT_TARGET_DATA)) - || code == OMP_TARGET_ENTER_DATA - || code == OMP_TARGET_EXIT_DATA)) - { - /* Turn a GOMP_MAP_ATTACH_DETACH clause into a GOMP_MAP_ATTACH or - GOMP_MAP_DETACH clause after we have detected a case that needs a - GOMP_MAP_STRUCT mapping added. */ - gomp_map_kind k - = ((code == OACC_EXIT_DATA || code == OMP_TARGET_EXIT_DATA) - ? GOMP_MAP_DETACH : GOMP_MAP_ATTACH); - 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) - return NULL_TREE; - - /* Nor for attach_detach for OpenMP. */ + /* FIXME: If we're not mapping the base pointer in some other clause on this + directive, I think we want to create ALLOC/RELEASE here -- i.e. not + early-exit. */ if (openmp && 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); - } + return NULL; - return NULL_TREE; - } - - if ((DECL_P (decl) && (n == NULL || (n->value & GOVD_MAP) == 0)) - || (!DECL_P (decl) - && (!struct_map_to_clause - || struct_map_to_clause->get (decl) == NULL))) + if (!struct_map_to_clause || struct_map_to_clause->get (base) == NULL) { - tree l = build_omp_clause (OMP_CLAUSE_LOCATION (c), OMP_CLAUSE_MAP); + tree l = build_omp_clause (OMP_CLAUSE_LOCATION (grp_end), OMP_CLAUSE_MAP); gomp_map_kind k = attach ? GOMP_MAP_FORCE_PRESENT : GOMP_MAP_STRUCT; OMP_CLAUSE_SET_MAP_KIND (l, k); OMP_CLAUSE_DECL (l) = unshare_expr (base); - if (openmp - && !DECL_P (OMP_CLAUSE_DECL (l)) - && (gimplify_expr (&OMP_CLAUSE_DECL (l), pre_p, NULL, - is_gimple_lvalue, fb_lvalue) == GS_ERROR)) - return error_mark_node; + OMP_CLAUSE_SIZE (l) = (!attach ? size_int (1) : (DECL_P (OMP_CLAUSE_DECL (l)) @@ -9864,19 +9801,17 @@ build_struct_group (struct gimplify_omp_ctx *ctx, : TYPE_SIZE_UNIT (TREE_TYPE (OMP_CLAUSE_DECL (l))))); if (struct_map_to_clause == NULL) struct_map_to_clause = new hash_map; - struct_map_to_clause->put (decl, l); + struct_map_to_clause->put (base, l); if (ptr || attach_detach) { tree extra_node; tree alloc_node - = build_struct_comp_nodes (code, *prev_list_p, c, &extra_node); + = build_struct_comp_nodes (code, *grp_start_p, grp_end, + &extra_node); OMP_CLAUSE_CHAIN (l) = alloc_node; - tree **sc = (struct_seen_clause - ? struct_seen_clause->get (decl) - : NULL); - tree *insert_node_pos = sc ? *sc : prev_list_p; + tree *insert_node_pos = grp_start_p; if (extra_node) { @@ -9887,131 +9822,89 @@ build_struct_group (struct gimplify_omp_ctx *ctx, OMP_CLAUSE_CHAIN (alloc_node) = *insert_node_pos; *insert_node_pos = l; - prev_list_p = NULL; } else - list_p = insert_node_after (l, list_p); - - bool base_ref - = (TREE_CODE (base) == INDIRECT_REF - && ((TREE_CODE (TREE_TYPE (TREE_OPERAND (base, 0))) - == REFERENCE_TYPE) - || ((TREE_CODE (TREE_OPERAND (base, 0)) == INDIRECT_REF) - && (TREE_CODE (TREE_TYPE (TREE_OPERAND - (TREE_OPERAND (base, 0), 0))) - == REFERENCE_TYPE)))); - bool base_ind = ((TREE_CODE (base) == INDIRECT_REF - || (TREE_CODE (base) == MEM_REF - && integer_zerop (TREE_OPERAND (base, 1)))) - && (TREE_CODE (TREE_TYPE (TREE_OPERAND (base, 0))) - == POINTER_TYPE)); - - /* Handle pointers to structs and references to structs: these cases - have an additional GOMP_MAP_FIRSTPRIVATE_{REFERENCE,POINTER} node - inserted after the GOMP_MAP_STRUCT node. References to pointers - use GOMP_MAP_FIRSTPRIVATE_REFERENCE. */ - if (base_ref && code == OMP_TARGET) { - tree c2 = build_omp_clause (OMP_CLAUSE_LOCATION (c), - OMP_CLAUSE_MAP); - enum gomp_map_kind mkind - = GOMP_MAP_FIRSTPRIVATE_REFERENCE; - OMP_CLAUSE_SET_MAP_KIND (c2, mkind); - OMP_CLAUSE_DECL (c2) = decl; - OMP_CLAUSE_SIZE (c2) = size_zero_node; - OMP_CLAUSE_CHAIN (c2) = OMP_CLAUSE_CHAIN (l); - OMP_CLAUSE_CHAIN (l) = c2; + gcc_assert (*grp_start_p == grp_end); + grp_start_p = insert_node_after (l, grp_start_p); } - else if (!openmp - && (base_ind || base_ref) - && (region_type & ORT_TARGET)) + + tree noind = strip_indirections (base); + + if (!openmp + && (region_type & ORT_TARGET) + && TREE_CODE (noind) == COMPONENT_REF) { - tree c2 = build_omp_clause (OMP_CLAUSE_LOCATION (c), OMP_CLAUSE_MAP); + /* The base for this component access is a struct component access + itself. Insert a node to be processed on the next iteration of + our caller's loop, which will subsequently be turned into a new, + inner GOMP_MAP_STRUCT mapping. + + We need to do this else the non-DECL_P base won't be + rewritten correctly in the offloaded region. */ + tree c2 = build_omp_clause (OMP_CLAUSE_LOCATION (grp_end), + OMP_CLAUSE_MAP); + OMP_CLAUSE_SET_MAP_KIND (c2, GOMP_MAP_FORCE_PRESENT); + OMP_CLAUSE_DECL (c2) = unshare_expr (noind); + OMP_CLAUSE_SIZE (c2) = TYPE_SIZE_UNIT (TREE_TYPE (noind)); + *inner = c2; + return NULL; + } + + tree sdecl = strip_components_and_deref (base); + + if (POINTER_TYPE_P (TREE_TYPE (sdecl)) && (region_type & ORT_TARGET)) + { + tree c2 = build_omp_clause (OMP_CLAUSE_LOCATION (grp_end), + OMP_CLAUSE_MAP); + bool base_ref + = (TREE_CODE (base) == INDIRECT_REF + && ((TREE_CODE (TREE_TYPE (TREE_OPERAND (base, 0))) + == REFERENCE_TYPE) + || ((TREE_CODE (TREE_OPERAND (base, 0)) + == INDIRECT_REF) + && (TREE_CODE (TREE_TYPE (TREE_OPERAND + (TREE_OPERAND (base, 0), 0))) + == REFERENCE_TYPE)))); enum gomp_map_kind mkind = base_ref ? GOMP_MAP_FIRSTPRIVATE_REFERENCE : GOMP_MAP_FIRSTPRIVATE_POINTER; OMP_CLAUSE_SET_MAP_KIND (c2, mkind); - OMP_CLAUSE_SIZE (c2) = size_zero_node; - tree sdecl = strip_components_and_deref (decl); - if (DECL_P (decl) - && (POINTER_TYPE_P (TREE_TYPE (sdecl)) - || TREE_CODE (TREE_TYPE (sdecl)) == REFERENCE_TYPE)) - { - /* Insert after struct node. */ - OMP_CLAUSE_CHAIN (c2) = OMP_CLAUSE_CHAIN (l); - OMP_CLAUSE_DECL (c2) = decl; - OMP_CLAUSE_CHAIN (l) = c2; - } - else - { - /* If the ultimate base for this component access is not a - pointer or reference, that means it is a struct component - access itself. Insert a node to be processed on the next - iteration of our caller's loop, which will subsequently be - turned into a new GOMP_MAP_STRUCT mapping itself. - - We need to do this else the non-DECL_P base won't be - rewritten correctly in the offloaded region. */ - tree c2 = build_omp_clause (OMP_CLAUSE_LOCATION (c), - OMP_CLAUSE_MAP); - OMP_CLAUSE_SET_MAP_KIND (c2, GOMP_MAP_FORCE_PRESENT); - OMP_CLAUSE_DECL (c2) = unshare_expr (decl); - OMP_CLAUSE_SIZE (c2) = (DECL_P (decl) - ? DECL_SIZE_UNIT (decl) - : TYPE_SIZE_UNIT (TREE_TYPE (decl))); - tree *next_node = &OMP_CLAUSE_CHAIN (*list_p); - OMP_CLAUSE_CHAIN (c2) = *next_node; - *next_node = c2; - return NULL_TREE; - } - } - *flags = GOVD_MAP | GOVD_EXPLICIT; - if (GOMP_MAP_ALWAYS_P (OMP_CLAUSE_MAP_KIND (c)) || ptr || attach_detach) - *flags |= GOVD_SEEN; - if (has_attachments) - *flags |= GOVD_MAP_HAS_ATTACHMENTS; - - /* If this is a *pointer-to-struct expression, make sure a - firstprivate map of the base-pointer exists. */ - if (openmp - && component_ref_p - && ((TREE_CODE (decl) == MEM_REF - && integer_zerop (TREE_OPERAND (decl, 1))) - || INDIRECT_REF_P (decl)) - && DECL_P (TREE_OPERAND (decl, 0)) - && !splay_tree_lookup (ctx->variables, - ((splay_tree_key) TREE_OPERAND (decl, 0)))) - { - decl = TREE_OPERAND (decl, 0); - tree c2 = build_omp_clause (OMP_CLAUSE_LOCATION (c), OMP_CLAUSE_MAP); - enum gomp_map_kind mkind = GOMP_MAP_FIRSTPRIVATE_POINTER; - OMP_CLAUSE_SET_MAP_KIND (c2, mkind); - OMP_CLAUSE_DECL (c2) = decl; - OMP_CLAUSE_SIZE (c2) = size_zero_node; - OMP_CLAUSE_CHAIN (c2) = OMP_CLAUSE_CHAIN (c); - OMP_CLAUSE_CHAIN (c) = c2; + OMP_CLAUSE_DECL (c2) = sdecl; + tree baddr = build_fold_addr_expr (base); + baddr = fold_convert_loc (OMP_CLAUSE_LOCATION (grp_end), + ptrdiff_type_node, baddr); + /* This isn't going to be good enough when we add support for more + complicated lvalue expressions. FIXME. */ + if (TREE_CODE (TREE_TYPE (sdecl)) == REFERENCE_TYPE + && TREE_CODE (TREE_TYPE (TREE_TYPE (sdecl))) == POINTER_TYPE) + sdecl = build_simple_mem_ref (sdecl); + tree decladdr = fold_convert_loc (OMP_CLAUSE_LOCATION (grp_end), + ptrdiff_type_node, sdecl); + OMP_CLAUSE_SIZE (c2) + = fold_build2_loc (OMP_CLAUSE_LOCATION (grp_end), MINUS_EXPR, + ptrdiff_type_node, baddr, decladdr); + /* Insert after struct node. */ + OMP_CLAUSE_CHAIN (c2) = OMP_CLAUSE_CHAIN (l); + OMP_CLAUSE_CHAIN (l) = c2; } - return decl; + return NULL; } else if (struct_map_to_clause) { - tree *osc = struct_map_to_clause->get (decl); + tree *osc = struct_map_to_clause->get (base); tree *sc = NULL, *scp = NULL; - if (n && (GOMP_MAP_ALWAYS_P (OMP_CLAUSE_MAP_KIND (c)) - || ptr - || attach_detach)) - n->value |= GOVD_SEEN; sc = &OMP_CLAUSE_CHAIN (*osc); /* The struct mapping might be immediately followed by a FIRSTPRIVATE_POINTER and/or FIRSTPRIVATE_REFERENCE -- if it's an indirect access or a reference, or both. (This added node is removed in omp-low.c after it has been processed there.) */ - if (*sc != c + if (*sc != grp_end && (OMP_CLAUSE_MAP_KIND (*sc) == GOMP_MAP_FIRSTPRIVATE_POINTER || OMP_CLAUSE_MAP_KIND (*sc) == GOMP_MAP_FIRSTPRIVATE_REFERENCE)) sc = &OMP_CLAUSE_CHAIN (*sc); - for (; *sc != c; sc = &OMP_CLAUSE_CHAIN (*sc)) - if ((ptr || attach_detach) && sc == prev_list_p) + for (; *sc != grp_end; sc = &OMP_CLAUSE_CHAIN (*sc)) + if ((ptr || attach_detach) && sc == grp_start_p) break; else if (TREE_CODE (OMP_CLAUSE_DECL (*sc)) != COMPONENT_REF && TREE_CODE (OMP_CLAUSE_DECL (*sc)) != INDIRECT_REF @@ -10022,7 +9915,6 @@ build_struct_group (struct gimplify_omp_ctx *ctx, tree sc_decl = OMP_CLAUSE_DECL (*sc); poly_offset_int offset; poly_int64 bitpos; - tree tree_offset; if (TREE_CODE (sc_decl) == ARRAY_REF) { @@ -10038,8 +9930,7 @@ build_struct_group (struct gimplify_omp_ctx *ctx, == REFERENCE_TYPE)) sc_decl = TREE_OPERAND (sc_decl, 0); - tree base2 = extract_base_bit_offset (sc_decl, &bitpos, &offset, - &tree_offset); + tree base2 = extract_base_bit_offset (sc_decl, &bitpos, &offset); if (!base2 || !operand_equal_p (base2, base, 0)) break; if (scp) @@ -10049,7 +9940,7 @@ build_struct_group (struct gimplify_omp_ctx *ctx, /* This duplicate checking code is currently only enabled for OpenACC. */ tree d1 = OMP_CLAUSE_DECL (*sc); - tree d2 = OMP_CLAUSE_DECL (c); + tree d2 = OMP_CLAUSE_DECL (grp_end); while (TREE_CODE (d1) == ARRAY_REF) d1 = TREE_OPERAND (d1, 0); while (TREE_CODE (d2) == ARRAY_REF) @@ -10069,10 +9960,10 @@ build_struct_group (struct gimplify_omp_ctx *ctx, break; if (d1 == d2) { - error_at (OMP_CLAUSE_LOCATION (c), + error_at (OMP_CLAUSE_LOCATION (grp_end), "%qE appears more than once in map clauses", - OMP_CLAUSE_DECL (c)); - return error_mark_node; + OMP_CLAUSE_DECL (grp_end)); + return NULL; } } if (maybe_lt (coffset, offset) @@ -10092,15 +9983,15 @@ build_struct_group (struct gimplify_omp_ctx *ctx, if (ptr || attach_detach) { tree cl = NULL_TREE, extra_node; - tree alloc_node = build_struct_comp_nodes (code, *prev_list_p, c, - &extra_node); + tree alloc_node = build_struct_comp_nodes (code, *grp_start_p, + grp_end, &extra_node); tree *tail_chain = NULL; /* Here, we have: - c : the currently-processed node. - prev_list_p : pointer to the first node in a pointer mapping group - up to and including C. + grp_end : the last (or only) node in this group. + grp_start_p : pointer to the first node in a pointer mapping group + up to and including GRP_END. sc : pointer to the chain for the end of the struct component list. scp : pointer to the chain for the sorted position at which we @@ -10111,7 +10002,7 @@ build_struct_group (struct gimplify_omp_ctx *ctx, (the end of the struct component list). extra_node : a newly-synthesized node for an additional indirect pointer mapping or a Fortran pointer set, if needed. - cl : first node to prepend before prev_list_p. + cl : first node to prepend before grp_start_p. tail_chain : pointer to chain of last prepended node. The general idea is we move the nodes for this struct mapping @@ -10147,32 +10038,180 @@ build_struct_group (struct gimplify_omp_ctx *ctx, tail_chain = &OMP_CLAUSE_CHAIN (alloc_node); } - tree *continue_at - = cl ? move_concat_nodes_after (cl, tail_chain, prev_list_p, c, sc) - : move_nodes_after (prev_list_p, c, sc); - - prev_list_p = NULL; - - if (continue_at) - { - list_p = continue_at; - *cont = true; - } + continue_at + = cl ? move_concat_nodes_after (cl, tail_chain, grp_start_p, + grp_end, sc) + : move_nodes_after (grp_start_p, grp_end, sc); } - else if (*sc != c) + else if (*sc != grp_end) { - if (openmp - && (gimplify_expr (pd, pre_p, NULL, is_gimple_lvalue, fb_lvalue) - == GS_ERROR)) - return error_mark_node; + gcc_assert (*grp_start_p == grp_end); + + /* We are moving the current node back to a previous struct node: + the node that used to point to the current node will now point to + the next node. */ + continue_at = grp_start_p; /* In the non-pointer case, the mapping clause itself is moved into the correct position in the struct component list, which in this case is just SC. */ - move_node_after (c, list_p, sc); - *cont = true; + move_node_after (grp_end, grp_start_p, sc); } } - return NULL_TREE; + return continue_at; +} + +/* Scan through GROUPS, and create sorted structure sibling lists without + gimplifying. */ + +static bool +omp_build_struct_sibling_lists (enum tree_code code, + enum omp_region_type region_type, + vec *groups, + hash_map + *grpmap) +{ + unsigned i; + omp_mapping_group *grp; + hash_map *struct_map_to_clause = NULL; + bool success = true; + tree *new_next = NULL; + tree *tail = &OMP_CLAUSE_CHAIN ((*groups)[groups->length () - 1].grp_end); + + FOR_EACH_VEC_ELT (*groups, i, grp) + { + tree c = grp->grp_end; + tree decl = OMP_CLAUSE_DECL (c); + tree *grp_start_p = new_next ? new_next : grp->grp_start; + tree grp_end = grp->grp_end; + + new_next = NULL; + + if (DECL_P (decl)) + continue; + + if (OMP_CLAUSE_CHAIN (*grp_start_p) + && OMP_CLAUSE_CHAIN (*grp_start_p) != grp_end) + { + /* Don't process an array descriptor that isn't inside a derived type + as a struct (the GOMP_MAP_POINTER following will have the form + "var.data", but such mappings are handled specially). */ + tree grpmid = OMP_CLAUSE_CHAIN (*grp_start_p); + if (OMP_CLAUSE_CODE (grpmid) == OMP_CLAUSE_MAP + && OMP_CLAUSE_MAP_KIND (grpmid) == GOMP_MAP_TO_PSET + && DECL_P (OMP_CLAUSE_DECL (grpmid))) + continue; + } + + tree d = decl; + if (TREE_CODE (d) == ARRAY_REF) + { + while (TREE_CODE (d) == ARRAY_REF) + d = TREE_OPERAND (d, 0); + if (TREE_CODE (d) == COMPONENT_REF + && TREE_CODE (TREE_TYPE (d)) == ARRAY_TYPE) + decl = d; + } + if (d == decl + && TREE_CODE (decl) == INDIRECT_REF + && TREE_CODE (TREE_OPERAND (decl, 0)) == COMPONENT_REF + && (TREE_CODE (TREE_TYPE (TREE_OPERAND (decl, 0))) + == REFERENCE_TYPE) + && (OMP_CLAUSE_MAP_KIND (c) + != GOMP_MAP_POINTER_TO_ZERO_LENGTH_ARRAY_SECTION)) + decl = TREE_OPERAND (decl, 0); + + STRIP_NOPS (decl); + + if (TREE_CODE (decl) != COMPONENT_REF) + continue; + + omp_mapping_group **wholestruct = NULL; + tree wsdecl = omp_containing_struct (OMP_CLAUSE_DECL (c)); + + if (!(region_type & ORT_ACC) && wsdecl != OMP_CLAUSE_DECL (c)) + { + wholestruct = grpmap->get (wsdecl); + if (!wholestruct + && TREE_CODE (wsdecl) == MEM_REF + && integer_zerop (TREE_OPERAND (wsdecl, 1))) + { + tree deref = TREE_OPERAND (wsdecl, 0); + deref = build1 (INDIRECT_REF, TREE_TYPE (wsdecl), deref); + wholestruct = grpmap->get (deref); + } + } + + if (wholestruct) + { + + if (*grp_start_p == grp_end) + { + /* Remove the whole of this mapping -- redundant. */ + new_next = grp_start_p; + *grp_start_p = OMP_CLAUSE_CHAIN (grp_end); + } + + continue; + } + + if (OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_TO_PSET + && OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_ATTACH + && OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_DETACH + && code != OACC_UPDATE + && code != OMP_TARGET_UPDATE) + { + if (error_operand_p (decl)) + { + success = false; + goto error_out; + } + + tree stype = TREE_TYPE (decl); + if (TREE_CODE (stype) == REFERENCE_TYPE) + stype = TREE_TYPE (stype); + if (TYPE_SIZE_UNIT (stype) == NULL + || TREE_CODE (TYPE_SIZE_UNIT (stype)) != INTEGER_CST) + { + error_at (OMP_CLAUSE_LOCATION (c), + "mapping field %qE of variable length " + "structure", OMP_CLAUSE_DECL (c)); + success = false; + goto error_out; + } + + tree inner = NULL_TREE; + + new_next = build_struct_group (region_type, code, + struct_map_to_clause, grp_start_p, + grp_end, &inner); + + if (inner) + { + if (new_next && *new_next == NULL_TREE) + *new_next = inner; + else + *tail = inner; + + OMP_CLAUSE_CHAIN (inner) = NULL_TREE; + + omp_mapping_group newgrp; + newgrp.grp_start = new_next ? new_next : tail; + newgrp.grp_end = inner; + newgrp.mark = UNVISITED; + newgrp.sibling = NULL; + newgrp.next = NULL; + groups->safe_push (newgrp); + + tail = &OMP_CLAUSE_CHAIN (inner); + } + } + } + +error_out: + if (struct_map_to_clause) + delete struct_map_to_clause; + + return success; } /* Scan the OMP clauses in *LIST_P, installing mappings into a new @@ -10185,9 +10224,6 @@ 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; int nowait = -1; @@ -10219,14 +10255,10 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p, break; } - /* Topological sorting may fail if we have duplicate nodes, which - we should have detected and shown an error for already. Skip - sorting in that case. */ - if (!seen_error () - && (code == OMP_TARGET - || code == OMP_TARGET_DATA - || code == OMP_TARGET_ENTER_DATA - || code == OMP_TARGET_EXIT_DATA)) + if (code == OMP_TARGET + || code == OMP_TARGET_DATA + || code == OMP_TARGET_ENTER_DATA + || code == OMP_TARGET_EXIT_DATA) { vec *groups; groups = omp_gather_mapping_groups (list_p); @@ -10234,12 +10266,46 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p, { hash_map *grpmap; grpmap = omp_index_mapping_groups (groups); - omp_mapping_group *outlist - = omp_tsort_mapping_groups (groups, grpmap); - outlist = omp_segregate_mapping_groups (outlist); - list_p = omp_reorder_mapping_groups (groups, outlist, list_p); + + omp_build_struct_sibling_lists (code, region_type, groups, grpmap); + + omp_mapping_group *outlist = NULL; + + /* Topological sorting may fail if we have duplicate nodes, which + we should have detected and shown an error for already. Skip + sorting in that case. */ + if (seen_error ()) + goto failure; + delete grpmap; delete groups; + + /* Rebuild now we have struct sibling lists. */ + groups = omp_gather_mapping_groups (list_p); + grpmap = omp_index_mapping_groups (groups); + + outlist = omp_tsort_mapping_groups (groups, grpmap); + outlist = omp_segregate_mapping_groups (outlist); + list_p = omp_reorder_mapping_groups (groups, outlist, list_p); + + failure: + delete grpmap; + delete groups; + } + } + else if (region_type & ORT_ACC) + { + vec *groups; + groups = omp_gather_mapping_groups (list_p); + if (groups) + { + hash_map *grpmap; + grpmap = omp_index_mapping_groups (groups); + + omp_build_struct_sibling_lists (code, region_type, groups, grpmap); + + delete groups; + delete grpmap; } } @@ -10648,6 +10714,28 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p, GOVD_FIRSTPRIVATE | GOVD_SEEN); } + if (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_STRUCT) + { + tree base = strip_components_and_deref (decl); + if (DECL_P (base)) + { + decl = base; + splay_tree_node n + = splay_tree_lookup (ctx->variables, + (splay_tree_key) decl); + if (seen_error () + && n + && (n->value & (GOVD_MAP | GOVD_FIRSTPRIVATE)) != 0) + { + remove = true; + break; + } + flags = GOVD_MAP | GOVD_EXPLICIT; + + goto do_add_decl; + } + } + if (TREE_CODE (decl) == TARGET_EXPR) { if (gimplify_expr (&OMP_CLAUSE_DECL (c), pre_p, NULL, @@ -10678,143 +10766,6 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p, pd = &TREE_OPERAND (decl, 0); decl = TREE_OPERAND (decl, 0); } - 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) && TREE_CODE (decl) == COMPONENT_REF) - { - /* Strip off component refs from RHS of e.g. "a->b->c.d.e" - (which would leave "a->b" in that case). This is intended - to be equivalent to the base finding done by - get_inner_reference. */ - while (TREE_CODE (decl) == COMPONENT_REF - && (DECL_P (TREE_OPERAND (decl, 0)) - || (TREE_CODE (TREE_OPERAND (decl, 0)) - == COMPONENT_REF))) - decl = TREE_OPERAND (decl, 0); - - if (TREE_CODE (decl) == COMPONENT_REF) - decl = TREE_OPERAND (decl, 0); - - /* Strip off RHS from "a->b". */ - if ((TREE_CODE (decl) == INDIRECT_REF - || (TREE_CODE (decl) == MEM_REF - && integer_zerop (TREE_OPERAND (decl, 1)))) - && (TREE_CODE (TREE_TYPE (TREE_OPERAND (decl, 0))) - == POINTER_TYPE)) - decl = TREE_OPERAND (decl, 0); - - /* Strip off RHS from "a_ref.b" (where a_ref is - reference-typed). */ - if (TREE_CODE (decl) == INDIRECT_REF - && DECL_P (TREE_OPERAND (decl, 0)) - && (TREE_CODE (TREE_TYPE (TREE_OPERAND (decl, 0))) - == REFERENCE_TYPE)) - decl = TREE_OPERAND (decl, 0); - - STRIP_NOPS (decl); - } - else if ((region_type & (ORT_TARGET | ORT_TARGET_DATA)) != 0 - && TREE_CODE (*pd) == COMPONENT_REF - && OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ATTACH_DETACH - && code != OACC_UPDATE) - { - while (TREE_CODE (decl) == COMPONENT_REF) - { - decl = TREE_OPERAND (decl, 0); - component_ref_p = true; - if (((TREE_CODE (decl) == MEM_REF - && integer_zerop (TREE_OPERAND (decl, 1))) - || INDIRECT_REF_P (decl)) - && (TREE_CODE (TREE_TYPE (TREE_OPERAND (decl, 0))) - == POINTER_TYPE)) - { - indir_p = true; - indir_base = decl; - decl = TREE_OPERAND (decl, 0); - STRIP_NOPS (decl); - } - if (TREE_CODE (decl) == INDIRECT_REF - && DECL_P (TREE_OPERAND (decl, 0)) - && (TREE_CODE (TREE_TYPE (TREE_OPERAND (decl, 0))) - == REFERENCE_TYPE)) - { - decl_ref = decl; - decl = TREE_OPERAND (decl, 0); - } - } - } - else if (TREE_CODE (decl) == COMPONENT_REF - && (OMP_CLAUSE_MAP_KIND (c) - != GOMP_MAP_ATTACH_ZERO_LENGTH_ARRAY_SECTION)) - { - component_ref_p = true; - while (TREE_CODE (decl) == COMPONENT_REF) - decl = TREE_OPERAND (decl, 0); - if (TREE_CODE (decl) == INDIRECT_REF - && DECL_P (TREE_OPERAND (decl, 0)) - && (TREE_CODE (TREE_TYPE (TREE_OPERAND (decl, 0))) - == REFERENCE_TYPE)) - decl = TREE_OPERAND (decl, 0); - } - if (decl != orig_decl && DECL_P (decl) && indir_p) - { - gomp_map_kind k - = ((code == OACC_EXIT_DATA || code == OMP_TARGET_EXIT_DATA) - ? GOMP_MAP_DETACH : GOMP_MAP_ATTACH); - /* We have a dereference of a struct member. Make this an - attach/detach operation, and ensure the base pointer is - mapped as a FIRSTPRIVATE_POINTER. */ - OMP_CLAUSE_SET_MAP_KIND (c, k); - flags = GOVD_MAP | GOVD_SEEN | GOVD_EXPLICIT; - tree next_clause = OMP_CLAUSE_CHAIN (c); - if (k == GOMP_MAP_ATTACH - && code != OACC_ENTER_DATA - && code != OMP_TARGET_ENTER_DATA - && (!next_clause - || (OMP_CLAUSE_CODE (next_clause) != OMP_CLAUSE_MAP) - || (OMP_CLAUSE_MAP_KIND (next_clause) - != GOMP_MAP_POINTER) - || OMP_CLAUSE_DECL (next_clause) != decl) - && (!struct_deref_set - || !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 (); - /* As well as the attach, we also need a - FIRSTPRIVATE_POINTER clause to properly map the - pointer to the struct base. */ - tree c2 = build_omp_clause (OMP_CLAUSE_LOCATION (c), - OMP_CLAUSE_MAP); - OMP_CLAUSE_SET_MAP_KIND (c2, GOMP_MAP_ALLOC); - OMP_CLAUSE_MAP_MAYBE_ZERO_LENGTH_ARRAY_SECTION (c2) - = 1; - tree charptr_zero - = build_int_cst (build_pointer_type (char_type_node), - 0); - OMP_CLAUSE_DECL (c2) - = build2 (MEM_REF, char_type_node, - decl_ref ? decl_ref : decl, charptr_zero); - OMP_CLAUSE_SIZE (c2) = size_zero_node; - tree c3 = build_omp_clause (OMP_CLAUSE_LOCATION (c), - OMP_CLAUSE_MAP); - OMP_CLAUSE_SET_MAP_KIND (c3, - GOMP_MAP_FIRSTPRIVATE_POINTER); - OMP_CLAUSE_DECL (c3) = decl; - OMP_CLAUSE_SIZE (c3) = size_zero_node; - tree mapgrp = *prev_list_p; - *prev_list_p = c2; - OMP_CLAUSE_CHAIN (c3) = mapgrp; - OMP_CLAUSE_CHAIN (c2) = c3; - - struct_deref_set->add (decl); - } - goto do_add_decl; - } /* An "attach/detach" operation on an update directive should behave as a GOMP_MAP_ALWAYS_POINTER. Beware that unlike attach or detach map kinds, GOMP_MAP_ALWAYS_POINTER @@ -10822,91 +10773,49 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p, if (code == OACC_UPDATE && OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ATTACH_DETACH) OMP_CLAUSE_SET_MAP_KIND (c, GOMP_MAP_ALWAYS_POINTER); - if ((((region_type & ORT_ACC) && aggregate_base_p (decl)) - || (!(region_type & ORT_ACC) - && (DECL_P (decl) - || (component_ref_p - && (INDIRECT_REF_P (decl) - || 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 - && code != OACC_UPDATE - && code != OMP_TARGET_UPDATE) + + if (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ATTACH_DETACH) { - if (error_operand_p (decl)) + if (TREE_CODE (TREE_TYPE (OMP_CLAUSE_DECL (c))) + == ARRAY_TYPE) + remove = true; + else { - remove = true; - break; + 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); } - - tree stype = TREE_TYPE (decl); - if (TREE_CODE (stype) == REFERENCE_TYPE) - stype = TREE_TYPE (stype); - if (TYPE_SIZE_UNIT (stype) == NULL - || TREE_CODE (TYPE_SIZE_UNIT (stype)) != INTEGER_CST) - { - error_at (OMP_CLAUSE_LOCATION (c), - "mapping field %qE of variable length " - "structure", OMP_CLAUSE_DECL (c)); - remove = true; - break; - } - - if (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ALWAYS_POINTER - || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ATTACH_DETACH) - { - /* Error recovery. */ - if (prev_list_p == NULL) - { - remove = true; - break; - } - - /* 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) - { - remove = true; - break; - } - } - } - bool cont = false; - tree add_decl - = build_struct_group (ctx, region_type, code, decl, pd, - component_ref_p, &flags, c, - struct_map_to_clause, - struct_seen_clause, prev_list_p, - list_p, pre_p, &cont); - if (add_decl == error_mark_node) - { - remove = true; - break; - } - else if (add_decl && DECL_P (add_decl)) - { - decl = add_decl; - goto do_add_decl; - } - if (cont) - continue; } - else if (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ATTACH_DETACH) + + tree cref = decl; + + while (TREE_CODE (cref) == ARRAY_REF) + cref = TREE_OPERAND (cref, 0); + + if (TREE_CODE (cref) == INDIRECT_REF) + cref = TREE_OPERAND (cref, 0); + + if (TREE_CODE (cref) == COMPONENT_REF) { - 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); + tree base = cref; + while (base && !DECL_P (base)) + { + tree innerbase = omp_get_base_pointer (base); + if (!innerbase) + break; + base = innerbase; + } + if (base + && DECL_P (base) + && GOMP_MAP_ALWAYS_P (OMP_CLAUSE_MAP_KIND (c)) + && POINTER_TYPE_P (TREE_TYPE (base))) + { + splay_tree_node n + = splay_tree_lookup (ctx->variables, + (splay_tree_key) base); + n->value |= GOVD_SEEN; + } } if (code == OMP_TARGET && OMP_CLAUSE_MAP_IN_REDUCTION (c)) @@ -11024,24 +10933,6 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p, break; } - /* If this was of the form map(*pointer_to_struct), then the - 'pointer_to_struct' DECL should be considered deref'ed. */ - if ((OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ALLOC - || GOMP_MAP_COPY_TO_P (OMP_CLAUSE_MAP_KIND (c)) - || GOMP_MAP_COPY_FROM_P (OMP_CLAUSE_MAP_KIND (c))) - && INDIRECT_REF_P (orig_decl) - && DECL_P (TREE_OPERAND (orig_decl, 0)) - && TREE_CODE (TREE_TYPE (orig_decl)) == RECORD_TYPE) - { - tree ptr = TREE_OPERAND (orig_decl, 0); - if (!struct_deref_set || !struct_deref_set->contains (ptr)) - { - if (!struct_deref_set) - struct_deref_set = new hash_set (); - struct_deref_set->add (ptr); - } - } - if (!remove && OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_ALWAYS_POINTER && OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_ATTACH_DETACH @@ -11058,28 +10949,6 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p, break; } - else - { - /* DECL_P (decl) == true */ - tree *sc; - if (struct_map_to_clause - && (sc = struct_map_to_clause->get (decl)) != NULL - && OMP_CLAUSE_MAP_KIND (*sc) == GOMP_MAP_STRUCT - && decl == OMP_CLAUSE_DECL (*sc)) - { - /* We have found a map of the whole structure after a - leading GOMP_MAP_STRUCT has been created, so refill the - leading clause into a map of the whole structure - variable, and remove the current one. - TODO: we should be able to remove some maps of the - following structure element maps if they are of - compatible TO/FROM/ALLOC type. */ - OMP_CLAUSE_SET_MAP_KIND (*sc, OMP_CLAUSE_MAP_KIND (c)); - OMP_CLAUSE_SIZE (*sc) = unshare_expr (OMP_CLAUSE_SIZE (c)); - remove = true; - break; - } - } flags = GOVD_MAP | GOVD_EXPLICIT; if (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ALWAYS_TO || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ALWAYS_TOFROM) @@ -11721,12 +11590,6 @@ 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) - delete struct_deref_set; } /* Return true if DECL is a candidate for shared to firstprivate @@ -11875,8 +11738,6 @@ gimplify_adjust_omp_clauses_1 (splay_tree_node n, void *data) return 0; if ((flags & GOVD_SEEN) == 0) return 0; - if ((flags & GOVD_MAP_HAS_ATTACHMENTS) != 0) - return 0; if (flags & GOVD_DEBUG_PRIVATE) { gcc_assert ((flags & GOVD_DATA_SHARE_CLASS) == GOVD_SHARED); diff --git a/gcc/testsuite/g++.dg/goacc/member-array-acc.C b/gcc/testsuite/g++.dg/goacc/member-array-acc.C index e0c11570f5d..9993768ef20 100644 --- a/gcc/testsuite/g++.dg/goacc/member-array-acc.C +++ b/gcc/testsuite/g++.dg/goacc/member-array-acc.C @@ -10,4 +10,4 @@ struct Foo { }; int main() { Foo x; x.init(1024); } -/* { dg-final { scan-tree-dump {struct:\*\(struct Foo \*\) this \[len: 1\]\) map\(alloc:\(\(struct Foo \*\) this\)->a \[len: [0-9]+\]\) map\(alloc:\*_[0-9]+ \[len: _[0-9]+\]\) map\(attach:\(\(struct Foo \*\) this\)->a \[bias: 0\]\)} "gimple" } } */ +/* { dg-final { scan-tree-dump {struct:\*\(struct Foo \*\) this \[len: 1\]\) map\(alloc:this->a \[len: [0-9]+\]\) map\(alloc:\*_[0-9]+ \[len: _[0-9]+\]\) map\(attach:this->a \[bias: 0\]\)} "gimple" } } */ diff --git a/gcc/testsuite/g++.dg/gomp/target-3.C b/gcc/testsuite/g++.dg/gomp/target-3.C index f4d40ec8e4b..432f02614d8 100644 --- a/gcc/testsuite/g++.dg/gomp/target-3.C +++ b/gcc/testsuite/g++.dg/gomp/target-3.C @@ -33,4 +33,6 @@ T::bar (int x) template struct T<0>; -/* { dg-final { scan-tree-dump-times "map\\(struct:\\*this \\\[len: 2\\\]\\) map\\(alloc:this->a \\\[len: \[0-9\]+\\\]\\) map\\(alloc:this->b \\\[len: \[0-9\]+\\\]\\)" 4 "gimple" } } */ +/* { dg-final { scan-tree-dump-times "map\\(struct:\\*\\(struct S \\*\\) this \\\[len: 2\\\]\\) map\\(alloc:this->a \\\[len: \[0-9\]+\\\]\\) map\\(alloc:this->b \\\[len: \[0-9\]+\\\]\\)" 2 "gimple" } } */ + +/* { dg-final { scan-tree-dump-times "map\\(struct:\\*\\(struct T \\*\\) this \\\[len: 2\\\]\\) map\\(alloc:this->a \\\[len: \[0-9\]+\\\]\\) map\\(alloc:this->b \\\[len: \[0-9\]+\\\]\\)" 2 "gimple" } } */ diff --git a/gcc/testsuite/g++.dg/gomp/target-lambda-1.C b/gcc/testsuite/g++.dg/gomp/target-lambda-1.C index 150f286e312..bff7fa7c669 100644 --- a/gcc/testsuite/g++.dg/gomp/target-lambda-1.C +++ b/gcc/testsuite/g++.dg/gomp/target-lambda-1.C @@ -87,7 +87,7 @@ int main (void) return 0; } -/* { dg-final { scan-tree-dump {#pragma omp target num_teams.* firstprivate\(b\) map\(alloc:MEM.* \[len: 0\]\) map\(firstprivate:iptr \[pointer assign, bias: 0\]\) map\(alloc:MEM.* \[len: 0\]\) map\(firstprivate:this \[pointer assign, bias: 0\]\) map\(to:\*__closure \[len: [0-9]+\]\) map\(firstprivate:__closure \[pointer assign, bias: 0\]\) map\(struct:\*__closure \[len: 1\]\) map\(alloc:__closure->__this \[len: [0-9]+\]\) map\(tofrom:\*_[0-9]+ \[len: [0-9]+\]\) map\(always_pointer:__closure->__this \[pointer assign, bias: 0\]\) map\(from:mapped \[len: [0-9]+\]\) map\(alloc:\*_[0-9]+ \[len: 0\]\) map\(attach_zero_length_array_section:__closure->__iptr \[bias: 0\]\) map\(alloc:\*_[0-9]+ \[len: 0\]\) map\(attach_zero_length_array_section:_[0-9]+->ptr \[bias: 0\]\)} "gimple" } } */ +/* { dg-final { scan-tree-dump {#pragma omp target num_teams.* firstprivate\(b\) map\(alloc:MEM.* \[len: 0\]\) map\(firstprivate:iptr \[pointer assign, bias: 0\]\) map\(alloc:MEM.* \[len: 0\]\) map\(firstprivate:this \[pointer assign, bias: 0\]\) map\(to:\*__closure \[len: [0-9]+\]\) map\(firstprivate:__closure \[pointer assign, bias: 0\]\) map\(tofrom:\*_[0-9]+ \[len: [0-9]+\]\) map\(always_pointer:__closure->__this \[pointer assign, bias: 0\]\) map\(from:mapped \[len: [0-9]+\]\) map\(alloc:\*_[0-9]+ \[len: 0\]\) map\(attach_zero_length_array_section:__closure->__iptr \[bias: 0\]\) map\(alloc:\*_[0-9]+ \[len: 0\]\) map\(attach_zero_length_array_section:_[0-9]+->ptr \[bias: 0\]\)} "gimple" } } */ /* { dg-final { scan-tree-dump {#pragma omp target num_teams.* firstprivate\(end\) firstprivate\(begin\) map\(to:loop \[len: [0-9]+\]\) map\(alloc:\*_[0-9]+ \[len: 0\]\) map\(attach_zero_length_array_section:loop\.__data1 \[bias: 0\]\)} "gimple" } } */ diff --git a/gcc/testsuite/g++.dg/gomp/target-this-2.C b/gcc/testsuite/g++.dg/gomp/target-this-2.C index 679c85a54dd..cc08e7e8693 100644 --- a/gcc/testsuite/g++.dg/gomp/target-this-2.C +++ b/gcc/testsuite/g++.dg/gomp/target-this-2.C @@ -46,4 +46,4 @@ int main (void) return 0; } -/* { dg-final { scan-tree-dump {map\(to:\*__closure \[len: [0-9]+\]\) map\(firstprivate:__closure \[pointer assign, bias: 0\]\) map\(struct:\*__closure \[len: 1\]\) map\(alloc:__closure->__this \[len: [0-9]+\]\) map\(tofrom:\*_[0-9]+ \[len: [0-9]+\]\) map\(always_pointer:__closure->__this \[pointer assign, bias: 0\]\)} "gimple" } } */ +/* { dg-final { scan-tree-dump {map\(alloc:MEM\[\(char \*\)_[0-9]+\] \[len: [0-9]+\]\) map\(firstprivate:this \[pointer assign, bias: 0\]\) firstprivate\(m\) map\(to:\*__closure \[len: [0-9]+\]\) map\(firstprivate:__closure \[pointer assign, bias: 0\]\) map\(tofrom:\*_[0-9]+ \[len: [0-9]+\]\) map\(always_pointer:__closure->__this \[pointer assign, bias: 0\]\) map\(from:v \[len: [0-9]+\]\)} "gimple" } } */ diff --git a/gcc/testsuite/g++.dg/gomp/target-this-4.C b/gcc/testsuite/g++.dg/gomp/target-this-4.C index 8166f43ad42..9ade3cc0b2b 100644 --- a/gcc/testsuite/g++.dg/gomp/target-this-4.C +++ b/gcc/testsuite/g++.dg/gomp/target-this-4.C @@ -102,6 +102,6 @@ int main (void) return 0; } -/* { dg-final { scan-tree-dump {#pragma omp target num_teams.* firstprivate\(n\) map\(alloc:MEM.* \[len: 0\]\) map\(firstprivate:this \[pointer assign, bias: 0\]\) map\(to:\*__closure \[len: [0-9]+\]\) map\(firstprivate:__closure \[pointer assign, bias: 0\]\) map\(struct:\*__closure \[len: 1\]\) map\(alloc:__closure->__this \[len: [0-9]+\]\) map\(tofrom:\*_[0-9]+ \[len: [0-9]+\]\) map\(always_pointer:__closure->__this \[pointer assign, bias: 0\]\) map\(from:mapped \[len: 1\]\) map\(alloc:\*_[0-9]+ \[len: 0\]\) map\(attach_zero_length_array_section:_[0-9]+->ptr \[bias: 0\]\)} "gimple" } } */ +/* { dg-final { scan-tree-dump {#pragma omp target num_teams.* firstprivate\(n\) map\(alloc:MEM.* \[len: 0\]\) map\(firstprivate:this \[pointer assign, bias: 0\]\) map\(to:\*__closure \[len: [0-9]+\]\) map\(firstprivate:__closure \[pointer assign, bias: 0\]\) map\(tofrom:\*_[0-9]+ \[len: [0-9]+\]\) map\(always_pointer:__closure->__this \[pointer assign, bias: 0\]\) map\(from:mapped \[len: 1\]\) map\(alloc:\*_[0-9]+ \[len: 0\]\) map\(attach_zero_length_array_section:_[0-9]+->ptr \[bias: 0\]\)} "gimple" } } */ -/* { dg-final { scan-tree-dump {#pragma omp target num_teams.* firstprivate\(n\) map\(alloc:MEM.* \[len: 0\]\) map\(firstprivate:this \[pointer assign, bias: 0\]\) map\(to:\*__closure \[len: [0-9]+\]\) map\(firstprivate:__closure \[pointer assign, bias: 0\]\) map\(struct:\*__closure \[len: 1\]\) map\(alloc:__closure->__this \[len: [0-9]+\]\) map\(tofrom:\*_[0-9]+ \[len: [0-9]+\]\) map\(always_pointer:__closure->__this \[pointer assign, bias: 0\]\) map\(from:mapped \[len: [0-9]+\]\) map\(alloc:\*_[0-9]+ \[len: 0\]\) map\(alloc:\*_[0-9]+ \[pointer assign, zero-length array section, bias: 0\]\) map\(attach:_[0-9]+->refptr \[bias: 0\]\)} "gimple" } } */ +/* { dg-final { scan-tree-dump {#pragma omp target num_teams.* firstprivate\(n\) map\(alloc:MEM.* \[len: 0\]\) map\(firstprivate:this \[pointer assign, bias: 0\]\) map\(to:\*__closure \[len: [0-9]+\]\) map\(firstprivate:__closure \[pointer assign, bias: 0\]\) map\(tofrom:\*_[0-9]+ \[len: [0-9]+\]\) map\(always_pointer:__closure->__this \[pointer assign, bias: 0\]\) map\(from:mapped \[len: [0-9]+\]\) map\(alloc:\*_[0-9]+ \[len: 0\]\) map\(alloc:\*_[0-9]+ \[pointer assign, zero-length array section, bias: 0\]\) map\(attach:_[0-9]+->refptr \[bias: 0\]\)} "gimple" } } */