From patchwork Mon Oct 11 13:41: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: 46102 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 E028D385780A for ; Mon, 11 Oct 2021 14:09:31 +0000 (GMT) X-Original-To: gcc-patches@gcc.gnu.org Delivered-To: gcc-patches@gcc.gnu.org Received: from esa1.mentor.iphmx.com (esa1.mentor.iphmx.com [68.232.129.153]) by sourceware.org (Postfix) with ESMTPS id ED4673858D3C for ; Mon, 11 Oct 2021 13:41:32 +0000 (GMT) DMARC-Filter: OpenDMARC Filter v1.4.1 sourceware.org ED4673858D3C 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: nCeazqlYJAGSRSlZqMn5aG/rW8+zUq+WpM2CAuzRyH9QR+6Sj/FIMZomUnTdo2t0vZarFk+l+J qVLwTRWvhSvIT5XZpUlBqtpt4cKtOmyoVZIzDb3rVX5g0iVFW0p7G4ftJYOvIpaEmo+uO4rd8r OWrgUg1cA+j4TvZlOVtqZmIeUuUKCBZnnRGW1XJptwQApanfFXAHp3Nqw/zmH1nsZNi2MCqiiv wbrDKJuyNB8+ypi1dRCCH3xIezHXEdf5CrPeBpUyd/8RK0CxXacLMJLTztLy73baIjxO6bUqD0 ru6NtLO6bUHHwkng5KuuqbR9 X-IronPort-AV: E=Sophos;i="5.85,364,1624348800"; d="scan'208";a="69506254" Received: from orw-gwy-02-in.mentorg.com ([192.94.38.167]) by esa1.mentor.iphmx.com with ESMTP; 11 Oct 2021 05:41:32 -0800 IronPort-SDR: kjUABVLh5UlerVD2TjqtL4ZwwjNK4ceUQyQmy7TlRK2rhCGKQj6WQAWHuBi8M8X+4HFShxxh7w XzGOYCNZ7K4Z8isXpUwidXM4sQ2EshCI96Bx9mkD9n0wZ+WMfwY/qT+hJmtP7BtKPeYRavjBaT x7ZCI0us2F04cl3vm7lMqxQd/2/38JhLL8Xcs0D9eBwsPID6+dmgxZlTfQL8lUyCut1HSizMZi jvjt8+XKngJHi+6KWL3YSuNCeCgBOwY3bj4aW+3tnFkz76S7hXAgQryvIYlXzLuBddDAf44NDz CCo= From: Julian Brown To: Subject: [PATCH 1/2] OpenMP: Handle reference-typed struct members Date: Mon, 11 Oct 2021 06:41:07 -0700 Message-ID: <20211011134108.129856-1-julian@codesourcery.com> X-Mailer: git-send-email 2.29.2 MIME-Version: 1.0 X-Originating-IP: [137.202.0.90] X-ClientProxiedBy: SVR-IES-MBX-04.mgc.mentorg.com (139.181.222.4) To SVR-IES-MBX-04.mgc.mentorg.com (139.181.222.4) X-Spam-Status: No, score=-11.8 required=5.0 tests=BAYES_00, GIT_PATCH_0, HEADER_FROM_DIFFERENT_DOMAINS, KAM_DMARC_STATUS, KAM_SHORT, SPF_HELO_PASS, SPF_PASS, TXREP autolearn=ham autolearn_force=no version=3.4.4 X-Spam-Checker-Version: SpamAssassin 3.4.4 (2020-01-24) on server2.sourceware.org X-BeenThere: gcc-patches@gcc.gnu.org X-Mailman-Version: 2.1.29 Precedence: list List-Id: Gcc-patches mailing list List-Unsubscribe: , List-Archive: List-Post: List-Help: List-Subscribe: , Cc: Jakub Jelinek , Thomas Schwinge Errors-To: gcc-patches-bounces+patchwork=sourceware.org@gcc.gnu.org Sender: "Gcc-patches" This patch fixes the baseptrs-3.C test case introduced in the patch: https://gcc.gnu.org/pipermail/gcc-patches/2021-October/580729.html The problematic case concerns OpenMP mapping clauses containing struct members of reference type, e.g. "mystruct.myref.myptr[:N]". To be able to access the array slice through the reference in the middle, we need to perform an attach action for that reference, since it is represented internally as a pointer. I don't think the spec allows for this case explicitly. The closest clause is (OpenMP 5.0, "2.19.7.1 map Clause"): "If the type of a list item is a reference to a type T then the reference in the device data environment is initialized to refer to the object in the device data environment that corresponds to the object referenced by the list item. If mapping occurs, it occurs as though the object were mapped through a pointer with an array section of type T and length one." The patch as is allows the mapping to work with just "mystruct.myref.myptr[:N]", without an explicit "mystruct.myref" mapping also (because, would that refer to the hidden pointer used by the reference, or the automatically-dereferenced data itself?). An attach/detach operation is thus synthesised for the reference. Tested with offloading to NVPTX and bootstrapped. OK (pending previously-posted series?). Julian Brown 2021-10-11 Julian Brown gcc/cp/ * semantics.c (finish_omp_clauses): Handle reference-typed members. gcc/ * gimplify.c (build_struct_group): Arrange for attach/detach nodes to be created for reference-typed struct members for OpenMP. Only create firstprivate_pointer/firstprivate_reference nodes for innermost struct accesses, those with an optionally-indirected DECL_P base. (omp_build_struct_sibling_lists): Handle two-element chain for inner struct component returned from build_struct_group. libgomp/ * testsuite/libgomp.c++/baseptrs-3.C: Remove XFAILs and extend test. --- gcc/cp/semantics.c | 4 + gcc/gimplify.c | 56 +++++++++-- libgomp/testsuite/libgomp.c++/baseptrs-3.C | 109 +++++++++++++++++++-- 3 files changed, 154 insertions(+), 15 deletions(-) diff --git a/gcc/cp/semantics.c b/gcc/cp/semantics.c index a50ec0ad883..bb8577d0d36 100644 --- a/gcc/cp/semantics.c +++ b/gcc/cp/semantics.c @@ -7862,6 +7862,8 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort) STRIP_NOPS (t); if (TREE_CODE (t) == POINTER_PLUS_EXPR) t = TREE_OPERAND (t, 0); + if (REFERENCE_REF_P (t)) + t = TREE_OPERAND (t, 0); } } while (TREE_CODE (t) == COMPONENT_REF); @@ -7961,6 +7963,8 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort) { t = TREE_OPERAND (TREE_OPERAND (t, 0), 0); indir_component_ref_p = true; + if (REFERENCE_REF_P (t)) + t = TREE_OPERAND (t, 0); STRIP_NOPS (t); if (TREE_CODE (t) == POINTER_PLUS_EXPR) t = TREE_OPERAND (t, 0); diff --git a/gcc/gimplify.c b/gcc/gimplify.c index 3d444d1836f..d187dfe1ef2 100644 --- a/gcc/gimplify.c +++ b/gcc/gimplify.c @@ -10249,7 +10249,10 @@ build_struct_group (enum omp_region_type region_type, enum tree_code code, /* 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 (openmp + && attach_detach + && !(TREE_CODE (TREE_TYPE (ocd)) == REFERENCE_TYPE + && TREE_CODE (TREE_TYPE (TREE_TYPE (ocd))) != POINTER_TYPE)) return NULL; #ifdef NOISY_SIBLING_LISTS @@ -10317,9 +10320,32 @@ build_struct_group (enum omp_region_type region_type, enum tree_code code, tree noind = strip_indirections (base); - if (!openmp + if (openmp + && TREE_CODE (TREE_TYPE (noind)) == REFERENCE_TYPE && (region_type & ORT_TARGET) && TREE_CODE (noind) == COMPONENT_REF) + { + tree c2 = build_omp_clause (OMP_CLAUSE_LOCATION (grp_end), + OMP_CLAUSE_MAP); + OMP_CLAUSE_SET_MAP_KIND (c2, GOMP_MAP_TO); + OMP_CLAUSE_DECL (c2) = unshare_expr (base); + OMP_CLAUSE_SIZE (c2) = TYPE_SIZE_UNIT (TREE_TYPE (noind)); + + tree c3 = build_omp_clause (OMP_CLAUSE_LOCATION (grp_end), + OMP_CLAUSE_MAP); + OMP_CLAUSE_SET_MAP_KIND (c3, GOMP_MAP_ATTACH_DETACH); + OMP_CLAUSE_DECL (c3) = unshare_expr (noind); + OMP_CLAUSE_SIZE (c3) = size_zero_node; + + OMP_CLAUSE_CHAIN (c2) = c3; + OMP_CLAUSE_CHAIN (c3) = NULL_TREE; + + *inner = c2; + return NULL; + } + else if (!openmp + && (region_type & ORT_TARGET) + && TREE_CODE (noind) == COMPONENT_REF) { /* The base for this component access is a struct component access itself. Insert a node to be processed on the next iteration of @@ -10333,13 +10359,30 @@ build_struct_group (enum omp_region_type region_type, enum tree_code code, 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)); + OMP_CLAUSE_CHAIN (c2) = NULL_TREE; *inner = c2; return NULL; } - tree sdecl = strip_components_and_deref (base); + tree sdecl = base; + /* There are too many places we need to do things like this. */ + if (TREE_CODE (sdecl) == INDIRECT_REF + || TREE_CODE (sdecl) == MEM_REF) + { + sdecl = TREE_OPERAND (sdecl, 0); + if (TREE_CODE (sdecl) == INDIRECT_REF + && (TREE_CODE (TREE_TYPE (TREE_OPERAND (sdecl, 0))) + == REFERENCE_TYPE)) + sdecl = TREE_OPERAND (sdecl, 0); + } - if (POINTER_TYPE_P (TREE_TYPE (sdecl)) && (region_type & ORT_TARGET)) + while (TREE_CODE (sdecl) == COMPONENT_REF + || TREE_CODE (sdecl) == POINTER_PLUS_EXPR) + sdecl = TREE_OPERAND (sdecl, 0); + + if (DECL_P (sdecl) + && POINTER_TYPE_P (TREE_TYPE (sdecl)) + && (region_type & ORT_TARGET)) { tree c2 = build_omp_clause (OMP_CLAUSE_LOCATION (grp_end), OMP_CLAUSE_MAP); @@ -10740,11 +10783,10 @@ omp_build_struct_sibling_lists (enum tree_code code, 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.grp_end = (OMP_CLAUSE_CHAIN (inner) + ? OMP_CLAUSE_CHAIN (inner) : inner); newgrp.mark = UNVISITED; newgrp.sibling = NULL; newgrp.next = NULL; diff --git a/libgomp/testsuite/libgomp.c++/baseptrs-3.C b/libgomp/testsuite/libgomp.c++/baseptrs-3.C index cabeb7c2b7a..39a48a40920 100644 --- a/libgomp/testsuite/libgomp.c++/baseptrs-3.C +++ b/libgomp/testsuite/libgomp.c++/baseptrs-3.C @@ -1,9 +1,58 @@ -/* { dg-xfail-if "fails to parse correctly" { *-*-* } } */ - #include #include #include +struct sa0 +{ + int *ptr; +}; + +struct sb0 +{ + int arr[10]; +}; + +struct sc0 +{ + sa0 a; + sb0 b; + sc0 (sa0 &my_a, sb0 &my_b) : a(my_a), b(my_b) {} +}; + +void +foo0 () +{ + sa0 my_a; + sb0 my_b; + + my_a.ptr = (int *) malloc (sizeof (int) * 10); + sc0 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); +} + struct sa { int *ptr; @@ -90,6 +139,49 @@ bar () free (my_a.ptr); } +struct scp0 +{ + sa *a; + sb *b; + scp0 (sa *my_a, sb *my_b) : a(my_a), b(my_b) {} +}; + +void +foop0 () +{ + sa *my_a = new sa; + sb *my_b = new sb; + + my_a->ptr = new int[10]; + scp0 *my_c = new scp0(my_a, my_b); + + memset (my_c->a->ptr, 0, sizeof (int) * 10); + + #pragma omp target map (my_c->a, my_c->a[:1], 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, my_c->b[:1], 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; +} + struct scp { sa *&a; @@ -108,7 +200,7 @@ foop () memset (my_c->a->ptr, 0, sizeof (int) * 10); - #pragma omp target map (my_c->a->ptr, my_c->a->ptr[:10]) + #pragma omp target map (my_c->a, my_c->a[:1], my_c->a->ptr, my_c->a->ptr[:10]) { for (int i = 0; i < 10; i++) my_c->a->ptr[i] = i; @@ -119,8 +211,7 @@ foop () memset (my_c->b->arr, 0, sizeof (int) * 10); -/* FIXME: This currently ICEs. */ -/* #pragma omp target map (my_c->b->arr[:10]) */ + #pragma omp target map (my_c->b, my_c->b[:1], my_c->b->arr[:10]) { for (int i = 0; i < 10; i++) my_c->b->arr[i] = i; @@ -146,7 +237,8 @@ barp () memset (my_cref->a->ptr, 0, sizeof (int) * 10); - #pragma omp target map (my_cref->a->ptr, my_cref->a->ptr[:10]) + #pragma omp target map (my_cref->a, my_cref->a[:1], my_cref->a->ptr, \ + my_cref->a->ptr[:10]) { for (int i = 0; i < 10; i++) my_cref->a->ptr[i] = i; @@ -157,8 +249,7 @@ barp () memset (my_cref->b->arr, 0, sizeof (int) * 10); -/* FIXME: This currently ICEs. */ -/* #pragma omp target map (my_cref->b->arr[:10]) */ + #pragma omp target map (my_cref->b, my_cref->b[:1], my_cref->b->arr[:10]) { for (int i = 0; i < 10; i++) my_cref->b->arr[i] = i; @@ -174,8 +265,10 @@ barp () int main (int argc, char *argv[]) { + foo0 (); foo (); bar (); + foop0 (); foop (); barp (); return 0; From patchwork Mon Oct 11 13:41: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: 46103 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 1CDA53857819 for ; Mon, 11 Oct 2021 14:10:01 +0000 (GMT) X-Original-To: gcc-patches@gcc.gnu.org Delivered-To: gcc-patches@gcc.gnu.org Received: from esa1.mentor.iphmx.com (esa1.mentor.iphmx.com [68.232.129.153]) by sourceware.org (Postfix) with ESMTPS id 5DD243858D3C for ; Mon, 11 Oct 2021 13:41:35 +0000 (GMT) DMARC-Filter: OpenDMARC Filter v1.4.1 sourceware.org 5DD243858D3C 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: jwuvpUwIK7luDQ2St/hKqIcNsjWR+5LKfZiygy34Z2k3biHFhe5yuUOQkkEttd4J0CB63kQkRp Xiib+ZKHCRjLV8e2vi4RjytRtuzD956Stgi1ygAEANRQ6oBPo+spqlZeajv9Ss54Bguk4tu3GW FbbBYEQYYAXuqtRfJURU91RSUJzi/iZ5YVmaG/l597w/EfNH/E7bSxJVMXPSwk3Cm5hGK6HSLb MME/qI/qYpNvBtENMtXrvIU9ezMGQlMv9+aLLJD3LfgOPPrkW+UTFxjJR6uxQ0j8MDmowRU+v9 N8nX+aLR/4G/jFrCRi5RRdxs X-IronPort-AV: E=Sophos;i="5.85,364,1624348800"; d="scan'208";a="69506255" Received: from orw-gwy-02-in.mentorg.com ([192.94.38.167]) by esa1.mentor.iphmx.com with ESMTP; 11 Oct 2021 05:41:32 -0800 IronPort-SDR: OI2tli6MzLPDgb1+pLRNZsQN2MOxbdbSXdEnA/7ySLl+0i5vdKbpb8QOz3GIqfODKXX8CQpmXv 1z4ZL07EKK7LW/E9WaitShavY6z5AsLeIZg1+P93RP42A9evV2AhmsqShIYDVVFyzBM1apdmkg A4jHfvhttHIkchAd+vQssn7b3cCd2xsfO7m7hjW5LVV5YS5ZFeqIiUQCYzJKqf1kAgmZ5IwOGs 80H6Byl8G1KEJmCjZlw51RW5sDgHEsJxkzuu5TxkTsqMz+tDDB8DwfN1tmEJiG1DKp4kxWlHyR Efs= From: Julian Brown To: Subject: [PATCH 2/2] OpenACC: Make deep-copy-arrayofstruct.c a libgomp/runtime test Date: Mon, 11 Oct 2021 06:41:08 -0700 Message-ID: <20211011134108.129856-2-julian@codesourcery.com> X-Mailer: git-send-email 2.29.2 In-Reply-To: <20211011134108.129856-1-julian@codesourcery.com> References: <20211011134108.129856-1-julian@codesourcery.com> MIME-Version: 1.0 X-Originating-IP: [137.202.0.90] X-ClientProxiedBy: SVR-IES-MBX-04.mgc.mentorg.com (139.181.222.4) To SVR-IES-MBX-04.mgc.mentorg.com (139.181.222.4) X-Spam-Status: No, score=-11.8 required=5.0 tests=BAYES_00, GIT_PATCH_0, HEADER_FROM_DIFFERENT_DOMAINS, KAM_DMARC_STATUS, 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" I noticed that the test in question now compiles properly, and in fact runs properly too. Thus it's more useful as a runtime test than a passing compilation test that otherwise doesn't do much. This patch moves it to libgomp. Tested with offloading to NVPTX and bootstrapped. (I can probably self-approve as a testsuite-only change, but the patch depends on previously-posted series). Julian 2021-10-11 Julian Brown gcc/testsuite/ * libgomp.oacc-c-c++-common/deep-copy-arrayofstruct.c: Move test from here. libgomp/ * testsuite/libgomp.oacc-c-c++-common/deep-copy-arrayofstruct.c: Move test to here. --- .../libgomp.oacc-c-c++-common}/deep-copy-arrayofstruct.c | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) rename {gcc/testsuite/c-c++-common/goacc => libgomp/testsuite/libgomp.oacc-c-c++-common}/deep-copy-arrayofstruct.c (98%) diff --git a/gcc/testsuite/c-c++-common/goacc/deep-copy-arrayofstruct.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/deep-copy-arrayofstruct.c similarity index 98% rename from gcc/testsuite/c-c++-common/goacc/deep-copy-arrayofstruct.c rename to libgomp/testsuite/libgomp.oacc-c-c++-common/deep-copy-arrayofstruct.c index 4247607b61c..a11c64749cc 100644 --- a/gcc/testsuite/c-c++-common/goacc/deep-copy-arrayofstruct.c +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/deep-copy-arrayofstruct.c @@ -1,4 +1,4 @@ -/* { dg-do compile } */ +/* { dg-do run } */ #include #include