From patchwork Fri Aug 18 22:47:51 2023 Content-Type: text/plain; charset="utf-8" MIME-Version: 1.0 Content-Transfer-Encoding: 7bit X-Patchwork-Submitter: Julian Brown X-Patchwork-Id: 1823195 Return-Path: X-Original-To: incoming@patchwork.ozlabs.org Delivered-To: patchwork-incoming@legolas.ozlabs.org Authentication-Results: legolas.ozlabs.org; spf=pass (sender SPF authorized) smtp.mailfrom=gcc.gnu.org (client-ip=2620:52:3:1:0:246e:9693:128c; helo=server2.sourceware.org; envelope-from=gcc-patches-bounces+incoming=patchwork.ozlabs.org@gcc.gnu.org; receiver=patchwork.ozlabs.org) Received: from server2.sourceware.org (server2.sourceware.org [IPv6:2620:52:3:1:0:246e:9693:128c]) (using TLSv1.3 with cipher TLS_AES_256_GCM_SHA384 (256/256 bits) key-exchange X25519 server-signature ECDSA (secp384r1) server-digest SHA384) (No client certificate requested) by legolas.ozlabs.org (Postfix) with ESMTPS id 4RSHBC5LQ5z1yfT for ; Sat, 19 Aug 2023 08:49:39 +1000 (AEST) Received: from server2.sourceware.org (localhost [IPv6:::1]) by sourceware.org (Postfix) with ESMTP id BB241388B6B4 for ; Fri, 18 Aug 2023 22:49:36 +0000 (GMT) X-Original-To: gcc-patches@gcc.gnu.org Delivered-To: gcc-patches@gcc.gnu.org Received: from esa1.mentor.iphmx.com (esa1.mentor.iphmx.com [68.232.129.153]) by sourceware.org (Postfix) with ESMTPS id DE344382C134; Fri, 18 Aug 2023 22:49:12 +0000 (GMT) DMARC-Filter: OpenDMARC Filter v1.4.2 sourceware.org DE344382C134 Authentication-Results: sourceware.org; dmarc=none (p=none dis=none) header.from=codesourcery.com Authentication-Results: sourceware.org; spf=pass smtp.mailfrom=mentor.com X-IronPort-AV: E=Sophos;i="6.01,184,1684828800"; d="scan'208";a="16759180" Received: from orw-gwy-02-in.mentorg.com ([192.94.38.167]) by esa1.mentor.iphmx.com with ESMTP; 18 Aug 2023 14:49:11 -0800 IronPort-SDR: bd/fSDsZbnoQWlkegihQojnSafXUU/JWDE5GAGxtkxhaY5RrG5xtXHcjpKvXx79u5s7mRu+TCh x0x4uZI+ksSKf68st2pMoWOd84GMvB8bSdlh1qSeKRVHmzGPC9BrdgInfB1KnWXQ+7ZoUp/GHB NAE5N76gZpJ75PTVEXGVzC6NTdCBYDbzlh7QPI5gecHW6bgiLuP4EIO8jtPhbfR8K5YEgo4z8C Cqwcp2sspCmBslHCIStJHGA6Y18TefAdGu81bZTkuX1k7bo4xmRgq0CF4mA/C+nOoNczz6cPgx dC0= From: Julian Brown To: CC: , , Subject: [PATCH v7 5/5] OpenMP/OpenACC: Reorganise OMP map clause handling in gimplify.cc Date: Fri, 18 Aug 2023 15:47:51 -0700 Message-ID: X-Mailer: git-send-email 2.41.0 In-Reply-To: References: MIME-Version: 1.0 X-Originating-IP: [137.202.0.90] X-ClientProxiedBy: svr-ies-mbx-14.mgc.mentorg.com (139.181.222.14) To svr-ies-mbx-11.mgc.mentorg.com (139.181.222.11) 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.6 X-Spam-Checker-Version: SpamAssassin 3.4.6 (2021-04-09) on server2.sourceware.org X-BeenThere: gcc-patches@gcc.gnu.org X-Mailman-Version: 2.1.29 Precedence: list List-Id: Gcc-patches mailing list List-Unsubscribe: , List-Archive: List-Post: List-Help: List-Subscribe: , Errors-To: gcc-patches-bounces+incoming=patchwork.ozlabs.org@gcc.gnu.org Sender: "Gcc-patches" This patch has been separated out from the C++ "declare mapper" support patch. It contains just the gimplify.cc rearrangement work, mostly moving gimplification from gimplify_scan_omp_clauses to gimplify_adjust_omp_clauses for map clauses. The motivation for doing this was that we don't know if we need to instantiate mappers implicitly until the body of an offload region has been scanned, i.e. in gimplify_adjust_omp_clauses, but we also need the un-gimplified form of clauses to sort by base-pointer dependencies after mapper instantiation has taken place. The patch also reimplements the "present" clause sorting code to avoid another sorting pass on mapping nodes. This version of the patch is based on the version posted for og13, and additionally incorporates a follow-on fix for DECL_VALUE_EXPR handling in gimplify_adjust_omp_clauses: "OpenMP/OpenACC: Reorganise OMP map clause handling in gimplify.cc" https://gcc.gnu.org/pipermail/gcc-patches/2023-June/622223.html Parts of: "OpenMP: OpenMP 5.2 semantics for pointers with unmapped target" https://gcc.gnu.org/pipermail/gcc-patches/2023-June/623351.html 2023-08-18 Julian Brown gcc/ * gimplify.cc (omp_segregate_mapping_groups): Handle "present" groups. (gimplify_scan_omp_clauses): Use mapping group functionality to iterate through mapping nodes. Remove most gimplification of OMP_CLAUSE_MAP nodes from here, but still populate ctx->variables splay tree. (gimplify_adjust_omp_clauses): Move most gimplification of OMP_CLAUSE_MAP nodes here. gcc/testsuite/ * gfortran.dg/gomp/map-12.f90: Adjust scan output. --- gcc/gimplify.cc | 667 +++++++++++++--------- gcc/testsuite/gfortran.dg/gomp/map-12.f90 | 2 +- 2 files changed, 386 insertions(+), 283 deletions(-) diff --git a/gcc/gimplify.cc b/gcc/gimplify.cc index e682583054b0..1e32ad48b844 100644 --- a/gcc/gimplify.cc +++ b/gcc/gimplify.cc @@ -9804,10 +9804,15 @@ omp_tsort_mapping_groups (vec *groups, return outlist; } -/* Split INLIST into two parts, moving groups corresponding to - ALLOC/RELEASE/DELETE mappings to one list, and other mappings to another. - The former list is then appended to the latter. Each sub-list retains the - order of the original list. +/* Split INLIST into four parts: + + - "present" to/from groups + - "present" alloc groups + - other to/from groups + - other alloc/release/delete groups + + These sub-lists are then concatenated together to form the final list. + Each sub-list retains the order of the original list. Note that ATTACH nodes are later moved to the end of the list in gimplify_adjust_omp_clauses, for target regions. */ @@ -9815,7 +9820,9 @@ static omp_mapping_group * omp_segregate_mapping_groups (omp_mapping_group *inlist) { omp_mapping_group *ard_groups = NULL, *tf_groups = NULL; + omp_mapping_group *pa_groups = NULL, *ptf_groups = NULL; omp_mapping_group **ard_tail = &ard_groups, **tf_tail = &tf_groups; + omp_mapping_group **pa_tail = &pa_groups, **ptf_tail = &ptf_groups; for (omp_mapping_group *w = inlist; w;) { @@ -9834,6 +9841,20 @@ omp_segregate_mapping_groups (omp_mapping_group *inlist) ard_tail = &w->next; break; + case GOMP_MAP_PRESENT_ALLOC: + *pa_tail = w; + w->next = NULL; + pa_tail = &w->next; + break; + + case GOMP_MAP_PRESENT_FROM: + case GOMP_MAP_PRESENT_TO: + case GOMP_MAP_PRESENT_TOFROM: + *ptf_tail = w; + w->next = NULL; + ptf_tail = &w->next; + break; + default: *tf_tail = w; w->next = NULL; @@ -9845,8 +9866,10 @@ omp_segregate_mapping_groups (omp_mapping_group *inlist) /* Now splice the lists together... */ *tf_tail = ard_groups; + *pa_tail = tf_groups; + *ptf_tail = pa_groups; - return tf_groups; + return ptf_groups; } /* Given a list LIST_P containing groups of mappings given by GROUPS, reorder @@ -11698,119 +11721,30 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p, break; } - 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); - if (groups) - { - hash_map *grpmap; - grpmap = omp_index_mapping_groups (groups); + vec *groups = omp_gather_mapping_groups (list_p); + hash_map *grpmap = NULL; + unsigned grpnum = 0; + tree *grp_start_p = NULL, grp_end = NULL_TREE; - omp_resolve_clause_dependencies (code, groups, grpmap); - omp_build_struct_sibling_lists (code, region_type, groups, &grpmap, - list_p); - - omp_mapping_group *outlist = NULL; - bool enter_exit = (code == OMP_TARGET_ENTER_DATA - || code == OMP_TARGET_EXIT_DATA); - - /* 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, enter_exit); - outlist = omp_segregate_mapping_groups (outlist); - list_p = omp_reorder_mapping_groups (groups, outlist, list_p); - - failure: - delete grpmap; - delete groups; - } - - /* OpenMP map clauses with 'present' need to go in front of those - without. */ - tree present_map_head = NULL; - tree *present_map_tail_p = &present_map_head; - tree *first_map_clause_p = NULL; - - for (tree *c_p = list_p; *c_p; ) - { - tree c = *c_p; - tree *next_c_p = &OMP_CLAUSE_CHAIN (c); - - if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP) - { - if (!first_map_clause_p) - first_map_clause_p = c_p; - switch (OMP_CLAUSE_MAP_KIND (c)) - { - case GOMP_MAP_PRESENT_ALLOC: - case GOMP_MAP_PRESENT_FROM: - case GOMP_MAP_PRESENT_TO: - case GOMP_MAP_PRESENT_TOFROM: - next_c_p = c_p; - *c_p = OMP_CLAUSE_CHAIN (c); - - OMP_CLAUSE_CHAIN (c) = NULL; - *present_map_tail_p = c; - present_map_tail_p = &OMP_CLAUSE_CHAIN (c); - - break; - - default: - break; - } - } - - c_p = next_c_p; - } - if (first_map_clause_p && present_map_head) - { - tree next = *first_map_clause_p; - *first_map_clause_p = present_map_head; - *present_map_tail_p = next; - } - } - 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); - - oacc_resolve_clause_dependencies (groups, grpmap); - omp_build_struct_sibling_lists (code, region_type, groups, &grpmap, - list_p); - - delete groups; - delete grpmap; - } - } + if (groups) + grpmap = omp_index_mapping_groups (groups); while ((c = *list_p) != NULL) { bool remove = false; bool notice_outer = true; + bool map_descriptor; const char *check_non_private = NULL; unsigned int flags; tree decl; auto_vec addr_tokens; + if (grp_end && c == OMP_CLAUSE_CHAIN (grp_end)) + { + grp_start_p = NULL; + grp_end = NULL_TREE; + } + switch (OMP_CLAUSE_CODE (c)) { case OMP_CLAUSE_PRIVATE: @@ -12115,45 +12049,26 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p, goto do_add; case OMP_CLAUSE_MAP: + if (!grp_start_p) + { + grp_start_p = list_p; + grp_end = (*groups)[grpnum].grp_end; + grpnum++; + } decl = OMP_CLAUSE_DECL (c); + if (error_operand_p (decl)) + { + remove = true; + break; + } + if (!omp_parse_expr (addr_tokens, decl)) { remove = true; break; } - if (error_operand_p (decl)) - remove = true; - switch (code) - { - case OMP_TARGET: - break; - case OACC_DATA: - if (TREE_CODE (TREE_TYPE (decl)) != ARRAY_TYPE) - break; - goto check_firstprivate; - case OACC_ENTER_DATA: - case OACC_EXIT_DATA: - if (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ATTACH_DETACH - && addr_tokens[0]->type == ARRAY_BASE) - remove = true; - /* FALLTHRU */ - case OMP_TARGET_DATA: - case OMP_TARGET_ENTER_DATA: - case OMP_TARGET_EXIT_DATA: - case OACC_HOST_DATA: - check_firstprivate: - if (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_FIRSTPRIVATE_POINTER - || (OMP_CLAUSE_MAP_KIND (c) - == GOMP_MAP_FIRSTPRIVATE_REFERENCE)) - /* For target {,enter ,exit }data only the array slice is - mapped, but not the pointer to it. */ - remove = true; - break; - default: - break; - } if (remove) break; if (DECL_P (decl) && outer_ctx && (region_type & ORT_ACC)) @@ -12172,41 +12087,61 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p, DECL_NAME (decl)); } } - if (OMP_CLAUSE_SIZE (c) == NULL_TREE) - OMP_CLAUSE_SIZE (c) = DECL_P (decl) ? DECL_SIZE_UNIT (decl) - : TYPE_SIZE_UNIT (TREE_TYPE (decl)); - if (gimplify_expr (&OMP_CLAUSE_SIZE (c), pre_p, - NULL, is_gimple_val, fb_rvalue) == GS_ERROR) - { - remove = true; - break; - } - else if ((OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_FIRSTPRIVATE_POINTER - || (OMP_CLAUSE_MAP_KIND (c) - == GOMP_MAP_FIRSTPRIVATE_REFERENCE) - || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ATTACH_DETACH) - && TREE_CODE (OMP_CLAUSE_SIZE (c)) != INTEGER_CST) - { - OMP_CLAUSE_SIZE (c) - = get_initialized_tmp_var (OMP_CLAUSE_SIZE (c), pre_p, NULL, - false); - if ((region_type & ORT_TARGET) != 0) - omp_add_variable (ctx, OMP_CLAUSE_SIZE (c), - GOVD_FIRSTPRIVATE | GOVD_SEEN); - } - if ((OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_STRUCT - || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_STRUCT_UNORD) - && (addr_tokens[0]->type == STRUCTURE_BASE - || addr_tokens[0]->type == ARRAY_BASE) - && addr_tokens[0]->u.structure_base_kind == BASE_DECL) + map_descriptor = false; + + /* This condition checks if we're mapping an array descriptor that + isn't inside a derived type -- these have special handling, and + are not handled as structs in omp_build_struct_sibling_lists. + See that function for further details. */ + if (*grp_start_p != grp_end + && OMP_CLAUSE_CHAIN (*grp_start_p) + && OMP_CLAUSE_CHAIN (*grp_start_p) != grp_end) + { + tree grp_mid = OMP_CLAUSE_CHAIN (*grp_start_p); + if (omp_map_clause_descriptor_p (grp_mid) + && DECL_P (OMP_CLAUSE_DECL (grp_mid))) + map_descriptor = true; + } + else if (OMP_CLAUSE_CODE (grp_end) == OMP_CLAUSE_MAP + && (OMP_CLAUSE_MAP_KIND (grp_end) == GOMP_MAP_RELEASE + || OMP_CLAUSE_MAP_KIND (grp_end) == GOMP_MAP_DELETE) + && OMP_CLAUSE_RELEASE_DESCRIPTOR (grp_end)) + map_descriptor = true; + + /* Adding the decl for a struct access: we haven't created + GOMP_MAP_STRUCT nodes yet, so this statement needs to predict + whether they will be created in gimplify_adjust_omp_clauses. + NOTE: Technically we should probably look through DECL_VALUE_EXPR + here because something that looks like a DECL_P may actually be a + struct access, e.g. variables in a lambda closure + (__closure->__foo) or class members (this->foo). Currently in both + those cases we map the whole of the containing object (directly in + the C++ FE) though, so struct nodes are not created. */ + if (c == grp_end + && addr_tokens[0]->type == STRUCTURE_BASE + && addr_tokens[0]->u.structure_base_kind == BASE_DECL + && !map_descriptor) { gcc_assert (addr_tokens[1]->type == ACCESS_METHOD); /* If we got to this struct via a chain of pointers, maybe we want to map it implicitly instead. */ if (omp_access_chain_p (addr_tokens, 1)) break; + omp_mapping_group *wholestruct; + if (!(region_type & ORT_ACC) + && omp_mapped_by_containing_struct (grpmap, + OMP_CLAUSE_DECL (c), + &wholestruct)) + break; decl = addr_tokens[1]->expr; + if (splay_tree_lookup (ctx->variables, (splay_tree_key) decl)) + break; + /* Standalone attach or detach clauses for a struct element + should not inhibit implicit mapping of the whole struct. */ + if (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ATTACH + || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_DETACH) + break; flags = GOVD_MAP | GOVD_EXPLICIT; gcc_assert (addr_tokens[1]->u.access_kind != ACCESS_DIRECT @@ -12214,14 +12149,7 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p, goto do_add_decl; } - if (TREE_CODE (decl) == TARGET_EXPR) - { - if (gimplify_expr (&OMP_CLAUSE_DECL (c), pre_p, NULL, - is_gimple_lvalue, fb_lvalue) - == GS_ERROR) - remove = true; - } - else if (!DECL_P (decl)) + if (!DECL_P (decl)) { tree d = decl, *pd; if (TREE_CODE (d) == ARRAY_REF) @@ -12244,56 +12172,20 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p, pd = &TREE_OPERAND (decl, 0); decl = TREE_OPERAND (decl, 0); } - /* 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 - depends on the previous mapping. */ - if (code == OACC_UPDATE - && OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ATTACH_DETACH) - OMP_CLAUSE_SET_MAP_KIND (c, GOMP_MAP_ALWAYS_POINTER); - if (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ATTACH_DETACH) + if (addr_tokens[0]->type == STRUCTURE_BASE + && addr_tokens[0]->u.structure_base_kind == BASE_DECL + && addr_tokens[1]->type == ACCESS_METHOD + && (addr_tokens[1]->u.access_kind == ACCESS_POINTER + || (addr_tokens[1]->u.access_kind + == ACCESS_POINTER_OFFSET)) + && GOMP_MAP_ALWAYS_P (OMP_CLAUSE_MAP_KIND (c))) { - if (TREE_CODE (TREE_TYPE (OMP_CLAUSE_DECL (c))) - == ARRAY_TYPE) - remove = true; - else - { - 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 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) - { - 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; - } + tree base = addr_tokens[1]->expr; + 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)) @@ -12404,56 +12296,12 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p, } } } - else if (gimplify_expr (pd, pre_p, NULL, is_gimple_lvalue, - fb_lvalue) == GS_ERROR) - { - remove = true; - break; - } 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) flags |= GOVD_MAP_ALWAYS_TO; - - if ((code == OMP_TARGET - || code == OMP_TARGET_DATA - || code == OMP_TARGET_ENTER_DATA - || code == OMP_TARGET_EXIT_DATA) - && OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ATTACH_DETACH) - { - for (struct gimplify_omp_ctx *octx = outer_ctx; octx; - octx = octx->outer_context) - { - splay_tree_node n - = splay_tree_lookup (octx->variables, - (splay_tree_key) OMP_CLAUSE_DECL (c)); - /* If this is contained in an outer OpenMP region as a - firstprivate value, remove the attach/detach. */ - if (n && (n->value & GOVD_FIRSTPRIVATE)) - { - OMP_CLAUSE_SET_MAP_KIND (c, GOMP_MAP_FIRSTPRIVATE_POINTER); - goto do_add; - } - } - - enum gomp_map_kind map_kind = (code == OMP_TARGET_EXIT_DATA - ? GOMP_MAP_DETACH - : GOMP_MAP_ATTACH); - OMP_CLAUSE_SET_MAP_KIND (c, map_kind); - } - else if ((code == OACC_ENTER_DATA - || code == OACC_EXIT_DATA - || code == OACC_PARALLEL) - && OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ATTACH_DETACH) - { - enum gomp_map_kind map_kind = (code == OACC_EXIT_DATA - ? GOMP_MAP_DETACH - : GOMP_MAP_ATTACH); - OMP_CLAUSE_SET_MAP_KIND (c, map_kind); - } - goto do_add; case OMP_CLAUSE_AFFINITY: @@ -13092,6 +12940,12 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p, list_p = &OMP_CLAUSE_CHAIN (c); } + if (groups) + { + delete grpmap; + delete groups; + } + ctx->clauses = *orig_list_p; gimplify_omp_ctxp = ctx; } @@ -13562,15 +13416,75 @@ gimplify_adjust_omp_clauses (gimple_seq *pre_p, gimple_seq body, tree *list_p, } } + 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); + hash_map *grpmap = NULL; + + if (groups) + { + grpmap = omp_index_mapping_groups (groups); + + omp_resolve_clause_dependencies (code, groups, grpmap); + omp_build_struct_sibling_lists (code, ctx->region_type, groups, + &grpmap, list_p); + + omp_mapping_group *outlist = NULL; + + delete grpmap; + delete groups; + + /* Rebuild now we have struct sibling lists. */ + groups = omp_gather_mapping_groups (list_p); + grpmap = omp_index_mapping_groups (groups); + + bool enter_exit = (code == OMP_TARGET_ENTER_DATA + || code == OMP_TARGET_EXIT_DATA); + + outlist = omp_tsort_mapping_groups (groups, grpmap, enter_exit); + outlist = omp_segregate_mapping_groups (outlist); + list_p = omp_reorder_mapping_groups (groups, outlist, list_p); + + delete grpmap; + delete groups; + } + } + else if (ctx->region_type & ORT_ACC) + { + vec *groups; + groups = omp_gather_mapping_groups (list_p); + if (groups) + { + hash_map *grpmap; + grpmap = omp_index_mapping_groups (groups); + + oacc_resolve_clause_dependencies (groups, grpmap); + omp_build_struct_sibling_lists (code, ctx->region_type, groups, + &grpmap, list_p); + + delete groups; + delete grpmap; + } + } + tree attach_list = NULL_TREE; tree *attach_tail = &attach_list; + tree *grp_start_p = NULL, grp_end = NULL_TREE; + while ((c = *list_p) != NULL) { splay_tree_node n; bool remove = false; bool move_attach = false; + if (grp_end && c == OMP_CLAUSE_CHAIN (grp_end)) + grp_end = NULL_TREE; + switch (OMP_CLAUSE_CODE (c)) { case OMP_CLAUSE_FIRSTPRIVATE: @@ -13725,6 +13639,12 @@ gimplify_adjust_omp_clauses (gimple_seq *pre_p, gimple_seq body, tree *list_p, break; case OMP_CLAUSE_MAP: + decl = OMP_CLAUSE_DECL (c); + if (!grp_end) + { + grp_start_p = list_p; + grp_end = *omp_group_last (grp_start_p); + } switch (OMP_CLAUSE_MAP_KIND (c)) { case GOMP_MAP_PRESENT_ALLOC: @@ -13736,26 +13656,62 @@ gimplify_adjust_omp_clauses (gimple_seq *pre_p, gimple_seq body, tree *list_p, default: break; } - if (code == OMP_TARGET_EXIT_DATA - && OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ALWAYS_POINTER) + switch (code) { + case OMP_TARGET: + break; + case OACC_DATA: + if (TREE_CODE (TREE_TYPE (decl)) != ARRAY_TYPE) + break; + goto check_firstprivate; + case OACC_ENTER_DATA: + case OACC_EXIT_DATA: + case OMP_TARGET_DATA: + case OMP_TARGET_ENTER_DATA: + case OMP_TARGET_EXIT_DATA: + case OACC_HOST_DATA: + check_firstprivate: + if (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_FIRSTPRIVATE_POINTER + || (OMP_CLAUSE_MAP_KIND (c) + == GOMP_MAP_FIRSTPRIVATE_REFERENCE)) + /* For target {,enter ,exit }data only the array slice is + mapped, but not the pointer to it. */ + remove = true; + if (code == OMP_TARGET_EXIT_DATA + && (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ALWAYS_POINTER + || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_POINTER)) + remove = true; + break; + default: + break; + } + if (remove) + break; + if (OMP_CLAUSE_SIZE (c) == NULL_TREE) + OMP_CLAUSE_SIZE (c) = DECL_P (decl) ? DECL_SIZE_UNIT (decl) + : TYPE_SIZE_UNIT (TREE_TYPE (decl)); + gimplify_omp_ctxp = ctx->outer_context; + if (gimplify_expr (&OMP_CLAUSE_SIZE (c), pre_p, NULL, + is_gimple_val, fb_rvalue) == GS_ERROR) + { + gimplify_omp_ctxp = ctx; remove = true; break; } - /* If we have a target region, we can push all the attaches to the - end of the list (we may have standalone "attach" operations - synthesized for GOMP_MAP_STRUCT nodes that must be processed after - the attachment point AND the pointed-to block have been mapped). - If we have something else, e.g. "enter data", we need to keep - "attach" nodes together with the previous node they attach to so - that separate "exit data" operations work properly (see - libgomp/target.c). */ - if ((ctx->region_type & ORT_TARGET) != 0 - && (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ATTACH - || (OMP_CLAUSE_MAP_KIND (c) - == GOMP_MAP_ATTACH_ZERO_LENGTH_ARRAY_SECTION))) - move_attach = true; - decl = OMP_CLAUSE_DECL (c); + else if ((OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_FIRSTPRIVATE_POINTER + || (OMP_CLAUSE_MAP_KIND (c) + == GOMP_MAP_FIRSTPRIVATE_REFERENCE) + || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ATTACH_DETACH) + && TREE_CODE (OMP_CLAUSE_SIZE (c)) != INTEGER_CST) + { + OMP_CLAUSE_SIZE (c) + = get_initialized_tmp_var (OMP_CLAUSE_SIZE (c), pre_p, NULL, + false); + if ((ctx->region_type & ORT_TARGET) != 0) + omp_add_variable (ctx, OMP_CLAUSE_SIZE (c), + GOVD_FIRSTPRIVATE | GOVD_SEEN); + } + gimplify_omp_ctxp = ctx; /* Data clauses associated with reductions must be compatible with present_or_copy. Warn and adjust the clause if that is not the case. */ @@ -13792,7 +13748,25 @@ gimplify_adjust_omp_clauses (gimple_seq *pre_p, gimple_seq body, tree *list_p, remove = true; break; } - if (!DECL_P (decl)) + /* If we have a DECL_VALUE_EXPR (e.g. this is a class member and/or + a variable captured in a lambda closure), look through that now + before the DECL_P check below. (A code other than COMPONENT_REF, + i.e. INDIRECT_REF, will be a VLA/variable-length array + section. A global var may be a variable in a common block. We + don't want to do this here for either of those.) */ + if ((ctx->region_type & ORT_ACC) == 0 + && DECL_P (decl) + && !is_global_var (decl) + && DECL_HAS_VALUE_EXPR_P (decl) + && TREE_CODE (DECL_VALUE_EXPR (decl)) == COMPONENT_REF) + decl = OMP_CLAUSE_DECL (c) = DECL_VALUE_EXPR (decl); + if (TREE_CODE (decl) == TARGET_EXPR) + { + if (gimplify_expr (&OMP_CLAUSE_DECL (c), pre_p, NULL, + is_gimple_lvalue, fb_lvalue) == GS_ERROR) + remove = true; + } + else if (!DECL_P (decl)) { if ((ctx->region_type & ORT_TARGET) != 0 && OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_FIRSTPRIVATE_POINTER) @@ -13815,8 +13789,122 @@ gimplify_adjust_omp_clauses (gimple_seq *pre_p, gimple_seq body, tree *list_p, } } } + + tree d = decl, *pd; + 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; + } + pd = &OMP_CLAUSE_DECL (c); + 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)) + { + pd = &TREE_OPERAND (decl, 0); + decl = TREE_OPERAND (decl, 0); + } + + if (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ATTACH_DETACH) + switch (code) + { + case OACC_ENTER_DATA: + case OACC_EXIT_DATA: + if (TREE_CODE (TREE_TYPE (OMP_CLAUSE_DECL (c))) + == ARRAY_TYPE) + remove = true; + else if (code == OACC_ENTER_DATA) + goto change_to_attach; + /* Fallthrough. */ + case OMP_TARGET_EXIT_DATA: + OMP_CLAUSE_SET_MAP_KIND (c, GOMP_MAP_DETACH); + break; + case OACC_UPDATE: + /* An "attach/detach" operation on an update directive + should behave as a GOMP_MAP_ALWAYS_POINTER. Note that + both GOMP_MAP_ATTACH_DETACH and GOMP_MAP_ALWAYS_POINTER + kinds depend on the previous mapping (for non-TARGET + regions). */ + OMP_CLAUSE_SET_MAP_KIND (c, GOMP_MAP_ALWAYS_POINTER); + break; + default: + change_to_attach: + OMP_CLAUSE_SET_MAP_KIND (c, GOMP_MAP_ATTACH); + if ((ctx->region_type & ORT_TARGET) != 0) + move_attach = true; + } + else if ((ctx->region_type & ORT_TARGET) != 0 + && (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ATTACH + || (OMP_CLAUSE_MAP_KIND (c) + == GOMP_MAP_ATTACH_ZERO_LENGTH_ARRAY_SECTION))) + move_attach = true; + + /* If we have e.g. map(struct: *var), don't gimplify the + argument since omp-low.cc wants to see the decl itself. */ + if (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_STRUCT) + break; + + /* We've already partly gimplified this in + gimplify_scan_omp_clauses. Don't do any more. */ + if (code == OMP_TARGET && OMP_CLAUSE_MAP_IN_REDUCTION (c)) + break; + + gimplify_omp_ctxp = ctx->outer_context; + if (gimplify_expr (pd, pre_p, NULL, is_gimple_lvalue, + fb_lvalue) == GS_ERROR) + remove = true; + gimplify_omp_ctxp = ctx; break; } + + if ((code == OMP_TARGET + || code == OMP_TARGET_DATA + || code == OMP_TARGET_ENTER_DATA + || code == OMP_TARGET_EXIT_DATA) + && OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ATTACH_DETACH) + { + bool firstprivatize = false; + + for (struct gimplify_omp_ctx *octx = ctx->outer_context; octx; + octx = octx->outer_context) + { + splay_tree_node n + = splay_tree_lookup (octx->variables, + (splay_tree_key) OMP_CLAUSE_DECL (c)); + /* If this is contained in an outer OpenMP region as a + firstprivate value, remove the attach/detach. */ + if (n && (n->value & GOVD_FIRSTPRIVATE)) + { + firstprivatize = true; + break; + } + } + + enum gomp_map_kind map_kind; + if (firstprivatize) + map_kind = GOMP_MAP_FIRSTPRIVATE_POINTER; + else if (code == OMP_TARGET_EXIT_DATA) + map_kind = GOMP_MAP_DETACH; + else + map_kind = GOMP_MAP_ATTACH; + OMP_CLAUSE_SET_MAP_KIND (c, map_kind); + } + else if ((ctx->region_type & ORT_ACC) != 0 + && OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ATTACH_DETACH) + { + enum gomp_map_kind map_kind = (code == OACC_EXIT_DATA + ? GOMP_MAP_DETACH + : GOMP_MAP_ATTACH); + OMP_CLAUSE_SET_MAP_KIND (c, map_kind); + } + n = splay_tree_lookup (ctx->variables, (splay_tree_key) decl); if ((ctx->region_type & ORT_TARGET) != 0 && !(n->value & GOVD_SEEN) @@ -13891,6 +13979,21 @@ gimplify_adjust_omp_clauses (gimple_seq *pre_p, gimple_seq body, tree *list_p, || ((n->value & (GOVD_PRIVATE | GOVD_FIRSTPRIVATE)) == 0)); } + + /* If we have a target region, we can push all the attaches to the + end of the list (we may have standalone "attach" operations + synthesized for GOMP_MAP_STRUCT nodes that must be processed after + the attachment point AND the pointed-to block have been mapped). + If we have something else, e.g. "enter data", we need to keep + "attach" nodes together with the previous node they attach to so + that separate "exit data" operations work properly (see + libgomp/target.c). */ + if ((ctx->region_type & ORT_TARGET) != 0 + && (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ATTACH + || (OMP_CLAUSE_MAP_KIND (c) + == GOMP_MAP_ATTACH_ZERO_LENGTH_ARRAY_SECTION))) + move_attach = true; + break; case OMP_CLAUSE_TO: diff --git a/gcc/testsuite/gfortran.dg/gomp/map-12.f90 b/gcc/testsuite/gfortran.dg/gomp/map-12.f90 index ac9a0f8aae04..433bf98911c8 100644 --- a/gcc/testsuite/gfortran.dg/gomp/map-12.f90 +++ b/gcc/testsuite/gfortran.dg/gomp/map-12.f90 @@ -60,7 +60,7 @@ end subroutine ! { dg-final { scan-tree-dump-times "#pragma omp target data map\\(always,present,tofrom:b1 \\\[len: 4\\\]\\)\[\r\n\]" 2 "omplower" } } ! { dg-final { scan-tree-dump-times "#pragma omp target enter data map\\(to:b \\\[len: 4\\\]\\) map\\(to:b1 \\\[len: 4\\\]\\) map\\(alloc:a \\\[len: 4\\\]\\)\[\r\n\]" 2 "omplower" } } ! { dg-final { scan-tree-dump-times "#pragma omp target enter data map\\(always,to:b \\\[len: 4\\\]\\) map\\(always,to:b1 \\\[len: 4\\\]\\) map\\(alloc:a \\\[len: 4\\\]\\)\[\r\n\]" 2 "omplower" } } -! { dg-final { scan-tree-dump-times "#pragma omp target enter data map\\(force_present:a \\\[len: 4\\\]\\) map\\(force_present:b \\\[len: 4\\\]\\) map\\(force_present:b1 \\\[len: 4\\\]\\)\[\r\n\]" 2 "omplower" } } +! { dg-final { scan-tree-dump-times "#pragma omp target enter data map\\(force_present:b \\\[len: 4\\\]\\) map\\(force_present:b1 \\\[len: 4\\\]\\) map\\(force_present:a \\\[len: 4\\\]\\)\[\r\n\]" 2 "omplower" } } ! { dg-final { scan-tree-dump-times "#pragma omp target enter data map\\(force_present:a \\\[len: 4\\\]\\) map\\(always,present,to:b \\\[len: 4\\\]\\) map\\(always,present,to:b1 \\\[len: 4\\\]\\)\[\r\n\]" 2 "omplower" } } ! { dg-final { scan-tree-dump-times "#pragma omp target exit data map\\(from:b1 \\\[len: 4\\\]\\) map\\(delete:a \\\[len: 4\\\]\\) map\\(release:b \\\[len: 4\\\]\\)\[\r\n\]" 2 "omplower" } } ! { dg-final { scan-tree-dump-times "#pragma omp target exit data map\\(always,from:b1 \\\[len: 4\\\]\\) map\\(delete:a \\\[len: 4\\\]\\) map\\(release:b \\\[len: 4\\\]\\)\[\r\n\]" 2 "omplower" } }