From patchwork Thu Nov 25 14: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: 48148 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 12A2A385C40C for ; Thu, 25 Nov 2021 14:15:28 +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 810073857C56 for ; Thu, 25 Nov 2021 14:10:23 +0000 (GMT) DMARC-Filter: OpenDMARC Filter v1.4.1 sourceware.org 810073857C56 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: n0YaoVga6K6zcfwDv4yacIDCq7JlM+ymzvceA4c69RytQJBSQ2yiVzUD5YkakF9D8lGE98DG55 mVgFlnrMCuCdTxarL9iX4QdusA6Lhpft0JD2tpeoGaVtGGflfzkMZ6/PlTR2t8D55kQ6jZxF17 Yf1DpZfTRfvb/wcQKd5Sx+7OGPw8zx2xTaPG7JgDoCFzanteHr/WJ16nUC33PjW0BipZDMyYeI YPvzNc27T6z92/sdn44mxppKfMLg2Ie+4bWH6faoX85NmtfAaJgHJMAJRnuSZV6I704TJYyEO4 EQnh/UMOC0Rm+XVCmWgUsOtK X-IronPort-AV: E=Sophos;i="5.87,263,1631606400"; d="scan'208";a="68920866" Received: from orw-gwy-02-in.mentorg.com ([192.94.38.167]) by esa2.mentor.iphmx.com with ESMTP; 25 Nov 2021 06:10:22 -0800 IronPort-SDR: i9AzXoZ45NUBQrNmgGOpDhdvotj+ZPWgPmgwNzPphqJQ3YHS0grcNf5+vO3bO+/fJ79K02ksEF o8jBOayfVBm8UBG6VDDaWc6gXFlY9s0X8ipSo0+LV0/K6fBjaMhYNcd491givGP4XUvrk8ypL+ GldoMdbiwJwWTLNMxqHi5D9iS96oNF7r3cLBYHSKXDu1OqMFKMsrIos3xzICS/3MGMblzsf7WR HmAhukUVxNkUR8WF5MZQep5oSjC/sjiMOH2M8tfwQ6PW0MdkuViSmF4N+ERqFSTQMJpCDCl97C DII= From: Julian Brown To: Subject: [PATCH 10/16] OpenMP: Fix non-zero attach/detach bias for struct dereferences Date: Thu, 25 Nov 2021 06:10:09 -0800 Message-ID: <20211125141013.113782-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-15.mgc.mentorg.com (139.181.222.15) 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? 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 d5dad99ff97..dd103d8eecd 100644 --- a/gcc/c-family/c-common.h +++ b/gcc/c-family/c-common.h @@ -1251,6 +1251,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 3f84fd1b5cb..a90696fe706 100644 --- a/gcc/c-family/c-omp.c +++ b/gcc/c-family/c-omp.c @@ -3113,6 +3113,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 4d156f6d3ec..cfac7d0a2b5 100644 --- a/gcc/c/c-typeck.c +++ b/gcc/c/c-typeck.c @@ -13799,9 +13799,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)) @@ -13834,10 +13840,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 e882c302f31..068c0c69e58 100644 --- a/gcc/cp/semantics.c +++ b/gcc/cp/semantics.c @@ -5620,9 +5620,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)) { @@ -5684,12 +5691,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; +}