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; +}