From patchwork Fri Oct 1 17:07:48 2021 Content-Type: text/plain; charset="utf-8" MIME-Version: 1.0 Content-Transfer-Encoding: 7bit X-Patchwork-Submitter: Julian Brown X-Patchwork-Id: 45697 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 E0352385742E for ; Fri, 1 Oct 2021 17:08:58 +0000 (GMT) X-Original-To: gcc-patches@gcc.gnu.org Delivered-To: gcc-patches@gcc.gnu.org Received: from esa3.mentor.iphmx.com (esa3.mentor.iphmx.com [68.232.137.180]) by sourceware.org (Postfix) with ESMTPS id 40E353857829 for ; Fri, 1 Oct 2021 17:08:09 +0000 (GMT) DMARC-Filter: OpenDMARC Filter v1.4.1 sourceware.org 40E353857829 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: 4GUL6GX+wCBTDV95KQmP8u+7eOy002QCYeSZEimhABErlzPWoT0j8pc21ekx7hixrD765W5DF5 U4EP8HCdHmWQzoUtJSGFuFzBlkb9zg9ECYIREDggpHQo3KCqZImF15UAXot9HOgvmHflQ6r/l6 31sM73jms+vjED4rCDRWgjpX/UaknPOp/h6c14MWbSNWgPMRYeGitKchfdIFTy7kEkWcyjJcy+ Cqr6ZXmDmX28WJ4TvN3CYf0lWFUWKVojvXqM5HHCwMSYScqAVt05wvATcgQq2JH8Wyv24JsYyW I95vONc+dWCaD9RwSC8RQKsb X-IronPort-AV: E=Sophos;i="5.85,339,1624348800"; d="scan'208";a="66548477" Received: from orw-gwy-01-in.mentorg.com ([192.94.38.165]) by esa3.mentor.iphmx.com with ESMTP; 01 Oct 2021 09:08:09 -0800 IronPort-SDR: p+n4wL0FlIzi5rgYSNsWorjUsx1PTtAB8iloCtvx0pLwcs4dU5kzzXv4iB4Zcflginx/q+UQ4n 40PwMy+RQhm7IbSWDH7Gp2D60CRBeyqrCR4f0QYledDBbDu2lnaMJ2zAKM8f+VXHCjHqQRiy1Y dzQS9G1RQWMTP9AzNe+dWywlCLG8IGSgcIjCK2WXWxlGyZVLGGAOj5xrnqTpwlSgI+ZHEMAbPJ N5itvnk58ja8b5ihcCx/64rgIOcuaCpgWH/e1qYqTo7dm/i/GoCn7Xm3sBpOZPsO0oUin2sBq7 BwY= From: Julian Brown To: Subject: [PATCH 01/11] libgomp: Release device lock on cbuf error path Date: Fri, 1 Oct 2021 10:07:48 -0700 Message-ID: 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-07.mgc.mentorg.com (139.181.222.7) 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 releases the device lock on a sanity-checking error path in transfer combining (cbuf) handling in libgomp:target.c. This shouldn't happen when handling well-formed mapping clauses, but erroneous clauses can currently cause a hang if the condition triggers. Tested with offloading to NVPTX. OK? 2021-09-29 Julian Brown libgomp/ * target.c (gomp_copy_host2dev): Release device lock on cbuf error path. --- libgomp/target.c | 5 ++++- 1 file changed, 4 insertions(+), 1 deletion(-) diff --git a/libgomp/target.c b/libgomp/target.c index 65bb40100e5..84c6fdf2c47 100644 --- a/libgomp/target.c +++ b/libgomp/target.c @@ -385,7 +385,10 @@ gomp_copy_host2dev (struct gomp_device_descr *devicep, else if (cbuf->chunks[middle].start <= doff) { if (doff + sz > cbuf->chunks[middle].end) - gomp_fatal ("internal libgomp cbuf error"); + { + gomp_mutex_unlock (&devicep->lock); + gomp_fatal ("internal libgomp cbuf error"); + } memcpy ((char *) cbuf->buf + (doff - cbuf->chunks[0].start), h, sz); return; From patchwork Fri Oct 1 17:07:49 2021 Content-Type: text/plain; charset="utf-8" MIME-Version: 1.0 Content-Transfer-Encoding: 7bit X-Patchwork-Submitter: Julian Brown X-Patchwork-Id: 45698 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 7BBB83857007 for ; Fri, 1 Oct 2021 17:09:28 +0000 (GMT) X-Original-To: gcc-patches@gcc.gnu.org Delivered-To: gcc-patches@gcc.gnu.org Received: from esa3.mentor.iphmx.com (esa3.mentor.iphmx.com [68.232.137.180]) by sourceware.org (Postfix) with ESMTPS id 457AD3857C7E for ; Fri, 1 Oct 2021 17:08:13 +0000 (GMT) DMARC-Filter: OpenDMARC Filter v1.4.1 sourceware.org 457AD3857C7E 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: PS8fzPsrf7m7X1tJVcw+OjNL1NqdCdP/GtpkNiP9Ljx7yeJ9FDD1i0kirVKqW6qGTwvx4/Ax+N YivyuhGrGOo8W53zRNI1S7fA6wkocbxZ+HIxMjUYTNynwL7yReWisfmIWOM9/MpoCxZ/528P79 xOOY7YkcxTAUKW2jJQ818/Dd83nanU7o8MS70dFhl26IenuYMCtLojFWyiuZkDKFQze0zI8opQ uNr7UArsGWPppGcPIkho0V0+lDixPN1+PZ61XTl2t10i+Cdbl1LgZUDsfJCARFCtpymLqe+f6U Pkqr5JEdVRPX12FyK+hBtCFL X-IronPort-AV: E=Sophos;i="5.85,339,1624348800"; d="scan'208";a="66548482" Received: from orw-gwy-01-in.mentorg.com ([192.94.38.165]) by esa3.mentor.iphmx.com with ESMTP; 01 Oct 2021 09:08:13 -0800 IronPort-SDR: hKFIWeJI0a2f6FVD3jtl0fNymQKG7nhMnoNt+/YwRzTbebF8FwniXSJ0zWxpA7UEdUKrEgh+Tg 4CqtvH/jBCVvxNaBSQ6tKWACEnMcomB68oNv72x45wh0+jNmoCOqgGjJ70Edpy7I75NXEFHF9c 6B8nTLioIaDw70zPVuaPfz9vdUOli89XiWk49UDA/3dcR8EIbU52D1OGgHhZb+wPO/Q7ZHeWxq 1cio2c5bc8QyUOpvo4vjvgbv+8UGVvxLhpXeMH76sBTttiRQYbsSdg3Hxps0b0b7Yt753nfX/e lps= From: Julian Brown To: Subject: [PATCH 02/11] Remove base_ind/base_ref handling from extract_base_bit_offset Date: Fri, 1 Oct 2021 10:07:49 -0700 Message-ID: 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-07.mgc.mentorg.com (139.181.222.7) 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" In preparation for follow-up patches extending struct dereference handling for OpenMP, this patch removes base_ind/base_ref handling from gimplify.c:extract_base_bit_offset. This arguably simplifies some of the code around the callers of the function also, though subsequent patches modify those parts further. OK for mainline? Thanks, Julian 2021-09-29 Julian Brown gcc/ * gimplify.c (extract_base_bit_offset): Remove BASE_IND, BASE_REF and OPENMP parameters. (strip_indirections): New function. (build_struct_group): Update calls to extract_base_bit_offset. Rearrange indirect/reference handling accordingly. Use extracted base instead of passed-in decl when grouping component accesses together. --- gcc/gimplify.c | 109 ++++++++++++++++++++++++++----------------------- 1 file changed, 57 insertions(+), 52 deletions(-) diff --git a/gcc/gimplify.c b/gcc/gimplify.c index 92f8a7b4073..ece22b7a4ae 100644 --- a/gcc/gimplify.c +++ b/gcc/gimplify.c @@ -8641,9 +8641,8 @@ build_struct_comp_nodes (enum tree_code code, tree grp_start, tree grp_end, has array type, else return NULL. */ static tree -extract_base_bit_offset (tree base, tree *base_ind, tree *base_ref, - poly_int64 *bitposp, poly_offset_int *poffsetp, - tree *offsetp, bool openmp) +extract_base_bit_offset (tree base, poly_int64 *bitposp, + poly_offset_int *poffsetp, tree *offsetp) { tree offset; poly_int64 bitsize, bitpos; @@ -8651,38 +8650,12 @@ extract_base_bit_offset (tree base, tree *base_ind, tree *base_ref, int unsignedp, reversep, volatilep = 0; poly_offset_int poffset; - if (base_ind) - *base_ind = NULL_TREE; - - if (base_ref) - *base_ref = NULL_TREE; + STRIP_NOPS (base); base = get_inner_reference (base, &bitsize, &bitpos, &offset, &mode, &unsignedp, &reversep, &volatilep); - if (!openmp - && (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) - { - if (base_ind) - *base_ind = base; - base = TREE_OPERAND (base, 0); - } - if ((TREE_CODE (base) == INDIRECT_REF - || (TREE_CODE (base) == MEM_REF - && integer_zerop (TREE_OPERAND (base, 1)))) - && DECL_P (TREE_OPERAND (base, 0)) - && TREE_CODE (TREE_TYPE (TREE_OPERAND (base, 0))) == REFERENCE_TYPE) - { - if (base_ref) - *base_ref = base; - base = TREE_OPERAND (base, 0); - } - - if (!openmp) - STRIP_NOPS (base); + STRIP_NOPS (base); if (offset && poly_int_tree_p (offset)) { @@ -8739,6 +8712,17 @@ strip_components_and_deref (tree expr) return expr; } +static tree +strip_indirections (tree expr) +{ + while (TREE_CODE (expr) == INDIRECT_REF + || (TREE_CODE (expr) == MEM_REF + && integer_zerop (TREE_OPERAND (expr, 1)))) + expr = TREE_OPERAND (expr, 0); + + return expr; +} + /* Return TRUE if EXPR is something we will use as the base of an aggregate access, either: @@ -9232,7 +9216,7 @@ build_struct_group (struct gimplify_omp_ctx *ctx, { poly_offset_int coffset; poly_int64 cbitpos; - tree base_ind, base_ref, tree_coffset; + tree tree_coffset; tree ocd = OMP_CLAUSE_DECL (c); bool openmp = !(region_type & ORT_ACC); @@ -9242,10 +9226,25 @@ 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, &base_ind, &base_ref, &cbitpos, - &coffset, &tree_coffset, openmp); + tree base = extract_base_bit_offset (ocd, &cbitpos, &coffset, &tree_coffset); + tree sbase; - bool do_map_struct = (base == decl && !tree_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); + + 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" @@ -9305,19 +9304,12 @@ build_struct_group (struct gimplify_omp_ctx *ctx, OMP_CLAUSE_SET_MAP_KIND (l, k); - if (!openmp && base_ind) - OMP_CLAUSE_DECL (l) = unshare_expr (base_ind); - else if (base_ref) - OMP_CLAUSE_DECL (l) = unshare_expr (base_ref); - else - { - OMP_CLAUSE_DECL (l) = unshare_expr (decl); - 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_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)) @@ -9353,6 +9345,20 @@ build_struct_group (struct gimplify_omp_ctx *ctx, 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 @@ -9485,10 +9491,9 @@ build_struct_group (struct gimplify_omp_ctx *ctx, == REFERENCE_TYPE)) sc_decl = TREE_OPERAND (sc_decl, 0); - tree base = extract_base_bit_offset (sc_decl, NULL, NULL, - &bitpos, &offset, - &tree_offset, openmp); - if (!base || !operand_equal_p (base, decl, 0)) + tree base2 = extract_base_bit_offset (sc_decl, &bitpos, &offset, + &tree_offset); + if (!base2 || !operand_equal_p (base2, base, 0)) break; if (scp) continue; From patchwork Fri Oct 1 17:07:50 2021 Content-Type: text/plain; charset="utf-8" MIME-Version: 1.0 Content-Transfer-Encoding: 7bit X-Patchwork-Submitter: Julian Brown X-Patchwork-Id: 45699 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 BD812385740B for ; Fri, 1 Oct 2021 17:09:58 +0000 (GMT) X-Original-To: gcc-patches@gcc.gnu.org Delivered-To: gcc-patches@gcc.gnu.org Received: from esa3.mentor.iphmx.com (esa3.mentor.iphmx.com [68.232.137.180]) by sourceware.org (Postfix) with ESMTPS id DF5403857405 for ; Fri, 1 Oct 2021 17:08:14 +0000 (GMT) DMARC-Filter: OpenDMARC Filter v1.4.1 sourceware.org DF5403857405 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: aJ2oOSg9DwOO2r14T7U4OZv18Hteoxdew2e0BxAjwJQnWK/Z3V4gkvybfyuoV/mo/Sv5FvBCXu NbSdMDPhT+YqbuVFUsIAcVQAbGWUaSkBQaZI7V3GvAHYNNs2OjsMSsr0DN7SGxQpA+mWEAnC2E rnovec+2tIBioc19vVIgSheg/YWKaIpFq9rApKb0I+bFZoUfTd6fq8yH+pPabfj/c1CLRsuHnK SMYSmrT4SgxuVCaiZM9X0x0d4YN2UR8c+iiwKqgcxLi/kiBywZvvntFOlZpR08HhxIwsktXK6b lGjenifQDkRU8TNzgwI7Eg/r X-IronPort-AV: E=Sophos;i="5.85,339,1624348800"; d="scan'208";a="66548486" Received: from orw-gwy-01-in.mentorg.com ([192.94.38.165]) by esa3.mentor.iphmx.com with ESMTP; 01 Oct 2021 09:08:14 -0800 IronPort-SDR: ZroEWvpuVlT75ui4b2BSs5ZMThv3r/cLseWL1md0XCUjLou1sNBBySwqhHmvP9uWTL9SjbhiHC rBNNFXvKewpY/O5+Kbl81kAS2GHyTbJZxoFJxB4TdAFDZqpmFDrG+GHSR33mxS7kXyjW1MTLfs WSEP+KUt8u2ZBzFFwryj6kuW9HTXWJq4KhQN1DI1JyvXihc0s6M4PUABobozdOJIAH5zql2b9P 1XzjxJ52ARvZB7tOmSO7+TMo46xmcvBbmpkru9Ed2BroDswBi0vpPp/wJg2bz9x3rwPa26EvGl xCw= From: Julian Brown To: Subject: [PATCH 03/11] OpenMP 5.0: Clause ordering for OpenMP 5.0 (topological sorting by base pointer) Date: Fri, 1 Oct 2021 10:07:50 -0700 Message-ID: <339014dc992f87bfaaa9307e80b19194a6c1fd45.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-07.mgc.mentorg.com (139.181.222.7) 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: , Cc: Jakub Jelinek , Thomas Schwinge Errors-To: gcc-patches-bounces+patchwork=sourceware.org@gcc.gnu.org Sender: "Gcc-patches" This patch reimplements the omp_target_reorder_clauses function in anticipation of supporting "deeper" struct mappings (that is, with several structure dereference operators, or similar). The idea is that in place of the (possibly quadratic) algorithm in omp_target_reorder_clauses that greedily moves clauses containing addresses that are subexpressions of other addresses before those other addresses, we employ a topological sort algorithm to calculate a proper order for map clauses. This should run in linear time, and hopefully handles degenerate cases where multiple "levels" of indirect accesses are present on a given directive. The new method also takes care to keep clause groups together, addressing the concerns raised in: https://gcc.gnu.org/pipermail/gcc-patches/2021-May/570501.html To figure out if some given clause depends on a base pointer in another clause, we strip off the outer layers of the address expression, and check (via a tree_operand_hash hash table we have built) if the result is a "base pointer" as defined in OpenMP 5.0 (1.2.6 Data Terminology). There are some subtleties involved, however: - We must treat MEM_REF with zero offset the same as INDIRECT_REF. This should probably be fixed in the front ends instead so we always use a canonical form (probably INDIRECT_REF). The following patch shows one instance of the problem, but there may be others: https://gcc.gnu.org/pipermail/gcc-patches/2021-May/571382.html - Mapping a whole struct implies mapping each of that struct's elements, which may be base pointers. Because those base pointers aren't necessarily explicitly referenced in the directive in question, we treat the whole-struct mapping as a dependency instead. This version of the patch is significantly improved over the version posted previously in order to support the subsequent patches in this series. OK for mainline? Thanks, Julian 2021-09-29 Julian Brown gcc/ * gimplify.c (is_or_contains_p, omp_target_reorder_clauses): Delete functions. (omp_tsort_mark): Add enum. (omp_mapping_group): Add struct. (omp_get_base_pointer, omp_get_attachment, omp_group_last, omp_gather_mapping_groups, omp_group_base, omp_index_mapping_groups, omp_containing_struct, omp_tsort_mapping_groups_1, omp_tsort_mapping_groups, omp_segregate_mapping_groups, omp_reorder_mapping_groups): New functions. (gimplify_scan_omp_clauses): Call above functions instead of omp_target_reorder_clauses, unless we've seen an error. * omp-low.c (scan_sharing_clauses): Avoid strict test if we haven't sorted mapping groups. gcc/testsuite/ * g++.dg/gomp/target-lambda-1.C: Adjust expected output. * g++.dg/gomp/target-this-3.C: Likewise. * g++.dg/gomp/target-this-4.C: Likewise. --- gcc/gimplify.c | 804 +++++++++++++++++++- gcc/omp-low.c | 7 +- gcc/testsuite/g++.dg/gomp/target-lambda-1.C | 6 +- gcc/testsuite/g++.dg/gomp/target-this-3.C | 4 +- gcc/testsuite/g++.dg/gomp/target-this-4.C | 4 +- 5 files changed, 788 insertions(+), 37 deletions(-) diff --git a/gcc/gimplify.c b/gcc/gimplify.c index ece22b7a4ae..d3346fc8d35 100644 --- a/gcc/gimplify.c +++ b/gcc/gimplify.c @@ -8675,29 +8675,6 @@ extract_base_bit_offset (tree base, poly_int64 *bitposp, return base; } -/* Returns true if EXPR is or contains (as a sub-component) BASE_PTR. */ - -static bool -is_or_contains_p (tree expr, tree 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); -} - /* Remove COMPONENT_REFS and indirections from EXPR. */ static tree @@ -8751,6 +8728,7 @@ aggregate_base_p (tree expr) return false; } +#if 0 /* Implement OpenMP 5.x map ordering rules for target directives. There are several rules, and with some level of ambiguity, hopefully we can at least collect the complexity here in one place. */ @@ -8930,6 +8908,758 @@ omp_target_reorder_clauses (tree *list_p) } } } +#endif + + +enum omp_tsort_mark { + UNVISITED, + TEMPORARY, + PERMANENT +}; + +struct omp_mapping_group { + tree *grp_start; + tree grp_end; + omp_tsort_mark mark; + struct omp_mapping_group *sibling; + struct omp_mapping_group *next; +}; + +__attribute__((used)) static void +debug_mapping_group (omp_mapping_group *grp) +{ + tree tmp = OMP_CLAUSE_CHAIN (grp->grp_end); + OMP_CLAUSE_CHAIN (grp->grp_end) = NULL; + debug_generic_expr (*grp->grp_start); + OMP_CLAUSE_CHAIN (grp->grp_end) = tmp; +} + +/* Return the OpenMP "base pointer" of an expression EXPR, or NULL if there + isn't one. This needs improvement. */ + +static tree +omp_get_base_pointer (tree expr) +{ + while (TREE_CODE (expr) == ARRAY_REF) + expr = TREE_OPERAND (expr, 0); + + while (TREE_CODE (expr) == COMPONENT_REF + && (DECL_P (TREE_OPERAND (expr, 0)) + || (TREE_CODE (TREE_OPERAND (expr, 0)) == 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))) + || TREE_CODE (TREE_OPERAND (expr, 0)) == ARRAY_REF)) + { + expr = TREE_OPERAND (expr, 0); + + while (TREE_CODE (expr) == ARRAY_REF) + expr = TREE_OPERAND (expr, 0); + + if (TREE_CODE (expr) == INDIRECT_REF || TREE_CODE (expr) == MEM_REF) + break; + } + + if (DECL_P (expr)) + return NULL_TREE; + + if (TREE_CODE (expr) == INDIRECT_REF + || TREE_CODE (expr) == MEM_REF) + { + expr = TREE_OPERAND (expr, 0); + while (TREE_CODE (expr) == COMPOUND_EXPR) + expr = TREE_OPERAND (expr, 1); + if (TREE_CODE (expr) == POINTER_PLUS_EXPR) + expr = TREE_OPERAND (expr, 0); + if (TREE_CODE (expr) == SAVE_EXPR) + expr = TREE_OPERAND (expr, 0); + STRIP_NOPS (expr); + return expr; + } + + return NULL_TREE; +} + +/* An attach or detach operation depends directly on the address being + attached/detached. Return that address, or none if there are no + attachments/detachments. */ + +static tree +omp_get_attachment (omp_mapping_group *grp) +{ + tree node = *grp->grp_start; + + switch (OMP_CLAUSE_MAP_KIND (node)) + { + case GOMP_MAP_TO: + case GOMP_MAP_FROM: + case GOMP_MAP_TOFROM: + case GOMP_MAP_ALWAYS_FROM: + case GOMP_MAP_ALWAYS_TO: + case GOMP_MAP_ALWAYS_TOFROM: + case GOMP_MAP_FORCE_FROM: + case GOMP_MAP_FORCE_TO: + case GOMP_MAP_FORCE_TOFROM: + case GOMP_MAP_FORCE_PRESENT: + case GOMP_MAP_ALLOC: + case GOMP_MAP_RELEASE: + case GOMP_MAP_DELETE: + case GOMP_MAP_FORCE_ALLOC: + if (node == grp->grp_end) + return NULL_TREE; + + node = OMP_CLAUSE_CHAIN (node); + if (node && OMP_CLAUSE_MAP_KIND (node) == GOMP_MAP_TO_PSET) + { + gcc_assert (node != grp->grp_end); + node = OMP_CLAUSE_CHAIN (node); + } + if (node) + switch (OMP_CLAUSE_MAP_KIND (node)) + { + case GOMP_MAP_POINTER: + case GOMP_MAP_ALWAYS_POINTER: + case GOMP_MAP_FIRSTPRIVATE_POINTER: + case GOMP_MAP_FIRSTPRIVATE_REFERENCE: + case GOMP_MAP_POINTER_TO_ZERO_LENGTH_ARRAY_SECTION: + return NULL_TREE; + + case GOMP_MAP_ATTACH_DETACH: + case GOMP_MAP_ATTACH_ZERO_LENGTH_ARRAY_SECTION: + return OMP_CLAUSE_DECL (node); + + default: + internal_error ("unexpected mapping node"); + } + return error_mark_node; + + case GOMP_MAP_TO_PSET: + gcc_assert (node != grp->grp_end); + node = OMP_CLAUSE_CHAIN (node); + if (OMP_CLAUSE_MAP_KIND (node) == GOMP_MAP_ATTACH + || OMP_CLAUSE_MAP_KIND (node) == GOMP_MAP_DETACH) + return OMP_CLAUSE_DECL (node); + else + internal_error ("unexpected mapping node"); + return error_mark_node; + + case GOMP_MAP_ATTACH: + case GOMP_MAP_DETACH: + node = OMP_CLAUSE_CHAIN (node); + if (!node || *grp->grp_start == grp->grp_end) + return OMP_CLAUSE_DECL (*grp->grp_start); + if (OMP_CLAUSE_MAP_KIND (node) == GOMP_MAP_FIRSTPRIVATE_POINTER + || OMP_CLAUSE_MAP_KIND (node) == GOMP_MAP_FIRSTPRIVATE_REFERENCE) + return OMP_CLAUSE_DECL (*grp->grp_start); + else + internal_error ("unexpected mapping node"); + return error_mark_node; + + case GOMP_MAP_STRUCT: + case GOMP_MAP_FORCE_DEVICEPTR: + case GOMP_MAP_DEVICE_RESIDENT: + case GOMP_MAP_LINK: + case GOMP_MAP_IF_PRESENT: + case GOMP_MAP_FIRSTPRIVATE: + case GOMP_MAP_FIRSTPRIVATE_INT: + case GOMP_MAP_USE_DEVICE_PTR: + case GOMP_MAP_ATTACH_ZERO_LENGTH_ARRAY_SECTION: + return NULL_TREE; + + default: + internal_error ("unexpected mapping node"); + } + + return error_mark_node; +} + +/* Given a pointer START_P to the start of a group of related (e.g. pointer) + mappings, return the chain pointer to the end of that group in the list. */ + +static tree * +omp_group_last (tree *start_p) +{ + tree c = *start_p, nc, *grp_last_p = start_p; + + gcc_assert (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP); + + nc = OMP_CLAUSE_CHAIN (c); + + if (!nc || OMP_CLAUSE_CODE (nc) != OMP_CLAUSE_MAP) + return grp_last_p; + + switch (OMP_CLAUSE_MAP_KIND (c)) + { + default: + while (nc + && OMP_CLAUSE_CODE (nc) == OMP_CLAUSE_MAP + && (OMP_CLAUSE_MAP_KIND (nc) == GOMP_MAP_FIRSTPRIVATE_REFERENCE + || OMP_CLAUSE_MAP_KIND (nc) == GOMP_MAP_FIRSTPRIVATE_POINTER + || OMP_CLAUSE_MAP_KIND (nc) == GOMP_MAP_ATTACH_DETACH + || OMP_CLAUSE_MAP_KIND (nc) == GOMP_MAP_POINTER + || (OMP_CLAUSE_MAP_KIND (nc) + == GOMP_MAP_POINTER_TO_ZERO_LENGTH_ARRAY_SECTION) + || (OMP_CLAUSE_MAP_KIND (nc) + == GOMP_MAP_ATTACH_ZERO_LENGTH_ARRAY_SECTION) + || OMP_CLAUSE_MAP_KIND (nc) == GOMP_MAP_ALWAYS_POINTER + || OMP_CLAUSE_MAP_KIND (nc) == GOMP_MAP_TO_PSET)) + { + grp_last_p = &OMP_CLAUSE_CHAIN (c); + c = nc; + tree nc2 = OMP_CLAUSE_CHAIN (nc); + if (nc2 + && OMP_CLAUSE_CODE (nc2) == OMP_CLAUSE_MAP + && (OMP_CLAUSE_MAP_KIND (nc) + == GOMP_MAP_POINTER_TO_ZERO_LENGTH_ARRAY_SECTION) + && OMP_CLAUSE_MAP_KIND (nc2) == GOMP_MAP_ATTACH) + { + grp_last_p = &OMP_CLAUSE_CHAIN (nc); + c = nc2; + nc2 = OMP_CLAUSE_CHAIN (nc2); + } + nc = nc2; + } + break; + + case GOMP_MAP_ATTACH: + case GOMP_MAP_DETACH: + /* This is a weird artifact of how directives are parsed: bare attach or + detach clauses get a subsequent (meaningless) FIRSTPRIVATE_POINTER or + FIRSTPRIVATE_REFERENCE node. FIXME. */ + if (nc + && OMP_CLAUSE_CODE (nc) == OMP_CLAUSE_MAP + && (OMP_CLAUSE_MAP_KIND (nc) == GOMP_MAP_FIRSTPRIVATE_REFERENCE + || OMP_CLAUSE_MAP_KIND (nc) == GOMP_MAP_FIRSTPRIVATE_POINTER)) + grp_last_p = &OMP_CLAUSE_CHAIN (c); + break; + + case GOMP_MAP_TO_PSET: + if (OMP_CLAUSE_CODE (nc) == OMP_CLAUSE_MAP + && (OMP_CLAUSE_MAP_KIND (nc) == GOMP_MAP_ATTACH + || OMP_CLAUSE_MAP_KIND (nc) == GOMP_MAP_DETACH)) + grp_last_p = &OMP_CLAUSE_CHAIN (c); + break; + } + + return grp_last_p; +} + +/* Walk through LIST_P, and return a list of groups of mappings found (e.g. + OMP_CLAUSE_MAP with GOMP_MAP_{TO/FROM/TOFROM} followed by one or two + associated GOMP_MAP_POINTER mappings). Return a vector of omp_mapping_group + if we have more than one such group, else return NULL. */ + +static vec * +omp_gather_mapping_groups (tree *list_p) +{ + vec *groups = new vec (); + + for (tree *cp = list_p; *cp; cp = &OMP_CLAUSE_CHAIN (*cp)) + { + if (OMP_CLAUSE_CODE (*cp) != OMP_CLAUSE_MAP) + continue; + + tree *grp_last_p = omp_group_last (cp); + omp_mapping_group grp; + + grp.grp_start = cp; + grp.grp_end = *grp_last_p; + grp.mark = UNVISITED; + grp.sibling = NULL; + grp.next = NULL; + groups->safe_push (grp); + + cp = grp_last_p; + } + + if (groups->length () > 0) + return groups; + else + { + delete groups; + return NULL; + } +} + +/* A pointer mapping group GRP may define a block of memory starting at some + base address, and maybe also define a firstprivate pointer or firstprivate + reference that points to that block. The return value is a node containing + the former, and the *FIRSTPRIVATE pointer is set if we have the latter. + If we define several base pointers, i.e. for a GOMP_MAP_STRUCT mapping, + return the number of consecutive chained nodes in CHAINED. */ + +static tree +omp_group_base (omp_mapping_group *grp, unsigned int *chained, + tree *firstprivate) +{ + tree node = *grp->grp_start; + + *firstprivate = NULL_TREE; + *chained = 1; + + switch (OMP_CLAUSE_MAP_KIND (node)) + { + case GOMP_MAP_TO: + case GOMP_MAP_FROM: + case GOMP_MAP_TOFROM: + case GOMP_MAP_ALWAYS_FROM: + case GOMP_MAP_ALWAYS_TO: + case GOMP_MAP_ALWAYS_TOFROM: + case GOMP_MAP_FORCE_FROM: + case GOMP_MAP_FORCE_TO: + case GOMP_MAP_FORCE_TOFROM: + case GOMP_MAP_FORCE_PRESENT: + case GOMP_MAP_ALLOC: + case GOMP_MAP_RELEASE: + case GOMP_MAP_DELETE: + case GOMP_MAP_FORCE_ALLOC: + if (node == grp->grp_end) + return node; + + node = OMP_CLAUSE_CHAIN (node); + if (node && OMP_CLAUSE_MAP_KIND (node) == GOMP_MAP_TO_PSET) + { + gcc_assert (node != grp->grp_end); + node = OMP_CLAUSE_CHAIN (node); + } + if (node) + switch (OMP_CLAUSE_MAP_KIND (node)) + { + case GOMP_MAP_POINTER: + case GOMP_MAP_FIRSTPRIVATE_POINTER: + case GOMP_MAP_FIRSTPRIVATE_REFERENCE: + case GOMP_MAP_POINTER_TO_ZERO_LENGTH_ARRAY_SECTION: + *firstprivate = OMP_CLAUSE_DECL (node); + return *grp->grp_start; + + case GOMP_MAP_ALWAYS_POINTER: + case GOMP_MAP_ATTACH_DETACH: + case GOMP_MAP_ATTACH_ZERO_LENGTH_ARRAY_SECTION: + return *grp->grp_start; + + default: + internal_error ("unexpected mapping node"); + } + else + internal_error ("unexpected mapping node"); + return error_mark_node; + + case GOMP_MAP_TO_PSET: + gcc_assert (node != grp->grp_end); + node = OMP_CLAUSE_CHAIN (node); + if (OMP_CLAUSE_MAP_KIND (node) == GOMP_MAP_ATTACH + || OMP_CLAUSE_MAP_KIND (node) == GOMP_MAP_DETACH) + return NULL_TREE; + else + internal_error ("unexpected mapping node"); + return error_mark_node; + + case GOMP_MAP_ATTACH: + case GOMP_MAP_DETACH: + node = OMP_CLAUSE_CHAIN (node); + if (!node || *grp->grp_start == grp->grp_end) + return NULL_TREE; + if (OMP_CLAUSE_MAP_KIND (node) == GOMP_MAP_FIRSTPRIVATE_POINTER + || OMP_CLAUSE_MAP_KIND (node) == GOMP_MAP_FIRSTPRIVATE_REFERENCE) + { + /* We're mapping the base pointer itself in a bare attach or detach + node. This is a side effect of how parsing works, and the mapping + will be removed anyway (at least for enter/exit data directives). + We should ignore the mapping here. FIXME. */ + return NULL_TREE; + } + else + internal_error ("unexpected mapping node"); + return error_mark_node; + + case GOMP_MAP_FORCE_DEVICEPTR: + case GOMP_MAP_DEVICE_RESIDENT: + case GOMP_MAP_LINK: + case GOMP_MAP_IF_PRESENT: + case GOMP_MAP_FIRSTPRIVATE: + case GOMP_MAP_FIRSTPRIVATE_INT: + case GOMP_MAP_USE_DEVICE_PTR: + case GOMP_MAP_ATTACH_ZERO_LENGTH_ARRAY_SECTION: + return NULL_TREE; + + case GOMP_MAP_FIRSTPRIVATE_POINTER: + case GOMP_MAP_FIRSTPRIVATE_REFERENCE: + case GOMP_MAP_POINTER: + case GOMP_MAP_ALWAYS_POINTER: + case GOMP_MAP_POINTER_TO_ZERO_LENGTH_ARRAY_SECTION: + /* These shouldn't appear by themselves. */ + if (!seen_error ()) + internal_error ("unexpected pointer mapping node"); + return error_mark_node; + + default: + gcc_unreachable (); + } + + return error_mark_node; +} + +/* Given a vector of omp_mapping_groups, build a hash table so we can look up + nodes by tree_operand_hash. */ + +static hash_map * +omp_index_mapping_groups (vec *groups) +{ + hash_map *grpmap + = new hash_map; + + omp_mapping_group *grp; + unsigned int i; + + FOR_EACH_VEC_ELT (*groups, i, grp) + { + tree fpp; + unsigned int chained; + tree node = omp_group_base (grp, &chained, &fpp); + + if (node == error_mark_node || (!node && !fpp)) + continue; + + for (unsigned j = 0; + node && j < chained; + node = OMP_CLAUSE_CHAIN (node), j++) + { + tree decl = OMP_CLAUSE_DECL (node); + + /* Sometimes we see zero-offset MEM_REF instead of INDIRECT_REF, + meaning node-hash lookups don't work. This is a workaround for + that, but ideally we should just create the INDIRECT_REF at + source instead. FIXME. */ + if (TREE_CODE (decl) == MEM_REF + && integer_zerop (TREE_OPERAND (decl, 1))) + decl = build1 (INDIRECT_REF, TREE_TYPE (decl), + TREE_OPERAND (decl, 0)); + + omp_mapping_group **prev = grpmap->get (decl); + + if (prev && *prev == grp) + /* Empty. */; + else if (prev) + { + /* Mapping the same thing twice is normally diagnosed as an error, + but can happen under some circumstances, e.g. in pr99928-16.c, + the directive: + + #pragma omp target simd reduction(+:a[:3]) \ + map(always, tofrom: a[:6]) + ... + + will result in two "a[0]" mappings (of different sizes). */ + + grp->sibling = (*prev)->sibling; + (*prev)->sibling = grp; + } + else + grpmap->put (decl, grp); + } + + if (!fpp) + continue; + + omp_mapping_group **prev = grpmap->get (fpp); + if (prev) + { + grp->sibling = (*prev)->sibling; + (*prev)->sibling = grp; + } + else + grpmap->put (fpp, grp); + } + return grpmap; +} + +/* Find the immediately-containing struct for a component ref (etc.) + expression EXPR. */ + +static tree +omp_containing_struct (tree expr) +{ + tree expr0 = expr; + + STRIP_NOPS (expr); + + tree expr1 = expr; + + /* FIXME: other types of accessors? */ + while (TREE_CODE (expr) == ARRAY_REF) + expr = TREE_OPERAND (expr, 0); + + if (TREE_CODE (expr) == COMPONENT_REF) + { + if (DECL_P (TREE_OPERAND (expr, 0)) + || TREE_CODE (TREE_OPERAND (expr, 0)) == 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))) + || TREE_CODE (TREE_OPERAND (expr, 0)) == ARRAY_REF) + expr = TREE_OPERAND (expr, 0); + else + internal_error ("unhandled component"); + } + + return (expr == expr1) ? expr0 : expr; +} + +/* Helper function for omp_tsort_mapping_groups. Returns TRUE on success, or + FALSE on error. */ + +static bool +omp_tsort_mapping_groups_1 (omp_mapping_group ***outlist, + vec *groups, + hash_map + *grpmap, + omp_mapping_group *grp) +{ + if (grp->mark == PERMANENT) + return true; + if (grp->mark == TEMPORARY) + { + fprintf (stderr, "when processing group:\n"); + debug_mapping_group (grp); + internal_error ("base pointer cycle detected"); + return false; + } + grp->mark = TEMPORARY; + + tree attaches_to = omp_get_attachment (grp); + + if (attaches_to) + { + omp_mapping_group **basep = grpmap->get (attaches_to); + + if (basep) + { + gcc_assert (*basep != grp); + for (omp_mapping_group *w = *basep; w; w = w->sibling) + if (!omp_tsort_mapping_groups_1 (outlist, groups, grpmap, w)) + return false; + } + } + + tree decl = OMP_CLAUSE_DECL (*grp->grp_start); + + while (decl) + { + tree base = omp_get_base_pointer (decl); + + if (!base) + break; + + omp_mapping_group **innerp = grpmap->get (base); + + /* We should treat whole-structure mappings as if all (pointer, in this + case) members are mapped as individual list items. Check if we have + such a whole-structure mapping, if we don't have an explicit reference + to the pointer member itself. */ + if (!innerp && TREE_CODE (base) == COMPONENT_REF) + { + base = omp_containing_struct (base); + innerp = grpmap->get (base); + + if (!innerp + && TREE_CODE (base) == MEM_REF + && integer_zerop (TREE_OPERAND (base, 1))) + { + tree ind = TREE_OPERAND (base, 0); + ind = build1 (INDIRECT_REF, TREE_TYPE (base), ind); + innerp = grpmap->get (ind); + } + } + + if (innerp && *innerp != grp) + { + for (omp_mapping_group *w = *innerp; w; w = w->sibling) + if (!omp_tsort_mapping_groups_1 (outlist, groups, grpmap, w)) + return false; + break; + } + + decl = base; + } + + grp->mark = PERMANENT; + + /* Emit grp to output list. */ + + **outlist = grp; + *outlist = &grp->next; + + return true; +} + +/* Topologically sort GROUPS, so that OMP 5.0-defined base pointers come + before mappings that use those pointers. This is an implementation of the + depth-first search algorithm, described e.g. at: + + https://en.wikipedia.org/wiki/Topological_sorting +*/ + +static omp_mapping_group * +omp_tsort_mapping_groups (vec *groups, + hash_map + *grpmap) +{ + omp_mapping_group *grp, *outlist = NULL, **cursor; + unsigned int i; + + cursor = &outlist; + + FOR_EACH_VEC_ELT (*groups, i, grp) + { + if (grp->mark != PERMANENT) + if (!omp_tsort_mapping_groups_1 (&cursor, groups, grpmap, grp)) + return NULL; + } + + 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. */ + +static omp_mapping_group * +omp_segregate_mapping_groups (omp_mapping_group *inlist) +{ + omp_mapping_group *ard_groups = NULL, *tf_groups = NULL; + omp_mapping_group **ard_tail = &ard_groups, **tf_tail = &tf_groups; + + for (omp_mapping_group *w = inlist; w;) + { + tree c = *w->grp_start; + omp_mapping_group *next = w->next; + + gcc_assert (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP); + + switch (OMP_CLAUSE_MAP_KIND (c)) + { + case GOMP_MAP_ALLOC: + case GOMP_MAP_RELEASE: + case GOMP_MAP_DELETE: + *ard_tail = w; + w->next = NULL; + ard_tail = &w->next; + break; + + default: + *tf_tail = w; + w->next = NULL; + tf_tail = &w->next; + } + + w = next; + } + + /* Now splice the lists together... */ + *tf_tail = ard_groups; + + return tf_groups; +} + +/* Given a list LIST_P containing groups of mappings given by GROUPS, reorder + those groups based on the output list of omp_tsort_mapping_groups -- + singly-linked, threaded through each element's NEXT pointer starting at + HEAD. Each list element appears exactly once in that linked list. + + Each element of GROUPS may correspond to one or several mapping nodes. + Node groups are kept together, and in the reordered list, the positions of + the original groups are reused for the positions of the reordered list. + Hence if we have e.g. + + {to ptr ptr} firstprivate {tofrom ptr} ... + ^ ^ ^ + first group non-"map" second group + + and say the second group contains a base pointer for the first so must be + moved before it, the resulting list will contain: + + {tofrom ptr} firstprivate {to ptr ptr} ... + ^ prev. second group ^ prev. first group +*/ + +static tree * +omp_reorder_mapping_groups (vec *groups, + omp_mapping_group *head, + tree *list_p) +{ + omp_mapping_group *grp; + unsigned int i; + unsigned numgroups = groups->length (); + auto_vec old_heads (numgroups); + auto_vec new_heads (numgroups); + auto_vec old_succs (numgroups); + bool map_at_start = (list_p == (*groups)[0].grp_start); + + tree *new_grp_tail = NULL; + + /* Stash the start & end nodes of each mapping group before we start + modifying the list. */ + FOR_EACH_VEC_ELT (*groups, i, grp) + { + old_heads.quick_push (*grp->grp_start); + old_succs.quick_push (OMP_CLAUSE_CHAIN (grp->grp_end)); + } + + /* And similarly, the heads of the groups in the order we want to rearrange + the list to. */ + for (omp_mapping_group *w = head; w; w = w->next) + new_heads.quick_push (*w->grp_start); + + FOR_EACH_VEC_ELT (*groups, i, grp) + { + gcc_assert (head); + + if (new_grp_tail && old_succs[i - 1] == old_heads[i]) + { + /* a {b c d} {e f g} h i j (original) + --> + a {k l m} {e f g} h i j (inserted new group on last iter) + --> + a {k l m} {n o p} h i j (this time, chain last group to new one) + ^new_grp_tail + */ + *new_grp_tail = new_heads[i]; + } + else if (new_grp_tail) + { + /* a {b c d} e {f g h} i j k (original) + --> + a {l m n} e {f g h} i j k (gap after last iter's group) + --> + a {l m n} e {o p q} h i j (chain last group to old successor) + ^new_grp_tail + */ + *new_grp_tail = old_succs[i - 1]; + } + else + { + /* The first inserted group -- point to new group, and leave end + open. + a {b c d} e f + --> + a {g h i... + */ + *grp->grp_start = new_heads[i]; + } + + new_grp_tail = &OMP_CLAUSE_CHAIN (head->grp_end); + + head = head->next; + } + + if (new_grp_tail) + *new_grp_tail = old_succs[numgroups - 1]; + + gcc_assert (!head); + + return map_at_start ? (*groups)[0].grp_start : list_p; +} /* DECL is supposed to have lastprivate semantics in the outer contexts of combined/composite constructs, starting with OCTX. @@ -9672,11 +10402,29 @@ 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) - omp_target_reorder_clauses (list_p); + /* 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)) + { + vec *groups; + groups = omp_gather_mapping_groups (list_p); + if (groups) + { + 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); + delete grpmap; + delete groups; + } + } while ((c = *list_p) != NULL) { diff --git a/gcc/omp-low.c b/gcc/omp-low.c index a63362c0511..d7d7f5310d0 100644 --- a/gcc/omp-low.c +++ b/gcc/omp-low.c @@ -1530,8 +1530,11 @@ scan_sharing_clauses (tree clauses, omp_context *ctx) { /* If this is an offloaded region, an attach operation should only exist when the pointer variable is mapped in a prior - clause. */ - if (is_gimple_omp_offloaded (ctx->stmt)) + clause. + If we had an error, we may not have attempted to sort clauses + properly, so avoid the test. */ + if (is_gimple_omp_offloaded (ctx->stmt) + && !seen_error ()) gcc_assert (maybe_lookup_decl (decl, ctx) || (is_global_var (maybe_lookup_decl_in_outer_ctx (decl, ctx)) diff --git a/gcc/testsuite/g++.dg/gomp/target-lambda-1.C b/gcc/testsuite/g++.dg/gomp/target-lambda-1.C index e5a24d7abc4..150f286e312 100644 --- a/gcc/testsuite/g++.dg/gomp/target-lambda-1.C +++ b/gcc/testsuite/g++.dg/gomp/target-lambda-1.C @@ -87,8 +87,8 @@ 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\(attach_zero_length_array_section:__closure->__iptr \[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\(attach_zero_length_array_section:_[0-9]+->ptr \[bias: 0\]\) map\(from:mapped \[len: [0-9]+\]\) map\(alloc:\*_[0-9]+ \[len: 0\]\) map\(alloc:\*_[0-9]+ \[len: 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\(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\(end\) firstprivate\(begin\) map\(to:loop \[len: [0-9]+\]\) map\(attach_zero_length_array_section:loop\.__data1 \[bias: 0\]\) map\(alloc:\*_[0-9]+ \[len: 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" } } */ -/* { dg-final { scan-tree-dump {#pragma omp target num_teams.* firstprivate\(end\) firstprivate\(begin\) map\(to:loop \[len: [0-9]+\]\) map\(attach_zero_length_array_section:loop\.__data2 \[bias: 0\]\) map\(alloc:\*_[0-9]+ \[len: 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\.__data2 \[bias: 0\]\)} "gimple" } } */ diff --git a/gcc/testsuite/g++.dg/gomp/target-this-3.C b/gcc/testsuite/g++.dg/gomp/target-this-3.C index 2755b4b58bd..bc2cc0b297d 100644 --- a/gcc/testsuite/g++.dg/gomp/target-this-3.C +++ b/gcc/testsuite/g++.dg/gomp/target-this-3.C @@ -100,6 +100,6 @@ int main (void) return 0; } -/* { dg-final { scan-tree-dump {#pragma omp target num_teams.* firstprivate\(n\) map\(tofrom:\*this \[len: [0-9]+\]\) map\(firstprivate:this \[pointer assign, bias: 0\]\) map\(alloc:\*_[0-9]+ \[pointer assign, zero-length array section, bias: 0\]\) map\(attach:this->refptr \[bias: 0\]\) map\(from:mapped \[len: [0-9]+\]\) map\(alloc:\*_[0-9+] \[len: 0\]\)} "gimple" } } */ +/* { dg-final { scan-tree-dump {#pragma omp target num_teams.* firstprivate\(n\) map\(tofrom:\*this \[len: [0-9]+\]\) map\(firstprivate: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:this->refptr \[bias: 0\]\)} "gimple" } } */ -/* { dg-final { scan-tree-dump {#pragma omp target num_teams.* firstprivate\(n\) map\(tofrom:\*this \[len: [0-9]+\]\) map\(firstprivate:this \[pointer assign, bias: 0\]\) map\(attach_zero_length_array_section:this->ptr \[bias: 0\]\) map\(from:mapped \[len: [0-9]+\]\) map\(alloc:\*_[0-9]+ \[len: 0\]\)} "gimple" } } */ +/* { dg-final { scan-tree-dump {#pragma omp target num_teams.* firstprivate\(n\) map\(tofrom:\*this \[len: [0-9]+\]\) map\(firstprivate:this \[pointer assign, bias: 0\]\) map\(from:mapped \[len: [0-9]+\]\) map\(alloc:\*_[0-9]+ \[len: 0\]\) map\(attach_zero_length_array_section:this->ptr \[bias: 0\]\)} "gimple" } } */ diff --git a/gcc/testsuite/g++.dg/gomp/target-this-4.C b/gcc/testsuite/g++.dg/gomp/target-this-4.C index 3703762f45a..8166f43ad42 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\(attach_zero_length_array_section:_[0-9]+->ptr \[bias: 0\]\) map\(from:mapped \[len: 1\]\) map\(alloc:\*_[0-9]+ \[len: 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: 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\(alloc:\*_[0-9]+ \[pointer assign, zero-length array section, bias: 0\]\) map\(attach:_[0-9]+->refptr \[bias: 0\]\) map\(from:mapped \[len: [0-9]+\]\) map\(alloc:\*_[0-9]+ \[len: 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" } } */ From patchwork Fri Oct 1 17:07:51 2021 Content-Type: text/plain; charset="utf-8" MIME-Version: 1.0 Content-Transfer-Encoding: 7bit X-Patchwork-Submitter: Julian Brown X-Patchwork-Id: 45700 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 6BF833857403 for ; Fri, 1 Oct 2021 17:10:28 +0000 (GMT) X-Original-To: gcc-patches@gcc.gnu.org Delivered-To: gcc-patches@gcc.gnu.org Received: from esa3.mentor.iphmx.com (esa3.mentor.iphmx.com [68.232.137.180]) by sourceware.org (Postfix) with ESMTPS id E086A3857401 for ; Fri, 1 Oct 2021 17:08:17 +0000 (GMT) DMARC-Filter: OpenDMARC Filter v1.4.1 sourceware.org E086A3857401 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: esD1RtjWooE3gfG5ppcv0z5O4EvtaIUxLap8r/Qcgze5qtV0pk6w0axvqXB1fXgIEHATjyD+2n jwrUerhIky9QcLdFQEFWKG3fANnpoSrWa1B2QbUbU/pxOtf/ShtCJtpvephZ5eUq8LLoIN83sn PqFKQPuklY+EZdvK+woBoW8drbQ3mlWXALJPhMBFSC0vsMWncFAyG48UEcAers5/dUq1ibtmLP /NNkypQgauJW1MntmJNvlmyxOpw92uEkoeIOS1m8zZGCEOtJMCMZ7rFPUJ+uN/0j7HkRCfVA8z lQBPZY8E1in1Dls5Lv7NJyII X-IronPort-AV: E=Sophos;i="5.85,339,1624348800"; d="scan'208";a="66548493" Received: from orw-gwy-01-in.mentorg.com ([192.94.38.165]) by esa3.mentor.iphmx.com with ESMTP; 01 Oct 2021 09:08:18 -0800 IronPort-SDR: 9MX8Ri+gwtYTDmHipz55pc1AU6WzLfY760+MkZaqxpG7q6JnKcjumu5sgLtTbvyAC7NeKHk+89 tMb0olxHj8+5IA0DYYiRbnapzX15AB+54YpKj9Y8pm6o0XvExp+T2dgS2zLYa7LuRlEDwvNFHi RukrI1+IWuQI0BQJtYw+AOccI6BO7bed4VTLH4RbcqB5mJNrzIY1JCmk6qd0RaZ1l56/gey211 DFpk6amk08JqoaYzi1J0LwVUtttQClZBcFAFrscdx46uV2WJvg8wie5k8YtUl5gzRgnkIdhZkn uvM= From: Julian Brown To: Subject: [PATCH 04/11] Remove omp_target_reorder_clauses Date: Fri, 1 Oct 2021 10:07:51 -0700 Message-ID: <8e784faf83b9973c98c052fd987d20ec04811852.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-07.mgc.mentorg.com (139.181.222.7) 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 has been split out from the previous one to avoid a confusingly-interleaved diff. The two patches should probably be committed squashed together. 2021-10-01 Julian Brown gcc/ * gimplify.c (omp_target_reorder_clauses): Delete. --- gcc/gimplify.c | 183 ------------------------------------------------- 1 file changed, 183 deletions(-) diff --git a/gcc/gimplify.c b/gcc/gimplify.c index d3346fc8d35..c10a3e8842a 100644 --- a/gcc/gimplify.c +++ b/gcc/gimplify.c @@ -8728,189 +8728,6 @@ aggregate_base_p (tree expr) return false; } -#if 0 -/* Implement OpenMP 5.x map ordering rules for target directives. There are - several rules, and with some level of ambiguity, hopefully we can at least - collect the complexity here in one place. */ - -static void -omp_target_reorder_clauses (tree *list_p) -{ - /* Collect refs to alloc/release/delete maps. */ - auto_vec ard; - tree *cp = list_p; - while (*cp != NULL_TREE) - if (OMP_CLAUSE_CODE (*cp) == OMP_CLAUSE_MAP - && (OMP_CLAUSE_MAP_KIND (*cp) == GOMP_MAP_ALLOC - || OMP_CLAUSE_MAP_KIND (*cp) == GOMP_MAP_RELEASE - || OMP_CLAUSE_MAP_KIND (*cp) == GOMP_MAP_DELETE)) - { - /* Unlink cp and push to ard. */ - tree c = *cp; - tree nc = OMP_CLAUSE_CHAIN (c); - *cp = nc; - ard.safe_push (c); - - /* Any associated pointer type maps should also move along. */ - while (*cp != NULL_TREE - && OMP_CLAUSE_CODE (*cp) == OMP_CLAUSE_MAP - && (OMP_CLAUSE_MAP_KIND (*cp) == GOMP_MAP_FIRSTPRIVATE_REFERENCE - || OMP_CLAUSE_MAP_KIND (*cp) == GOMP_MAP_FIRSTPRIVATE_POINTER - || OMP_CLAUSE_MAP_KIND (*cp) == GOMP_MAP_ATTACH_DETACH - || OMP_CLAUSE_MAP_KIND (*cp) == GOMP_MAP_POINTER - || OMP_CLAUSE_MAP_KIND (*cp) == GOMP_MAP_ALWAYS_POINTER - || OMP_CLAUSE_MAP_KIND (*cp) == GOMP_MAP_TO_PSET)) - { - c = *cp; - nc = OMP_CLAUSE_CHAIN (c); - *cp = nc; - ard.safe_push (c); - } - } - else - cp = &OMP_CLAUSE_CHAIN (*cp); - - /* Link alloc/release/delete maps to the end of list. */ - for (unsigned int i = 0; i < ard.length (); i++) - { - *cp = ard[i]; - cp = &OMP_CLAUSE_CHAIN (ard[i]); - } - *cp = NULL_TREE; - - /* OpenMP 5.0 requires that pointer variables are mapped before - its use as a base-pointer. */ - auto_vec atf; - 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 clause tree pointers. */ - 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) - atf.safe_push (cp); - } - - for (unsigned int i = 0; i < atf.length (); i++) - if (atf[i]) - { - tree *cp = atf[i]; - tree decl = OMP_CLAUSE_DECL (*cp); - if (TREE_CODE (decl) == INDIRECT_REF || TREE_CODE (decl) == MEM_REF) - { - tree base_ptr = TREE_OPERAND (decl, 0); - STRIP_TYPE_NOPS (base_ptr); - for (unsigned int j = i + 1; j < atf.length (); j++) - 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; - } - } - } - } - } -} -#endif - - enum omp_tsort_mark { UNVISITED, TEMPORARY, 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" } } */ From patchwork Fri Oct 1 17:09:04 2021 Content-Type: text/plain; charset="utf-8" MIME-Version: 1.0 Content-Transfer-Encoding: 7bit X-Patchwork-Submitter: Julian Brown X-Patchwork-Id: 45702 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 44A6E385B83F for ; Fri, 1 Oct 2021 17:11:39 +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 8CAE7385740A for ; Fri, 1 Oct 2021 17:09:25 +0000 (GMT) DMARC-Filter: OpenDMARC Filter v1.4.1 sourceware.org 8CAE7385740A 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: ktHJbcPR+/C+StzZK8ubowiLl1NgCQrPuRzzPmEMcXLDt9+kgwK7OdWuMiRnbjEyS9j4yTKkWK 0tD+DFkvVKMZH3FgUEqNmB2LIq4a2GPCZz9x8z85lqeBSQy1p+uqMoQLYEdN5Bk8uG34E5umYE 4RRazuug54gwnUXf2U41Sxt9U6GzsbH2qQQuraxYy5Lng0dC5YWjSYkgT84RItCRG1iQuaj8Vt yGpnInJywPGrgrGyU56QaHYDftwUsnvh6jg3K2yQSuHoyiettY3U5GIrKlKshHNirVy/geRwAm 58Q799E8FLQjA/7XkDtHuGoh X-IronPort-AV: E=Sophos;i="5.85,339,1624348800"; d="scan'208";a="66726631" 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: UJO5EL8VoMWCEQFJ9B62hta67LliBLsAtvyOaYFJ6wxDj6Hg4zPCkhI8zxFOB3T8VA67d+Jkic m8cycL2RYn+MFqYg5VJ0HAhIv5ZNYaIs3KgT+Sy636lhDnXtXo5frPlCi1Sc+7SDFhhJ4spgao tXU0iF+ERuvniOwpKmukosG0u6iCkeECO0xaG66bscZ9WPe2FItnoeReMjkgtRJIHfjBLXwXkW XY7/RaxZwC5k7Oxy82Nhv2p6iW+tvXgGVrYDPLKicnliWzF0WwuBsYttdMVo27yBOiN76HqgKO sAc= From: Julian Brown To: Subject: [PATCH 06/11] OpenMP: Allow array ref components for C & C++ Date: Fri, 1 Oct 2021 10:09:04 -0700 Message-ID: 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 fixes parsing for struct components that are array references in OMP clauses in both the C and C++ front ends. OK for mainline? Thanks, Julian 2021-09-29 Julian Brown gcc/c/ * c-typeck.c (c_finish_omp_clauses): Allow ARRAY_REF components. gcc/cp/ * semantics.c (finish_omp_clauses): Allow ARRAY_REF components. --- gcc/c/c-typeck.c | 3 ++- gcc/cp/semantics.c | 3 ++- 2 files changed, 4 insertions(+), 2 deletions(-) diff --git a/gcc/c/c-typeck.c b/gcc/c/c-typeck.c index e10e6aa8439..d0494cadf05 100644 --- a/gcc/c/c-typeck.c +++ b/gcc/c/c-typeck.c @@ -14815,7 +14815,8 @@ c_finish_omp_clauses (tree clauses, enum c_omp_region_type ort) { t = TREE_OPERAND (t, 0); if (TREE_CODE (t) == MEM_REF - || TREE_CODE (t) == INDIRECT_REF) + || TREE_CODE (t) == INDIRECT_REF + || TREE_CODE (t) == ARRAY_REF) { t = TREE_OPERAND (t, 0); STRIP_NOPS (t); diff --git a/gcc/cp/semantics.c b/gcc/cp/semantics.c index 6e954ca06a6..53bd8d236bb 100644 --- a/gcc/cp/semantics.c +++ b/gcc/cp/semantics.c @@ -7849,7 +7849,8 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort) if (REFERENCE_REF_P (t)) t = TREE_OPERAND (t, 0); if (TREE_CODE (t) == MEM_REF - || TREE_CODE (t) == INDIRECT_REF) + || TREE_CODE (t) == INDIRECT_REF + || TREE_CODE (t) == ARRAY_REF) { t = TREE_OPERAND (t, 0); STRIP_NOPS (t); From patchwork Fri Oct 1 17:09:05 2021 Content-Type: text/plain; charset="utf-8" MIME-Version: 1.0 Content-Transfer-Encoding: 7bit X-Patchwork-Submitter: Julian Brown X-Patchwork-Id: 45704 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 7C0AC3857027 for ; Fri, 1 Oct 2021 17:12:47 +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 1EE573857427 for ; Fri, 1 Oct 2021 17:09:26 +0000 (GMT) DMARC-Filter: OpenDMARC Filter v1.4.1 sourceware.org 1EE573857427 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: xzbYGI70A1QrrfGYvXTgS0nayFYSRheBaeZehjEBRz7ACTEXzRnpjGZlADdihNZejzMiKHpx32 I+9FNcdcWuL1TgBFZHqAOZxVp0uDSNuQixF+5LHxkMxcijAQNy7dvwFnkk1OoCya/T71Qj+S0A 1nVePJr6Asjir2IunK7sV48QplZI0QlwJJAF0Op/NDvQTwYXB2d6GxDlrwFEVsskKBIR3SjQES yo0nHjjTcJeyNIqPMzNLlRwy3u6ZYBRfRDQVQPWdML6d2N8N0e74Gy6Ijb4Uy0mxKYM2R3C+ro 2gmeVCJ0bthdx9hZbuqAtAl3 X-IronPort-AV: E=Sophos;i="5.85,339,1624348800"; d="scan'208";a="66726634" Received: from orw-gwy-01-in.mentorg.com ([192.94.38.165]) by esa4.mentor.iphmx.com with ESMTP; 01 Oct 2021 09:09:24 -0800 IronPort-SDR: SFaCIHQ0mZeow0/0cxISva/7Cv5XI/qFvctszVYVKxsGq5UhOTacL//awqFYhWJkBHL1QPP0bY uRfSIU9BRMw5NNDh1kO7S169MMNjxU2vQ1sk0hJVr1OdNnePhRomGLWwmk3kIUCQtyMy05sTtJ 7ewKgFuaa4RCVc1FrYyMTb+j5q9FN5eakcW/B+M5YObvl+1k1d4wWFW/T/ZBRXbFsYzegnxP6o OyieBlau6fF+OZKltOjJc0W7cS8TeeCoZtoBZp4t1yRAoKeGxdogrRUwEVWUItmCz/KeYyuMPI z+o= From: Julian Brown To: Subject: [PATCH 07/11] OpenMP: Fix non-zero attach/detach bias for struct dereferences Date: Fri, 1 Oct 2021 10:09:05 -0700 Message-ID: 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 fixes attach/detach operations for OpenMP that have a non-zero bias: these can occur if we have a mapping such as: #pragma omp target map(mystruct->a.b[idx].c[:arrsz]) i.e. where there is an offset between the attachment point ("mystruct" here) and the pointed-to data. (The "b" and "c" members would be array types here, not pointers themselves). In this example the difference (thus bias encoded in the attach/detach node) will be something like: (uintptr_t) &mystruct->a.b[idx].c[0] - (uintptr_t) &mystruct->a OK for mainline? Thanks, Julian 2021-09-29 Julian Brown gcc/c-family/ * c-common.h (c_omp_decompose_attachable_address): Add prototype. * c-omp.c (c_omp_decompose_attachable_address): New function. gcc/c/ * c-typeck.c (handle_omp_array_sections): Handle attach/detach for struct dereferences with non-zero bias. gcc/cp/ * semantics.c (handle_omp_array_section): Handle attach/detach for struct dereferences with non-zero bias. libgomp/ * testsuite/libgomp.c++/baseptrs-3.C: Add test (XFAILed for now). * testsuite/libgomp.c-c++-common/baseptrs-1.c: Add test. * testsuite/libgomp.c-c++-common/baseptrs-2.c: Add test. --- gcc/c-family/c-common.h | 1 + gcc/c-family/c-omp.c | 42 ++++ gcc/c/c-typeck.c | 12 +- gcc/cp/semantics.c | 14 +- libgomp/testsuite/libgomp.c++/baseptrs-3.C | 182 ++++++++++++++++++ .../libgomp.c-c++-common/baseptrs-1.c | 50 +++++ .../libgomp.c-c++-common/baseptrs-2.c | 70 +++++++ 7 files changed, 364 insertions(+), 7 deletions(-) create mode 100644 libgomp/testsuite/libgomp.c++/baseptrs-3.C create mode 100644 libgomp/testsuite/libgomp.c-c++-common/baseptrs-1.c create mode 100644 libgomp/testsuite/libgomp.c-c++-common/baseptrs-2.c diff --git a/gcc/c-family/c-common.h b/gcc/c-family/c-common.h index 849cefab882..dab2dd33573 100644 --- a/gcc/c-family/c-common.h +++ b/gcc/c-family/c-common.h @@ -1249,6 +1249,7 @@ extern tree c_omp_check_context_selector (location_t, tree); extern void c_omp_mark_declare_variant (location_t, tree, tree); extern const char *c_omp_map_clause_name (tree, bool); extern void c_omp_adjust_map_clauses (tree, bool); +extern tree c_omp_decompose_attachable_address (tree t, tree *virtbase); enum c_omp_directive_kind { C_OMP_DIR_STANDALONE, diff --git a/gcc/c-family/c-omp.c b/gcc/c-family/c-omp.c index 1f07a0a454b..fc50f57e768 100644 --- a/gcc/c-family/c-omp.c +++ b/gcc/c-family/c-omp.c @@ -3119,6 +3119,48 @@ c_omp_adjust_map_clauses (tree clauses, bool is_target) } } +tree +c_omp_decompose_attachable_address (tree t, tree *virtbase) +{ + *virtbase = t; + + /* It's already a pointer. Just use that. */ + if (POINTER_TYPE_P (TREE_TYPE (t))) + return NULL_TREE; + + /* Otherwise, look for a base pointer deeper within the expression. */ + + while (TREE_CODE (t) == COMPONENT_REF + && (TREE_CODE (TREE_OPERAND (t, 0)) == COMPONENT_REF + || TREE_CODE (TREE_OPERAND (t, 0)) == ARRAY_REF)) + { + t = TREE_OPERAND (t, 0); + while (TREE_CODE (t) == ARRAY_REF) + t = TREE_OPERAND (t, 0); + } + + + *virtbase = t; + + if (TREE_CODE (t) != COMPONENT_REF) + return NULL_TREE; + + t = TREE_OPERAND (t, 0); + + tree attach_pt = NULL_TREE; + + if ((TREE_CODE (t) == INDIRECT_REF + || TREE_CODE (t) == MEM_REF) + && TREE_CODE (TREE_TYPE (TREE_OPERAND (t, 0))) == POINTER_TYPE) + { + attach_pt = TREE_OPERAND (t, 0); + if (TREE_CODE (attach_pt) == POINTER_PLUS_EXPR) + attach_pt = TREE_OPERAND (attach_pt, 0); + } + + return attach_pt; +} + static const struct c_omp_directive omp_directives[] = { /* Keep this alphabetically sorted by the first word. Non-null second/third if any should precede null ones. */ diff --git a/gcc/c/c-typeck.c b/gcc/c/c-typeck.c index d0494cadf05..d1fd8be8e57 100644 --- a/gcc/c/c-typeck.c +++ b/gcc/c/c-typeck.c @@ -13696,9 +13696,15 @@ handle_omp_array_sections (tree c, enum c_omp_region_type ort) if (size) size = c_fully_fold (size, false, NULL); OMP_CLAUSE_SIZE (c) = size; + tree virtbase = t; + tree attach_pt + = ((ort != C_ORT_ACC) + ? c_omp_decompose_attachable_address (t, &virtbase) + : NULL_TREE); if (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_MAP || (TREE_CODE (t) == COMPONENT_REF - && TREE_CODE (TREE_TYPE (t)) == ARRAY_TYPE)) + && TREE_CODE (TREE_TYPE (t)) == ARRAY_TYPE + && !attach_pt)) return false; gcc_assert (OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_FORCE_DEVICEPTR); switch (OMP_CLAUSE_MAP_KIND (c)) @@ -13731,10 +13737,10 @@ handle_omp_array_sections (tree c, enum c_omp_region_type ort) if (OMP_CLAUSE_MAP_KIND (c2) != GOMP_MAP_FIRSTPRIVATE_POINTER && !c_mark_addressable (t)) return false; - OMP_CLAUSE_DECL (c2) = t; + OMP_CLAUSE_DECL (c2) = attach_pt ? attach_pt : t; t = build_fold_addr_expr (first); t = fold_convert_loc (OMP_CLAUSE_LOCATION (c), ptrdiff_type_node, t); - tree ptr = OMP_CLAUSE_DECL (c2); + tree ptr = virtbase; if (!POINTER_TYPE_P (TREE_TYPE (ptr))) ptr = build_fold_addr_expr (ptr); t = fold_build2_loc (OMP_CLAUSE_LOCATION (c), MINUS_EXPR, diff --git a/gcc/cp/semantics.c b/gcc/cp/semantics.c index 53bd8d236bb..a50ec0ad883 100644 --- a/gcc/cp/semantics.c +++ b/gcc/cp/semantics.c @@ -5608,9 +5608,16 @@ handle_omp_array_sections (tree c, enum c_omp_region_type ort) OMP_CLAUSE_SIZE (c) = size; if (TREE_CODE (t) == FIELD_DECL) t = finish_non_static_data_member (t, NULL_TREE, NULL_TREE); + + tree virtbase = t; + tree attach_pt + = ((ort != C_ORT_ACC) + ? c_omp_decompose_attachable_address (t, &virtbase) + : NULL_TREE); if (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_MAP || (TREE_CODE (t) == COMPONENT_REF - && TREE_CODE (TREE_TYPE (t)) == ARRAY_TYPE)) + && TREE_CODE (TREE_TYPE (t)) == ARRAY_TYPE + && !attach_pt)) return false; switch (OMP_CLAUSE_MAP_KIND (c)) { @@ -5670,12 +5677,11 @@ handle_omp_array_sections (tree c, enum c_omp_region_type ort) if (OMP_CLAUSE_MAP_KIND (c2) != GOMP_MAP_FIRSTPRIVATE_POINTER && !cxx_mark_addressable (t)) return false; - OMP_CLAUSE_DECL (c2) = t; + OMP_CLAUSE_DECL (c2) = attach_pt ? attach_pt : t; t = build_fold_addr_expr (first); t = fold_convert_loc (OMP_CLAUSE_LOCATION (c), ptrdiff_type_node, t); - tree ptr = OMP_CLAUSE_DECL (c2); - ptr = convert_from_reference (ptr); + tree ptr = convert_from_reference (virtbase); if (!INDIRECT_TYPE_P (TREE_TYPE (ptr))) ptr = build_fold_addr_expr (ptr); t = fold_build2_loc (OMP_CLAUSE_LOCATION (c), MINUS_EXPR, diff --git a/libgomp/testsuite/libgomp.c++/baseptrs-3.C b/libgomp/testsuite/libgomp.c++/baseptrs-3.C new file mode 100644 index 00000000000..cabeb7c2b7a --- /dev/null +++ b/libgomp/testsuite/libgomp.c++/baseptrs-3.C @@ -0,0 +1,182 @@ +/* { dg-xfail-if "fails to parse correctly" { *-*-* } } */ + +#include +#include +#include + +struct sa +{ + int *ptr; +}; + +struct sb +{ + int arr[10]; +}; + +struct sc +{ + sa &a; + sb &b; + sc (sa &my_a, sb &my_b) : a(my_a), b(my_b) {} +}; + +void +foo () +{ + sa my_a; + sb my_b; + + my_a.ptr = (int *) malloc (sizeof (int) * 10); + sc my_c(my_a, my_b); + + memset (my_c.a.ptr, 0, sizeof (int) * 10); + + #pragma omp target map (my_c.a.ptr, my_c.a.ptr[:10]) + { + for (int i = 0; i < 10; i++) + my_c.a.ptr[i] = i; + } + + for (int i = 0; i < 10; i++) + assert (my_c.a.ptr[i] == i); + + memset (my_c.b.arr, 0, sizeof (int) * 10); + + #pragma omp target map (my_c.b.arr[:10]) + { + for (int i = 0; i < 10; i++) + my_c.b.arr[i] = i; + } + + for (int i = 0; i < 10; i++) + assert (my_c.b.arr[i] == i); + + free (my_a.ptr); +} + +void +bar () +{ + sa my_a; + sb my_b; + + my_a.ptr = (int *) malloc (sizeof (int) * 10); + sc my_c(my_a, my_b); + sc &my_cref = my_c; + + memset (my_cref.a.ptr, 0, sizeof (int) * 10); + + #pragma omp target map (my_cref.a.ptr, my_cref.a.ptr[:10]) + { + for (int i = 0; i < 10; i++) + my_cref.a.ptr[i] = i; + } + + for (int i = 0; i < 10; i++) + assert (my_cref.a.ptr[i] == i); + + memset (my_cref.b.arr, 0, sizeof (int) * 10); + + #pragma omp target map (my_cref.b.arr[:10]) + { + for (int i = 0; i < 10; i++) + my_cref.b.arr[i] = i; + } + + for (int i = 0; i < 10; i++) + assert (my_cref.b.arr[i] == i); + + free (my_a.ptr); +} + +struct scp +{ + sa *&a; + sb *&b; + scp (sa *&my_a, sb *&my_b) : a(my_a), b(my_b) {} +}; + +void +foop () +{ + sa *my_a = new sa; + sb *my_b = new sb; + + my_a->ptr = new int[10]; + scp *my_c = new scp(my_a, my_b); + + memset (my_c->a->ptr, 0, sizeof (int) * 10); + + #pragma omp target map (my_c->a->ptr, my_c->a->ptr[:10]) + { + for (int i = 0; i < 10; i++) + my_c->a->ptr[i] = i; + } + + for (int i = 0; i < 10; i++) + assert (my_c->a->ptr[i] == i); + + memset (my_c->b->arr, 0, sizeof (int) * 10); + +/* FIXME: This currently ICEs. */ +/* #pragma omp target map (my_c->b->arr[:10]) */ + { + for (int i = 0; i < 10; i++) + my_c->b->arr[i] = i; + } + + for (int i = 0; i < 10; i++) + assert (my_c->b->arr[i] == i); + + delete[] my_a->ptr; + delete my_a; + delete my_b; +} + +void +barp () +{ + sa *my_a = new sa; + sb *my_b = new sb; + + my_a->ptr = new int[10]; + scp *my_c = new scp(my_a, my_b); + scp *&my_cref = my_c; + + memset (my_cref->a->ptr, 0, sizeof (int) * 10); + + #pragma omp target map (my_cref->a->ptr, my_cref->a->ptr[:10]) + { + for (int i = 0; i < 10; i++) + my_cref->a->ptr[i] = i; + } + + for (int i = 0; i < 10; i++) + assert (my_cref->a->ptr[i] == i); + + memset (my_cref->b->arr, 0, sizeof (int) * 10); + +/* FIXME: This currently ICEs. */ +/* #pragma omp target map (my_cref->b->arr[:10]) */ + { + for (int i = 0; i < 10; i++) + my_cref->b->arr[i] = i; + } + + for (int i = 0; i < 10; i++) + assert (my_cref->b->arr[i] == i); + + delete my_a->ptr; + delete my_a; + delete my_b; +} + +int main (int argc, char *argv[]) +{ + foo (); + bar (); + foop (); + barp (); + return 0; +} diff --git a/libgomp/testsuite/libgomp.c-c++-common/baseptrs-1.c b/libgomp/testsuite/libgomp.c-c++-common/baseptrs-1.c new file mode 100644 index 00000000000..073615625b7 --- /dev/null +++ b/libgomp/testsuite/libgomp.c-c++-common/baseptrs-1.c @@ -0,0 +1,50 @@ +#include +#include +#include +#include + +#define N 32 + +typedef struct { + int x2[10][N]; +} x1type; + +typedef struct { + x1type x1[10]; +} p2type; + +typedef struct { + p2type *p2; +} p1type; + +typedef struct { + p1type *p1; +} x0type; + +typedef struct { + x0type x0[10]; +} p0type; + +int main(int argc, char *argv[]) +{ + p0type *p0; + int k1 = 0, k2 = 0, k3 = 0, n = N; + + p0 = (p0type *) malloc (sizeof *p0); + p0->x0[0].p1 = (p1type *) malloc (sizeof *p0->x0[0].p1); + p0->x0[0].p1->p2 = (p2type *) malloc (sizeof *p0->x0[0].p1->p2); + memset (p0->x0[0].p1->p2, 0, sizeof *p0->x0[0].p1->p2); + +#pragma omp target map(tofrom: p0->x0[k1].p1->p2[k2].x1[k3].x2[4][0:n]) \ + map(to: p0->x0[k1].p1, p0->x0[k1].p1->p2) \ + map(to: p0->x0[k1].p1[0]) + { + for (int i = 0; i < n; i++) + p0->x0[k1].p1->p2[k2].x1[k3].x2[4][i] = i; + } + + for (int i = 0; i < n; i++) + assert (i == p0->x0[k1].p1->p2[k2].x1[k3].x2[4][i]); + + return 0; +} diff --git a/libgomp/testsuite/libgomp.c-c++-common/baseptrs-2.c b/libgomp/testsuite/libgomp.c-c++-common/baseptrs-2.c new file mode 100644 index 00000000000..e335d7da966 --- /dev/null +++ b/libgomp/testsuite/libgomp.c-c++-common/baseptrs-2.c @@ -0,0 +1,70 @@ +#include +#include +#include + +#define N 32 + +typedef struct { + int arr[N]; + int *ptr; +} sc; + +typedef struct { + sc *c; +} sb; + +typedef struct { + sb *b; + sc *c; +} sa; + +int main (int argc, char *argv[]) +{ + sa *p; + + p = (sa *) malloc (sizeof *p); + p->b = (sb *) malloc (sizeof *p->b); + p->b->c = (sc *) malloc (sizeof *p->b->c); + p->c = (sc *) malloc (sizeof *p->c); + p->b->c->ptr = (int *) malloc (N * sizeof (int)); + p->c->ptr = (int *) malloc (N * sizeof (int)); + + for (int i = 0; i < N; i++) + { + p->b->c->ptr[i] = 0; + p->c->ptr[i] = 0; + p->b->c->arr[i] = 0; + p->c->arr[i] = 0; + } + +#pragma omp target map(to: p->b, p->b[0], p->c, p->c[0], p->b->c, p->b->c[0]) \ + map(to: p->b->c->ptr, p->c->ptr) \ + map(tofrom: p->b->c->ptr[:N], p->c->ptr[:N]) + { + for (int i = 0; i < N; i++) + { + p->b->c->ptr[i] = i; + p->c->ptr[i] = i * 2; + } + } + +#pragma omp target map(to: p->b, p->b[0], p->b->c, p->c) \ + map(tofrom: p->c[0], p->b->c[0]) + { + for (int i = 0; i < N; i++) + { + p->b->c->arr[i] = i * 3; + p->c->arr[i] = i * 4; + } + } + + for (int i = 0; i < N; i++) + { + assert (p->b->c->ptr[i] == i); + assert (p->c->ptr[i] == i * 2); + assert (p->b->c->arr[i] == i * 3); + assert (p->c->arr[i] == i * 4); + } + + return 0; +} From patchwork Fri Oct 1 17:09:06 2021 Content-Type: text/plain; charset="utf-8" MIME-Version: 1.0 Content-Transfer-Encoding: 7bit X-Patchwork-Submitter: Julian Brown X-Patchwork-Id: 45703 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 03CA7385AC1D for ; Fri, 1 Oct 2021 17:12:18 +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 9D43C3857408 for ; Fri, 1 Oct 2021 17:09:26 +0000 (GMT) DMARC-Filter: OpenDMARC Filter v1.4.1 sourceware.org 9D43C3857408 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: ClBA5XTsN6CWroV6PF2DFn/RKnHkp5k3WFePVMNFWFvHvqubTofaS9QrBqJebGh2Ymm5VToUDM bzRaqAK1Gr7+fyj0ldQQZphVOKbStKheKnKk0XEL6mcc1JRHoyphQxXaKQ7/Arjklw5O9ECUmq uRqdQdzIlw/FzTuh8oF7+TOqf7HFcAUCFZNkecGXZSUckI47H1XgRU+Vtl+DTg0Fgx8BfMVVdI zo0u4eZDlQt/bvsCOhRcJwROLmcc38WJf2frtyPLHzC1W9BiS3AwRYAFfvec1VvAzuPmmJ1Fq8 GiNj91dSF+8tDbZzUrWWQVjB X-IronPort-AV: E=Sophos;i="5.85,339,1624348800"; d="scan'208";a="66726636" Received: from orw-gwy-01-in.mentorg.com ([192.94.38.165]) by esa4.mentor.iphmx.com with ESMTP; 01 Oct 2021 09:09:25 -0800 IronPort-SDR: OfwZiT0joizbfVkw3torQ+O/fXPhZeAACMWH/UFAvjuQ1hBjH5pTYjLm/edfatTQHo2O+2ZMwx a5DZ56ivatKogg9teEba25AfsyZtQ87NYSVDwt4fHIpYASuiyZKykE+ZBq8kJ/qHcKFQS6O20L XHZh04Ixb78LUTcb4TGWYwVC7lVP2K1VD/lwRUY9w+/WSKgp51B5gegqpXBeqNNq6DSui+aeKq ViQVkyYCv0vVg2OOaPcFu+t2TN7xQ7GA48GKCPJepUyL+DMHwZ9L790wIGUt5cVIVvc1XqRwLq Wjo= From: Julian Brown To: Subject: [PATCH 08/11] Not for committing: noisy topological sorting output Date: Fri, 1 Oct 2021 10:09:06 -0700 Message-ID: <2967aebbb316ed80326774f915b82134f0545f43.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" As a possible aid to review, this is my "printf-style" debugging cruft for the topological sorting implementation. We might want to rework this into something that emits scannable output into the gimple dump in order to write tests to make sure base pointer dependencies are being found properly, but that hasn't been done yet. This is not for committing. --- gcc/gimplify.c | 169 ++++++++++++++++++++++++++++++++++++++++++++++--- 1 file changed, 161 insertions(+), 8 deletions(-) diff --git a/gcc/gimplify.c b/gcc/gimplify.c index 31e2e4d9fe7..2ec83bf273b 100644 --- a/gcc/gimplify.c +++ b/gcc/gimplify.c @@ -70,6 +70,8 @@ along with GCC; see the file COPYING3. If not see #include "context.h" #include "tree-nested.h" +//#define NOISY_TOPOSORT + /* Hash set of poisoned variables in a bind expr. */ static hash_set *asan_poisoned_variables = NULL; @@ -8957,6 +8959,10 @@ omp_gather_mapping_groups (tree *list_p) { vec *groups = new vec (); +#ifdef NOISY_TOPOSORT + fprintf (stderr, "GATHER MAPPING GROUPS\n"); +#endif + for (tree *cp = list_p; *cp; cp = &OMP_CLAUSE_CHAIN (*cp)) { if (OMP_CLAUSE_CODE (*cp) != OMP_CLAUSE_MAP) @@ -8965,6 +8971,25 @@ omp_gather_mapping_groups (tree *list_p) tree *grp_last_p = omp_group_last (cp); omp_mapping_group grp; +#ifdef NOISY_TOPOSORT + if (cp == grp_last_p) + { + tree tmp = OMP_CLAUSE_CHAIN (*cp); + OMP_CLAUSE_CHAIN (*cp) = NULL_TREE; + fprintf (stderr, "found singleton clause:\n"); + debug_generic_expr (*cp); + OMP_CLAUSE_CHAIN (*cp) = tmp; + } + else + { + tree tmp = OMP_CLAUSE_CHAIN (*grp_last_p); + OMP_CLAUSE_CHAIN (*grp_last_p) = NULL_TREE; + fprintf (stderr, "found group:\n"); + debug_generic_expr (*cp); + OMP_CLAUSE_CHAIN (*grp_last_p) = tmp; + } +#endif + grp.grp_start = cp; grp.grp_end = *grp_last_p; grp.mark = UNVISITED; @@ -9129,14 +9154,44 @@ omp_index_mapping_groups (vec *groups) omp_mapping_group *grp; unsigned int i; +#ifdef NOISY_TOPOSORT + fprintf (stderr, "INDEX MAPPING GROUPS\n"); +#endif + FOR_EACH_VEC_ELT (*groups, i, grp) { +#ifdef NOISY_TOPOSORT + debug_mapping_group (grp); +#endif + tree fpp; unsigned int chained; tree node = omp_group_base (grp, &chained, &fpp); if (node == error_mark_node || (!node && !fpp)) - continue; + { +#ifdef NOISY_TOPOSORT + fprintf (stderr, " -- NULL base, not indexing.\n"); +#endif + continue; + } + +#ifdef NOISY_TOPOSORT + if (node) + { + fprintf (stderr, "base%s: ", chained > 1 ? " list" : ""); + + tree walk = node; + for (unsigned j = 0; j < chained; walk = OMP_CLAUSE_CHAIN (walk), j++) + debug_generic_expr (OMP_CLAUSE_DECL (walk)); + } + + if (fpp) + { + fprintf (stderr, "firstprivate pointer/reference: "); + debug_generic_expr (fpp); + } +#endif for (unsigned j = 0; node && j < chained; @@ -9156,7 +9211,11 @@ omp_index_mapping_groups (vec *groups) omp_mapping_group **prev = grpmap->get (decl); if (prev && *prev == grp) - /* Empty. */; + { +#ifdef NOISY_TOPOSORT + fprintf (stderr, " -- same node\n"); +#endif + } else if (prev) { /* Mapping the same thing twice is normally diagnosed as an error, @@ -9171,9 +9230,17 @@ omp_index_mapping_groups (vec *groups) grp->sibling = (*prev)->sibling; (*prev)->sibling = grp; +#ifdef NOISY_TOPOSORT + fprintf (stderr, " -- index as sibling\n"); +#endif } else - grpmap->put (decl, grp); + { +#ifdef NOISY_TOPOSORT + fprintf (stderr, " -- index as new decl\n"); +#endif + grpmap->put (decl, grp); + } } if (!fpp) @@ -9184,9 +9251,17 @@ omp_index_mapping_groups (vec *groups) { grp->sibling = (*prev)->sibling; (*prev)->sibling = grp; +#ifdef NOISY_TOPOSORT + fprintf (stderr, " -- index fpp as sibling\n"); +#endif } else - grpmap->put (fpp, grp); + { +#ifdef NOISY_TOPOSORT + fprintf (stderr, " -- index fpp as new decl\n"); +#endif + grpmap->put (fpp, grp); + } } return grpmap; } @@ -9233,6 +9308,11 @@ omp_tsort_mapping_groups_1 (omp_mapping_group ***outlist, *grpmap, omp_mapping_group *grp) { +#ifdef NOISY_TOPOSORT + fprintf (stderr, "processing node/group:\n"); + debug_mapping_group (grp); +#endif + if (grp->mark == PERMANENT) return true; if (grp->mark == TEMPORARY) @@ -9253,11 +9333,26 @@ omp_tsort_mapping_groups_1 (omp_mapping_group ***outlist, if (basep) { gcc_assert (*basep != grp); +#ifdef NOISY_TOPOSORT + fprintf (stderr, "has attachment to:\n"); + debug_mapping_group (*basep); +#endif for (omp_mapping_group *w = *basep; w; w = w->sibling) if (!omp_tsort_mapping_groups_1 (outlist, groups, grpmap, w)) return false; } +#ifdef NOISY_TOPOSORT + else + { + fprintf (stderr, "can't find base for attachment, tried:\n"); + debug_generic_expr (attaches_to); + } +#endif } +#ifdef NOISY_TOPOSORT + else + fprintf (stderr, "doesn't attach to anything\n"); +#endif tree decl = OMP_CLAUSE_DECL (*grp->grp_start); @@ -9266,7 +9361,13 @@ omp_tsort_mapping_groups_1 (omp_mapping_group ***outlist, tree base = omp_get_base_pointer (decl); if (!base) - break; + { +#ifdef NOISY_TOPOSORT + fprintf (stderr, "no base pointer for decl:\n"); + debug_generic_expr (decl); +#endif + break; + } omp_mapping_group **innerp = grpmap->get (base); @@ -9283,6 +9384,9 @@ omp_tsort_mapping_groups_1 (omp_mapping_group ***outlist, && TREE_CODE (base) == MEM_REF && integer_zerop (TREE_OPERAND (base, 1))) { +#ifdef NOISY_TOPOSORT + fprintf (stderr, "we have a mem_ref, retry as indirect_ref\n"); +#endif tree ind = TREE_OPERAND (base, 0); ind = build1 (INDIRECT_REF, TREE_TYPE (base), ind); innerp = grpmap->get (ind); @@ -9291,12 +9395,22 @@ omp_tsort_mapping_groups_1 (omp_mapping_group ***outlist, if (innerp && *innerp != grp) { +#ifdef NOISY_TOPOSORT + fprintf (stderr, "base pointer of node is mapped too:\n"); + debug_generic_expr (base); +#endif for (omp_mapping_group *w = *innerp; w; w = w->sibling) if (!omp_tsort_mapping_groups_1 (outlist, groups, grpmap, w)) return false; break; } - +#ifdef NOISY_TOPOSORT + else + { + fprintf (stderr, "base pointer of node is not mapped:\n"); + debug_generic_expr (base); + } +#endif decl = base; } @@ -9307,6 +9421,16 @@ omp_tsort_mapping_groups_1 (omp_mapping_group ***outlist, **outlist = grp; *outlist = &grp->next; +#ifdef NOISY_TOPOSORT + fprintf (stderr, "+++ EMIT NODE/GROUP:\n"); + { + tree tmp = OMP_CLAUSE_CHAIN (grp->grp_end); + OMP_CLAUSE_CHAIN (grp->grp_end) = NULL_TREE; + debug_generic_expr (*grp->grp_start); + OMP_CLAUSE_CHAIN (grp->grp_end) = tmp; + } +#endif + return true; } @@ -9327,6 +9451,10 @@ omp_tsort_mapping_groups (vec *groups, cursor = &outlist; +#ifdef NOISY_TOPOSORT + fprintf (stderr, "TOPOLOGICALLY SORT MAPPING GROUPS\n"); +#endif + FOR_EACH_VEC_ELT (*groups, i, grp) { if (grp->mark != PERMANENT) @@ -9348,6 +9476,10 @@ omp_segregate_mapping_groups (omp_mapping_group *inlist) omp_mapping_group *ard_groups = NULL, *tf_groups = NULL; omp_mapping_group **ard_tail = &ard_groups, **tf_tail = &tf_groups; +#ifdef NOISY_TOPOSORT + fprintf (stderr, "SEGREGATE MAPPING GROUPS\n"); +#endif + for (omp_mapping_group *w = inlist; w;) { tree c = *w->grp_start; @@ -9452,7 +9584,7 @@ omp_reorder_mapping_groups (vec *groups, --> a {l m n} e {o p q} h i j (chain last group to old successor) ^new_grp_tail - */ + */ *new_grp_tail = old_succs[i - 1]; } else @@ -10285,9 +10417,30 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p, grpmap = omp_index_mapping_groups (groups); outlist = omp_tsort_mapping_groups (groups, grpmap); +#ifdef NOISY_TOPOSORT + if (!outlist) + goto failure; + fprintf (stderr, "TOPO-SORTED ORDER:\n"); + for (omp_mapping_group *w = outlist; w; w = w->next) + debug_mapping_group (w); +#endif outlist = omp_segregate_mapping_groups (outlist); +#ifdef NOISY_TOPOSORT + fprintf (stderr, "FINAL ORDER:\n"); + for (omp_mapping_group *w = outlist; w; w = w->next) + debug_mapping_group (w); +#endif list_p = omp_reorder_mapping_groups (groups, outlist, list_p); - +#ifdef NOISY_TOPOSORT + fprintf (stderr, "REORDERED CLAUSES:\n"); + for (tree *w = list_p; *w; w = &OMP_CLAUSE_CHAIN (*w)) + { + tree tmp = OMP_CLAUSE_CHAIN (*w); + OMP_CLAUSE_CHAIN (*w) = NULL_TREE; + debug_generic_expr (*w); + OMP_CLAUSE_CHAIN (*w) = tmp; + } +#endif failure: delete grpmap; delete groups; From patchwork Fri Oct 1 17:09:07 2021 Content-Type: text/plain; charset="utf-8" MIME-Version: 1.0 Content-Transfer-Encoding: 7bit X-Patchwork-Submitter: Julian Brown X-Patchwork-Id: 45705 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 BCDFE385701A for ; Fri, 1 Oct 2021 17:13:24 +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 35EAD3857435 for ; Fri, 1 Oct 2021 17:09:29 +0000 (GMT) DMARC-Filter: OpenDMARC Filter v1.4.1 sourceware.org 35EAD3857435 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: l8whFU6vdXzL9Cp/NXWy80jylOJ3YxM5f/EkuNYBru3qaKkgHzd+iJZ/92JwSypCadGblXBJ3p wAY6Z87ETABOpTX9MWdWAREnWlbVOPjEf5z4bL4pZN+khOLkGO1gyPY2r59PX6M7I4WeUP15nd G9KhhflNEkUV5uC2tllIpgm7ZqLqAaz8ahC8AUfIoD+8JMnUm7DgKrXhNeR+uC2vQs9WUeELhM ykJeKdhx0chvnvFS60OKaNGEGN81Tx9Z2d100KI157dQ1pSD2NpIXrXUl6YhEVkGLqAcTMsXAE yPRojL0+bjPovzAtZ2ceBYbh X-IronPort-AV: E=Sophos;i="5.85,339,1624348800"; d="scan'208";a="66726637" Received: from orw-gwy-01-in.mentorg.com ([192.94.38.165]) by esa4.mentor.iphmx.com with ESMTP; 01 Oct 2021 09:09:28 -0800 IronPort-SDR: yOC50vmAy2scmu/VuiTsYQOGVrNBm+0obwCOSFz45ruakqK/cjc09e8a/rsdynjkIKUKJSeJRm ibvoTSCIjj4bvxaxxm0vFe0qZ5BSWJNspeoA6HKynF/qUoE07UeBRwdWZ+Gpd6WYG1Dg2AFAoU hwCGMxruBNyrPCxYZC2mgly5+pzkZoureVfJNRiGu/2ZC2wugfvmzzm516UY030Txe03xY/Syh arLUf0/AoxEzskSFyTYOhMFi1LsrKknM2ytt+yh2P99waAoWOER5dbH31pLus266kbYMiGVfJ8 Wbs= From: Julian Brown To: Subject: [PATCH 09/11] Not for committing: noisy sibling-list handling output Date: Fri, 1 Oct 2021 10:09:07 -0700 Message-ID: <2584f870ad013b62b92563b0f8edf9fa4348ac86.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" As a possible aid to review, this is my "printf-style" debugging cruft for the sibling list handling hoist/rework. It's not meant for committing. --- gcc/gimplify.c | 131 +++++++++++++++++++++++++++++++++++++++++++++++++ 1 file changed, 131 insertions(+) diff --git a/gcc/gimplify.c b/gcc/gimplify.c index 2ec83bf273b..ffb6eda5490 100644 --- a/gcc/gimplify.c +++ b/gcc/gimplify.c @@ -71,6 +71,7 @@ along with GCC; see the file COPYING3. If not see #include "tree-nested.h" //#define NOISY_TOPOSORT +//#define NOISY_SIBLING_LISTS /* Hash set of poisoned variables in a bind expr. */ static hash_set *asan_poisoned_variables = NULL; @@ -9895,6 +9896,11 @@ build_struct_group (enum omp_region_type region_type, enum tree_code code, bool openmp = !(region_type & ORT_ACC); tree *continue_at = NULL; +#ifdef NOISY_SIBLING_LISTS + fprintf (stderr, "DECL starts out as:\n"); + debug_generic_expr (ocd); +#endif + while (TREE_CODE (ocd) == ARRAY_REF) ocd = TREE_OPERAND (ocd, 0); @@ -9903,6 +9909,11 @@ build_struct_group (enum omp_region_type region_type, enum tree_code code, tree base = extract_base_bit_offset (ocd, &cbitpos, &coffset); +#ifdef NOISY_SIBLING_LISTS + fprintf (stderr, "BASE after extraction is (%p):\n", (void *) base); + debug_generic_expr (base); +#endif + 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) @@ -9917,6 +9928,25 @@ build_struct_group (enum omp_region_type region_type, enum tree_code code, if (openmp && attach_detach) return NULL; +#ifdef NOISY_SIBLING_LISTS + if (struct_map_to_clause) + { + fprintf (stderr, "s_m_t_c->get (base) = "); + debug_generic_expr (base); + tree *r = struct_map_to_clause->get (base); + fprintf (stderr, "returns: "); + if (r) + { + tree tmp = OMP_CLAUSE_CHAIN (*r); + OMP_CLAUSE_CHAIN (*r) = NULL_TREE; + debug_generic_expr (*r); + OMP_CLAUSE_CHAIN (*r) = tmp; + } + else + fprintf (stderr, "(nothing)\n"); + } +#endif + if (!struct_map_to_clause || struct_map_to_clause->get (base) == NULL) { tree l = build_omp_clause (OMP_CLAUSE_LOCATION (grp_end), OMP_CLAUSE_MAP); @@ -10026,6 +10056,11 @@ build_struct_group (enum omp_region_type region_type, enum tree_code code, { tree *osc = struct_map_to_clause->get (base); tree *sc = NULL, *scp = NULL; +#ifdef NOISY_SIBLING_LISTS + fprintf (stderr, "looked up osc %p for decl (%p)\n", (void *) osc, + (void *) base); + debug_generic_expr (base); +#endif sc = &OMP_CLAUSE_CHAIN (*osc); /* The struct mapping might be immediately followed by a FIRSTPRIVATE_POINTER and/or FIRSTPRIVATE_REFERENCE -- if it's an @@ -10098,6 +10133,17 @@ build_struct_group (enum omp_region_type region_type, enum tree_code code, return NULL; } } +#ifdef NOISY_SIBLING_LISTS + if (known_eq (coffset, offset) && known_eq (cbitpos, bitpos)) + { + fprintf (stderr, "duplicate offset!\n"); + tree o1 = OMP_CLAUSE_DECL (*sc); + tree o2 = OMP_CLAUSE_DECL (grp_end); + debug_generic_expr (o1); + debug_generic_expr (o2); + } + else +#endif if (maybe_lt (coffset, offset) || (known_eq (coffset, offset) && maybe_lt (cbitpos, bitpos))) @@ -10174,6 +10220,13 @@ build_struct_group (enum omp_region_type region_type, enum tree_code code, = cl ? move_concat_nodes_after (cl, tail_chain, grp_start_p, grp_end, sc) : move_nodes_after (grp_start_p, grp_end, sc); +#ifdef NOISY_SIBLING_LISTS + if (continue_at) + { + fprintf (stderr, "continue at (1):\n"); + debug_generic_expr (*continue_at); + } +#endif } else if (*sc != grp_end) { @@ -10187,6 +10240,10 @@ build_struct_group (enum omp_region_type region_type, enum tree_code code, the correct position in the struct component list, which in this case is just SC. */ move_node_after (grp_end, grp_start_p, sc); +#ifdef NOISY_SIBLING_LISTS + fprintf (stderr, "continue at (2):\n"); + debug_generic_expr (*continue_at); +#endif } } return continue_at; @@ -10218,6 +10275,16 @@ omp_build_struct_sibling_lists (enum tree_code code, new_next = NULL; +#ifdef NOISY_SIBLING_LISTS + { + tree *tmp = grp->grp_start; + grp->grp_start = grp_start_p; + fprintf (stderr, "processing group %u:\n", i); + debug_mapping_group (grp); + grp->grp_start = tmp; + } +#endif + if (DECL_P (decl)) continue; @@ -10252,6 +10319,11 @@ omp_build_struct_sibling_lists (enum tree_code code, != GOMP_MAP_POINTER_TO_ZERO_LENGTH_ARRAY_SECTION)) decl = TREE_OPERAND (decl, 0); +#ifdef NOISY_SIBLING_LISTS + fprintf (stderr, "using base pointer/decl:\n"); + debug_generic_expr (decl); +#endif + STRIP_NOPS (decl); if (TREE_CODE (decl) != COMPONENT_REF) @@ -10260,6 +10332,15 @@ omp_build_struct_sibling_lists (enum tree_code code, omp_mapping_group **wholestruct = NULL; tree wsdecl = omp_containing_struct (OMP_CLAUSE_DECL (c)); +#ifdef NOISY_SIBLING_LISTS + fprintf (stderr, "whole-struct decl is (%s original)\n", + wsdecl == OMP_CLAUSE_DECL (c) ? "same as" : "different from"); + + debug_generic_expr (wsdecl); + fprintf (stderr, "orig was:\n"); + debug_generic_expr (OMP_CLAUSE_DECL (c)); +#endif + if (!(region_type & ORT_ACC) && wsdecl != OMP_CLAUSE_DECL (c)) { wholestruct = grpmap->get (wsdecl); @@ -10275,16 +10356,27 @@ omp_build_struct_sibling_lists (enum tree_code code, if (wholestruct) { +#ifdef NOISY_SIBLING_LISTS + fprintf (stderr, "it looks like the whole struct is mapped by:\n"); + debug_mapping_group (*wholestruct); +#endif if (*grp_start_p == grp_end) { /* Remove the whole of this mapping -- redundant. */ +#ifdef NOISY_SIBLING_LISTS + fprintf (stderr, "removing this group\n"); +#endif new_next = grp_start_p; *grp_start_p = OMP_CLAUSE_CHAIN (grp_end); } continue; } +#ifdef NOISY_SIBLING_LISTS + else + fprintf (stderr, "whole struct is not mapped\n"); +#endif if (OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_TO_PSET && OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_ATTACH @@ -10360,6 +10452,11 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p, int handled_depend_iterators = -1; int nowait = -1; +#ifdef NOISY_SIBLING_LISTS + fprintf (stderr, "GIMPLIFY_SCAN_OMP_CLAUSES, list:\n"); + debug_generic_expr (*list_p); +#endif + ctx = new_omp_context (region_type); ctx->code = code; outer_ctx = ctx->outer_context; @@ -10399,7 +10496,14 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p, hash_map *grpmap; grpmap = omp_index_mapping_groups (groups); +#ifdef NOISY_SIBLING_LISTS + fprintf (stderr, "BUILDING STRUCT SIBLING LISTS\n"); +#endif omp_build_struct_sibling_lists (code, region_type, groups, grpmap); +#ifdef NOISY_SIBLING_LISTS + fprintf (stderr, "result:\n"); + debug_generic_expr (*list_p); +#endif omp_mapping_group *outlist = NULL; @@ -10455,6 +10559,9 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p, hash_map *grpmap; grpmap = omp_index_mapping_groups (groups); +#ifdef NOISY_SIBLING_LISTS + fprintf (stderr, "BUILDING STRUCT SIBLING LISTS\n"); +#endif omp_build_struct_sibling_lists (code, region_type, groups, grpmap); delete groups; @@ -10775,6 +10882,22 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p, case OMP_CLAUSE_MAP: decl = OMP_CLAUSE_DECL (c); + +#ifdef NOISY_SIBLING_LISTS + { + fprintf (stderr, "gimplify_scan_omp_clauses processing: "); + debug_generic_expr (c); + if (DECL_P (decl)) + { + splay_tree_node n + = splay_tree_lookup (ctx->variables, (splay_tree_key) decl); + fprintf (stderr, "decl has node: %p\n", (void *) n); + if (n) + fprintf (stderr, "flags are: %x\n", n->value); + } + } +#endif + if (error_operand_p (decl)) remove = true; switch (code) @@ -10870,6 +10993,10 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p, if (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_STRUCT) { tree base = strip_components_and_deref (decl); +#ifdef NOISY_SIBLING_LISTS + fprintf (stderr, "struct, base="); + debug_generic_expr (base); +#endif if (DECL_P (base)) { decl = base; @@ -10951,6 +11078,10 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p, if (TREE_CODE (cref) == COMPONENT_REF) { +#ifdef NOISY_SIBLING_LISTS + fprintf (stderr, "we see a component_ref for:\n"); + debug_generic_expr (c); +#endif tree base = cref; while (base && !DECL_P (base)) { From patchwork Fri Oct 1 17:10:08 2021 Content-Type: text/plain; charset="utf-8" MIME-Version: 1.0 Content-Transfer-Encoding: 7bit X-Patchwork-Submitter: Julian Brown X-Patchwork-Id: 45706 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 2184D385AC0A for ; Fri, 1 Oct 2021 17:13:54 +0000 (GMT) X-Original-To: gcc-patches@gcc.gnu.org Delivered-To: gcc-patches@gcc.gnu.org Received: from esa2.mentor.iphmx.com (esa2.mentor.iphmx.com [68.232.141.98]) by sourceware.org (Postfix) with ESMTPS id AFAF63857818 for ; Fri, 1 Oct 2021 17:10:18 +0000 (GMT) DMARC-Filter: OpenDMARC Filter v1.4.1 sourceware.org AFAF63857818 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: EN7KfAcIjhr4A08icnEUcXlMlzHFu2gil2Kc+j1qc1rOSsRubp75FWZ9J1Yu4Dp36bB96sKgzU 4E1UOsD+vu4z0zJpGSynEKAFFcs3Q+KGlcSQjzNdUibjgKtN6JG4oU1AWxiqLRQd8pcnO/U9kR pjyTXkFDE8psDr+evBkLSN3W46RrVH92hUnEuwjZ8ETmSHEfz1YxrTvNToHMpH7h/cnfNiySsw PmGWsDo1v7qLNIKVYlgJkf5Ao6MU8MCpgxJrcHX+/XKTfy/qcEE7gr5aRKAG/4lgGGQRnZe06U r0wuIFmI3nyezOGm0W0htnV9 X-IronPort-AV: E=Sophos;i="5.85,339,1624348800"; d="scan'208";a="66622521" Received: from orw-gwy-01-in.mentorg.com ([192.94.38.165]) by esa2.mentor.iphmx.com with ESMTP; 01 Oct 2021 09:10:17 -0800 IronPort-SDR: D0PLrSyaknZbAU3ZCjnggqJGJWn/DdW4buZCdYNyg2tSuQukLnyIR8wDu/fBYz4WG61ya0Nw4M VFAKoPY/x+nQoryOzNuj1KExfhdlma+uMPuzSG3yrDMwWGzaxc0TSAY/RMRllpn8NPEef415fB DAajC+UHIJ99HWaR2P5fsuEBHSX0xA5gno1vCq9WlGxLjVJ+gPchNnB4zKQELiWSWuw1v92JWr aCPbvPxSLw9d3FJAfWKZR6CQg/Nql9wpm0hTsfgbCUM+Rxp2XY8JZ1E3VH3douQd+r+QdoZCFF aYM= From: Julian Brown To: Subject: [PATCH 10/11] Not for committing: noisy mapping-group taxonomy Date: Fri, 1 Oct 2021 10:10:08 -0700 Message-ID: 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-02.mgc.mentorg.com (139.181.222.2) 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" As a possible aid to review, this is code that can be used to enumerate all the mapping group forms currently in use across the GCC/libgomp testsuites for OpenMP/OpenACC. These groups have been added somewhat organically, so there might be a couple of surprises: see e.g. the patch following this one. It's not meant for committing. --- gcc/gimplify.c | 327 +++++++++++++++++++++++++++++++++++++++++++++++++ 1 file changed, 327 insertions(+) diff --git a/gcc/gimplify.c b/gcc/gimplify.c index ffb6eda5490..d9fda21413d 100644 --- a/gcc/gimplify.c +++ b/gcc/gimplify.c @@ -72,6 +72,7 @@ along with GCC; see the file COPYING3. If not see //#define NOISY_TOPOSORT //#define NOISY_SIBLING_LISTS +//#define NOISY_TAXONOMY /* Hash set of poisoned variables in a bind expr. */ static hash_set *asan_poisoned_variables = NULL; @@ -9010,6 +9011,326 @@ omp_gather_mapping_groups (tree *list_p) } } +#ifdef NOISY_TAXONOMY + +static void +omp_mapping_group_taxonomy (vec *groups) +{ + int num = 0; + + for (auto &it : *groups) + { + tree node, grp_start = *it.grp_start, grp_end = it.grp_end; + gomp_map_kind kind0 = OMP_CLAUSE_MAP_KIND (grp_start), kind1, kind2, + kind3; + int count = 1; + node = grp_start; + if (node != grp_end) + { + node = OMP_CLAUSE_CHAIN (node); + kind1 = OMP_CLAUSE_MAP_KIND (node); + count++; + if (node != grp_end) + { + node = OMP_CLAUSE_CHAIN (node); + kind2 = OMP_CLAUSE_MAP_KIND (node); + count++; + if (node != grp_end) + { + node = OMP_CLAUSE_CHAIN (node); + kind3 = OMP_CLAUSE_MAP_KIND (node); + count++; + gcc_assert (node == grp_end); + } + } + } + + fprintf (stderr, "group %d: ", num); + + switch (count) + { + case 1: + if (kind0 == GOMP_MAP_TO + || kind0 == GOMP_MAP_FROM + || kind0 == GOMP_MAP_TOFROM) + fprintf (stderr, "scalar to/from\n"); + else if (kind0 == GOMP_MAP_ALLOC) + fprintf (stderr, "alloc\n"); + else if (kind0 == GOMP_MAP_POINTER) + fprintf (stderr, "pointer (by itself)\n"); + else if (kind0 == GOMP_MAP_TO_PSET) + fprintf (stderr, "map-to-pset (by itself)\n"); + else if (kind0 == GOMP_MAP_FORCE_PRESENT) + fprintf (stderr, "force present\n"); + else if (kind0 == GOMP_MAP_DELETE) + fprintf (stderr, "delete\n"); + else if (kind0 == GOMP_MAP_FORCE_DEVICEPTR) + fprintf (stderr, "force deviceptr\n"); + else if (kind0 == GOMP_MAP_DEVICE_RESIDENT) + fprintf (stderr, "device resident\n"); + else if (kind0 == GOMP_MAP_LINK) + fprintf (stderr, "link\n"); + else if (kind0 == GOMP_MAP_IF_PRESENT) + fprintf (stderr, "if present\n"); + else if (kind0 == GOMP_MAP_FIRSTPRIVATE) + fprintf (stderr, "firstprivate (by itself)\n"); + else if (kind0 == GOMP_MAP_FIRSTPRIVATE_INT) + fprintf (stderr, "firstprivate_int (by itself)\n"); + else if (kind0 == GOMP_MAP_USE_DEVICE_PTR) + fprintf (stderr, "use device ptr\n"); + else if (kind0 == GOMP_MAP_ZERO_LEN_ARRAY_SECTION) + fprintf (stderr, "zero-length array section (by itself)\n"); + else if (kind0 == GOMP_MAP_FORCE_ALLOC) + fprintf (stderr, "force alloc\n"); + else if (kind0 == GOMP_MAP_FORCE_TO + || kind0 == GOMP_MAP_FORCE_FROM + || kind0 == GOMP_MAP_FORCE_TOFROM) + fprintf (stderr, "force to/from (scalar)\n"); + else if (kind0 == GOMP_MAP_USE_DEVICE_PTR_IF_PRESENT) + fprintf (stderr, "use device ptr if present\n"); + else if (kind0 == GOMP_MAP_ALWAYS_TO + || kind0 == GOMP_MAP_ALWAYS_FROM + || kind0 == GOMP_MAP_ALWAYS_TOFROM) + fprintf (stderr, "always to/from (scalar)\n"); + else if (kind0 == GOMP_MAP_STRUCT) + fprintf (stderr, "struct\n"); + else if (kind0 == GOMP_MAP_ALWAYS_POINTER) + fprintf (stderr, "always pointer (by itself)\n"); + else if (kind0 == GOMP_MAP_POINTER_TO_ZERO_LENGTH_ARRAY_SECTION) + fprintf (stderr, "ptr to 0-length array section (by itself)\n"); + else if (kind0 == GOMP_MAP_RELEASE) + fprintf (stderr, "release\n"); + else if (kind0 == GOMP_MAP_ATTACH) + fprintf (stderr, "attach\n"); + else if (kind0 == GOMP_MAP_DETACH) + fprintf (stderr, "detach\n"); + else if (kind0 == GOMP_MAP_FORCE_DETACH) + fprintf (stderr, "force detach\n"); + else if (kind0 == GOMP_MAP_ATTACH_ZERO_LENGTH_ARRAY_SECTION) + fprintf (stderr, "attach 0-length array section\n"); + else if (kind0 == GOMP_MAP_FIRSTPRIVATE_POINTER) + fprintf (stderr, "firstprivate ptr (by itself)\n"); + else if (kind0 == GOMP_MAP_FIRSTPRIVATE_REFERENCE) + fprintf (stderr, "firstprivate ref (by itself)\n"); + else if (kind0 == GOMP_MAP_ATTACH_DETACH) + fprintf (stderr, "attach/detach (by itself)\n"); + else + fprintf (stderr, "unknown code %d\n", (int) kind0); + break; + + case 2: + if (kind0 == GOMP_MAP_TO + || kind0 == GOMP_MAP_FROM + || kind0 == GOMP_MAP_TOFROM) + { + if (kind1 == GOMP_MAP_POINTER) + fprintf (stderr, "to/from, pointer\n"); + else if (kind1 == GOMP_MAP_ALWAYS_POINTER) + fprintf (stderr, "to/from, always pointer\n"); + else if (kind1 == GOMP_MAP_ATTACH_DETACH) + fprintf (stderr, "to/from, attach/detach\n"); + else if (kind1 == GOMP_MAP_FIRSTPRIVATE_POINTER) + fprintf (stderr, "to/from, firstprivate pointer\n"); + else if (kind1 == GOMP_MAP_FIRSTPRIVATE_REFERENCE) + fprintf (stderr, "to/from, firstprivate reference\n"); + else + fprintf (stderr, "to/from, unknown code %d\n", (int) kind1); + } + else if (kind0 == GOMP_MAP_ALWAYS_FROM + || kind0 == GOMP_MAP_ALWAYS_TO + || kind0 == GOMP_MAP_ALWAYS_TOFROM) + { + if (kind1 == GOMP_MAP_FIRSTPRIVATE_POINTER) + fprintf (stderr, "always to/from, firstprivate pointer\n"); + else if (kind1 == GOMP_MAP_FIRSTPRIVATE_REFERENCE) + fprintf (stderr, "always to/from, firstprivate reference\n"); + else if (kind1 == GOMP_MAP_ATTACH_DETACH) + fprintf (stderr, "always to/from, attach/detach\n"); + else if (kind1 == GOMP_MAP_ALWAYS_POINTER) + fprintf (stderr, "always to/from, always pointer\n"); + else + fprintf (stderr, "always to/from, unknown code %d\n", + (int) kind1); + } + else if (kind0 == GOMP_MAP_FORCE_TO + || kind0 == GOMP_MAP_FORCE_FROM + || kind0 == GOMP_MAP_FORCE_TOFROM) + { + if (kind1 == GOMP_MAP_POINTER) + fprintf (stderr, "force to/from, pointer\n"); + else + fprintf (stderr, "force to/from, unknown code %d\n", + (int) kind1); + } + else if (kind0 == GOMP_MAP_FORCE_PRESENT) + { + if (kind1 == GOMP_MAP_POINTER) + fprintf (stderr, "force present, pointer\n"); + else if (kind1 == GOMP_MAP_FIRSTPRIVATE_POINTER) + fprintf (stderr, "force present, firstprivate pointer\n"); + else + fprintf (stderr, "force present, unknown code %d\n", + (int) kind1); + } + else if (kind0 == GOMP_MAP_ALLOC) + { + if (kind1 == GOMP_MAP_POINTER) + fprintf (stderr, "alloc, pointer\n"); + else if (kind1 == GOMP_MAP_ALWAYS_POINTER) + fprintf (stderr, "alloc, always pointer\n"); + else if (kind1 == GOMP_MAP_ATTACH_DETACH) + fprintf (stderr, "alloc, attach/detach\n"); + else if (kind1 == GOMP_MAP_FIRSTPRIVATE_POINTER) + fprintf (stderr, "alloc, firstprivate pointer\n"); + else if (kind1 == GOMP_MAP_FIRSTPRIVATE_REFERENCE) + fprintf (stderr, "alloc, firstprivate reference\n"); + else + fprintf (stderr, "alloc, unknown code %d\n", (int) kind1); + } + else if (kind0 == GOMP_MAP_RELEASE) + { + if (kind1 == GOMP_MAP_FIRSTPRIVATE_REFERENCE) + fprintf (stderr, "release, firstprivate reference\n"); + else if (kind1 == GOMP_MAP_ATTACH_DETACH) + fprintf (stderr, "release, attach/detach\n"); + else if (kind1 == GOMP_MAP_ALWAYS_POINTER) + fprintf (stderr, "release, always pointer\n"); + else if (kind1 == GOMP_MAP_POINTER) + fprintf (stderr, "release, pointer\n"); + else + fprintf (stderr, "release, unknown code %d\n", (int) kind1); + } + else if (kind0 == GOMP_MAP_DELETE) + { + if (kind1 == GOMP_MAP_ATTACH_DETACH) + fprintf (stderr, "delete, attach/detach\n"); + else + fprintf (stderr, "delete, unknown code %d\n", (int) kind1); + } + else if (kind0 == GOMP_MAP_TO_PSET) + { + if (kind1 == GOMP_MAP_ATTACH) + fprintf (stderr, "pset, attach\n"); + else if (kind1 == GOMP_MAP_DETACH) + fprintf (stderr, "pset, detach\n"); + else + fprintf (stderr, "pset, unknown code %d\n", (int) kind1); + } + else + fprintf (stderr, "unknown code %d, unknown code %d\n", + (int) kind0, (int) kind1); + break; + + case 3: + if (kind0 == GOMP_MAP_TO + || kind0 == GOMP_MAP_FROM + || kind0 == GOMP_MAP_TOFROM) + { + if (kind1 == GOMP_MAP_POINTER + && kind2 == GOMP_MAP_POINTER) + fprintf (stderr, "to/from, pointer, pointer\n"); + else if (kind1 == GOMP_MAP_ALWAYS_POINTER + && kind2 == GOMP_MAP_ALWAYS_POINTER) + fprintf (stderr, "to/from, always-pointer, always-pointer\n"); + else if (kind1 == GOMP_MAP_TO_PSET + && kind2 == GOMP_MAP_POINTER) + fprintf (stderr, "to/from, pset, pointer\n"); + else if (kind1 == GOMP_MAP_TO_PSET + && kind2 == GOMP_MAP_ALWAYS_POINTER) + fprintf (stderr, "to/from, pset, always-pointer\n"); + else if (kind1 == GOMP_MAP_TO_PSET + && kind2 == GOMP_MAP_ATTACH_DETACH) + fprintf (stderr, "to/from, pset, attach/detach\n"); + else if (kind1 == GOMP_MAP_POINTER + && kind2 == GOMP_MAP_ATTACH_DETACH) + fprintf (stderr, "to/from, pointer, attach/detach\n"); + else if (kind1 == GOMP_MAP_ALWAYS_POINTER + && kind2 == GOMP_MAP_ATTACH_DETACH) + fprintf (stderr, "to/from, always-pointer, attach/detach\n"); + else + fprintf (stderr, "to/from, unknown code %d, unknown code %d\n", + (int) kind1, (int) kind2); + } + else if (kind0 == GOMP_MAP_FORCE_TO + || kind0 == GOMP_MAP_FORCE_FROM + || kind0 == GOMP_MAP_FORCE_TOFROM) + { + if (kind1 == GOMP_MAP_TO_PSET + && kind2 == GOMP_MAP_POINTER) + fprintf (stderr, "force to/from, pset, pointer\n"); + else + fprintf (stderr, "force to/from, unknown code %d, " + "unknown code %d\n", (int) kind1, (int) kind2); + } + else if (kind0 == GOMP_MAP_ALLOC) + { + if (kind1 == GOMP_MAP_POINTER_TO_ZERO_LENGTH_ARRAY_SECTION + && kind2 == GOMP_MAP_ATTACH) + fprintf (stderr, "alloc, pointer to z-l-a-s, attach\n"); + else if (kind1 == GOMP_MAP_TO_PSET + && kind2 == GOMP_MAP_POINTER) + fprintf (stderr, "alloc, pset, pointer\n"); + else if (kind1 == GOMP_MAP_TO_PSET + && kind2 == GOMP_MAP_ALWAYS_POINTER) + fprintf (stderr, "alloc, pset, always-pointer\n"); + else + fprintf (stderr, "alloc, unknown code %d, unknown code %d\n", + (int) kind1, (int) kind2); + } + else if (kind0 == GOMP_MAP_DELETE) + { + if (kind1 == GOMP_MAP_TO_PSET + || kind1 == GOMP_MAP_POINTER) + fprintf (stderr, "delete, pset, pointer\n"); + else + fprintf (stderr, "delete, unknown code %d, unknown code %d\n", + (int) kind1, (int) kind2); + } + else + fprintf (stderr, "unknown code %d, unknown code %d, " + "unknown code %d\n", (int) kind0, (int) kind1, + (int) kind2); + break; + + case 4: + if (kind0 == GOMP_MAP_TO + || kind0 == GOMP_MAP_FROM + || kind0 == GOMP_MAP_TOFROM) + { + if (kind1 == GOMP_MAP_TO_PSET + && kind2 == GOMP_MAP_POINTER + && kind3 == GOMP_MAP_POINTER) + fprintf (stderr, "to/from, pset, pointer, pointer\n"); + else if (kind1 == GOMP_MAP_TO_PSET + && kind2 == GOMP_MAP_ALWAYS_POINTER + && kind3 == GOMP_MAP_POINTER) + fprintf (stderr, "to/from, pset, always pointer, pointer\n"); + else + fprintf (stderr, "to/from, unknown code %d, unknown code %d, " + "unknown code %d\n", (int) kind1, (int) kind2, + (int) kind3); + } + else + fprintf (stderr, "unknown code %d, unknown code %d, " + "unknown code %d, unknown code %d\n", (int) kind0, + (int) kind1, (int) kind2, (int) kind3); + break; + + default: + gcc_unreachable (); + } + + tree tmp = OMP_CLAUSE_CHAIN (grp_end); + OMP_CLAUSE_CHAIN (grp_end) = NULL_TREE; + debug_generic_expr (grp_start); + OMP_CLAUSE_CHAIN (grp_end) = tmp; + + num++; + } +} + +#endif + /* A pointer mapping group GRP may define a block of memory starting at some base address, and maybe also define a firstprivate pointer or firstprivate reference that points to that block. The return value is a node containing @@ -10493,6 +10814,9 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p, groups = omp_gather_mapping_groups (list_p); if (groups) { +#ifdef NOISY_TAXONOMY + omp_mapping_group_taxonomy (groups); +#endif hash_map *grpmap; grpmap = omp_index_mapping_groups (groups); @@ -10556,6 +10880,9 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p, groups = omp_gather_mapping_groups (list_p); if (groups) { +#ifdef NOISY_TAXONOMY + omp_mapping_group_taxonomy (groups); +#endif hash_map *grpmap; grpmap = omp_index_mapping_groups (groups); From patchwork Fri Oct 1 17:10:09 2021 Content-Type: text/plain; charset="utf-8" MIME-Version: 1.0 Content-Transfer-Encoding: 7bit X-Patchwork-Submitter: Julian Brown X-Patchwork-Id: 45707 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 44C9E3857407 for ; Fri, 1 Oct 2021 17:14:23 +0000 (GMT) X-Original-To: gcc-patches@gcc.gnu.org Delivered-To: gcc-patches@gcc.gnu.org Received: from esa2.mentor.iphmx.com (esa2.mentor.iphmx.com [68.232.141.98]) by sourceware.org (Postfix) with ESMTPS id 7A959385742C for ; Fri, 1 Oct 2021 17:10:21 +0000 (GMT) DMARC-Filter: OpenDMARC Filter v1.4.1 sourceware.org 7A959385742C 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: FmLmNgDZ5oqAoLLtju7YJukxrUlME1ivKMMjhVXeRmPCzXrAeFXt0IID2abU7uKDiMlbCn0u2R QNqKbHKatWZT7gTq+fBwz3ATpj7Q0Gi+oDsS6Ekz3iiZC+NKnujIUd6HtAOwWrQtqddhLmzCP0 DM2ZK3Xj/eMDIGQofoPEl7kMrLDLfAlSHBgoIsffyrvkVCXRd2oi4cI6WaF4S1pZrM7W0wgN2I ATG8y7W6nQyTfhJERx0/rfSO9Mql0MUg3PTnQRIWrsXv5gH7I9SXF/MzRujHpZO4ZzSK0zcQPW sPfAdj9xnEaJbXlomvUJc6NV X-IronPort-AV: E=Sophos;i="5.85,339,1624348800"; d="scan'208";a="66622524" Received: from orw-gwy-01-in.mentorg.com ([192.94.38.165]) by esa2.mentor.iphmx.com with ESMTP; 01 Oct 2021 09:10:20 -0800 IronPort-SDR: oLDiX+FSc9EM9T1qBPncyMZVKnbt3No/oQJlxiUubbuvDbAw+ZCBsSKK+W59CSr9cSfNWUdaK6 qDQXA8OUp4Gk3RdgibhFyETGLKfY4TWYMZ5q5YfAK9OnLBBvZxFV0SoyjCNgTnucF0ZfuQSKS1 ZCt2d7gkekFG09tDP4UUoXLCKU4i2IBHBTWjQ8Hw58QJbdh2bHQyCyO+UXi/e3/Ez1EAU0NTEY dLtt5kgZjliQgly91I6OLKgXvntNNnoOaWYjyc21Zk0qU2xgOQwq2t8pSBaIPGro8GpwKh4+cb OYk= From: Julian Brown To: Subject: [PATCH 11/11] OpenMP/OpenACC: [WIP] Add gcc_unreachable to apparently-dead path in build_struct_comp_nodes Date: Fri, 1 Oct 2021 10:10:09 -0700 Message-ID: <8bca087d3bfa223fc8413e7b0cc7e5294dfe4b64.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-02.mgc.mentorg.com (139.181.222.2) 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" The previous "not for committing" taxonomy patch shows that the path handling "extra nodes" in build_struct_comp_nodes is probably now dead, at least across the current testsuite. This patch adds gcc_unreachable on that path: this passes testing, which suggests that the extra node handling can probably be removed completely. (Otherwise we need test coverage for that path, ideally!) This is mostly posted as an FYI: a real patch would probably just remove the unused code path, if it really isn't needed any more. Thanks, Julian 2021-09-29 Julian Brown gcc/ * gimplify.c (build_struct_comp_nodes): Add gcc_unreachable on code path that appears to now be unused. --- gcc/gimplify.c | 3 +++ 1 file changed, 3 insertions(+) diff --git a/gcc/gimplify.c b/gcc/gimplify.c index d9fda21413d..3d444d1836f 100644 --- a/gcc/gimplify.c +++ b/gcc/gimplify.c @@ -8625,6 +8625,9 @@ build_struct_comp_nodes (enum tree_code code, tree grp_start, tree grp_end, OMP_CLAUSE_SIZE (c3) = TYPE_SIZE_UNIT (ptr_type_node); OMP_CLAUSE_CHAIN (c3) = NULL_TREE; + /* Apparently? */ + gcc_unreachable (); + *extra_node = c3; } else