From patchwork Mon Feb 21 15:18:57 2022 Content-Type: text/plain; charset="utf-8" MIME-Version: 1.0 Content-Transfer-Encoding: 7bit X-Patchwork-Submitter: Chung-Lin Tang X-Patchwork-Id: 51261 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 B90A7385DC09 for ; Mon, 21 Feb 2022 15:19: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 B98D53858D37 for ; Mon, 21 Feb 2022 15:19:11 +0000 (GMT) DMARC-Filter: OpenDMARC Filter v1.4.1 sourceware.org B98D53858D37 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: VE4f0W05xmtvL11Rj7688HkwD2DUcTRZDhf64j/vtoeYGiOJeX+O9cwDJEkm/z/BPqqCrsZS9i Wrzkhl9sez5VXSLkdAmc5mMaGjmq4PG8VzXr63L1LIbUBDVKBS605TbIfFPiOyWmo1dxEsRqzC JbTix7gDMvXMbLbwhX8IUAfcwIHHL8RcfuNBBk3boAWTV7Mpj6b1ipxIEw7GWvsnZbJQxLZP7R r9PwwyHV0s1TanuPnqm3VQnuPgm0t7VJ6Q56AJTIaa3KIYh26dzxNWM7zeZuwZczUVG6CbEdeO RCMAk3COs1ryXOMEmH9PW48w X-IronPort-AV: E=Sophos;i="5.88,386,1635235200"; d="scan'208";a="72094090" Received: from orw-gwy-02-in.mentorg.com ([192.94.38.167]) by esa3.mentor.iphmx.com with ESMTP; 21 Feb 2022 07:19:11 -0800 IronPort-SDR: Mac7oo6IRIJgM0uDAec60Z1ZbgaF4HFcdM03sKoZJXlBXYP0iLoHxxJ1kWkKanJ2TM5Vo7RvtX vbOLqFrVKh+sA+7dXX+31T45wEEmEX3+pAiBp5gffYD2bBvuQu+7ImBX0fEBkE9INmcBb0md6S kgJvJ9zkzMvGPpXXx7Vuv2IXkNFAxRb9kxN84KloU4CKcoB8BKImhb7BHvr7/nl/V5te/Lac4b ugz9stlLGMbMBtRDhVdHsqe2woJjwQ/HJ27Hm2wbAM9omhg8aLnwi1W/rFZEYyjiMebneM/X5A pAw= Message-ID: <65e1462e-d17f-2975-1401-358fe9c69e28@codesourcery.com> Date: Mon, 21 Feb 2022 23:18:57 +0800 MIME-Version: 1.0 User-Agent: Mozilla/5.0 (Macintosh; Intel Mac OS X 10.13; rv:91.0) Gecko/20100101 Thunderbird/91.6.0 Content-Language: en-US To: gcc-patches , Jakub Jelinek From: Chung-Lin Tang Subject: [PATCH, OpenMP, C/C++] Handle array reference base-pointers in array sections X-ClientProxiedBy: svr-orw-mbx-03.mgc.mentorg.com (147.34.90.203) To svr-orw-mbx-02.mgc.mentorg.com (147.34.90.202) X-Spam-Status: No, score=-10.5 required=5.0 tests=BAYES_00, GIT_PATCH_0, HEADER_FROM_DIFFERENT_DOMAINS, KAM_DMARC_STATUS, SPF_HELO_PASS, SPF_PASS, TXREP, T_SCC_BODY_TEXT_LINE 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: , Errors-To: gcc-patches-bounces+patchwork=sourceware.org@gcc.gnu.org Sender: "Gcc-patches" Hi Jakub, as encountered in cases where a program constructs its own deep-copying for arrays-of-pointers, e.g: #pragma omp target enter data map(to:level->vectors[:N]) for (i = 0; i < N; i++) #pragma omp target enter data map(to:level->vectors[i][:N]) We need to treat the part of the array reference before the array section as a base-pointer (here 'level->vectors[i]'), providing pointer-attachment behavior. This patch adds this inside handle_omp_array_sections(), tracing the whole sequence of array dimensions, creating a whole base-pointer reference iteratively using build_array_ref(). The conditions are that each of the "absorbed" dimensions must be length==1, and the final reference must be of pointer-type (so that pointer attachment makes sense). There's also a little patch in gimplify_scan_omp_clauses(), to make sure the array-ref base-pointer goes down the right path. This case was encountered when working to make 534.hpgmgfv_t from SPEChpc 2021 properly compile. Tested without regressions on trunk. Okay to go in once stage1 opens? Thanks, Chung-Lin 2022-02-21 Chung-Lin Tang gcc/c/ChangeLog: * c-typeck.cc (handle_omp_array_sections): Add handling for creating array-reference base-pointer attachment clause. gcc/cp/ChangeLog: * semantics.cc (handle_omp_array_sections): Add handling for creating array-reference base-pointer attachment clause. gcc/ChangeLog: * gimplify.cc (gimplify_scan_omp_clauses): Add case for attach/detach map kind for ARRAY_REF of POINTER_TYPE. gcc/testsuite/ChangeLog: * c-c++-common/gomp/target-enter-data-1.c: Adjust testcase. libgomp/testsuite/ChangeLog: * libgomp.c-c++-common/ptr-attach-2.c: New test. diff --git a/gcc/c/c-typeck.cc b/gcc/c/c-typeck.cc index 3075c883548..4257e373557 100644 --- a/gcc/c/c-typeck.cc +++ b/gcc/c/c-typeck.cc @@ -13649,6 +13649,10 @@ handle_omp_array_sections (tree c, enum c_omp_region_type ort) if (int_size_in_bytes (TREE_TYPE (first)) <= 0) maybe_zero_len = true; + struct dim { tree low_bound, length; }; + auto_vec dims (num); + dims.safe_grow (num); + for (i = num, t = OMP_CLAUSE_DECL (c); i > 0; t = TREE_CHAIN (t)) { @@ -13763,6 +13767,9 @@ handle_omp_array_sections (tree c, enum c_omp_region_type ort) else size = size_binop (MULT_EXPR, size, l); } + + dim d = { low_bound, length }; + dims[i] = d; } if (side_effects) size = build2 (COMPOUND_EXPR, sizetype, side_effects, size); @@ -13802,6 +13809,23 @@ handle_omp_array_sections (tree c, enum c_omp_region_type ort) OMP_CLAUSE_DECL (c) = t; return false; } + + tree aref = t; + for (i = 0; i < dims.length (); i++) + { + if (dims[i].length && integer_onep (dims[i].length)) + { + tree lb = dims[i].low_bound; + aref = build_array_ref (OMP_CLAUSE_LOCATION (c), aref, lb); + } + else + { + if (TREE_CODE (TREE_TYPE (aref)) == POINTER_TYPE) + t = aref; + break; + } + } + first = c_fully_fold (first, false, NULL); OMP_CLAUSE_DECL (c) = first; if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_HAS_DEVICE_ADDR) @@ -13836,7 +13860,8 @@ handle_omp_array_sections (tree c, enum c_omp_region_type ort) break; } tree c2 = build_omp_clause (OMP_CLAUSE_LOCATION (c), OMP_CLAUSE_MAP); - if (TREE_CODE (t) == COMPONENT_REF) + if (TREE_CODE (t) == COMPONENT_REF || TREE_CODE (t) == ARRAY_REF + || TREE_CODE (t) == INDIRECT_REF) OMP_CLAUSE_SET_MAP_KIND (c2, GOMP_MAP_ATTACH_DETACH); else OMP_CLAUSE_SET_MAP_KIND (c2, GOMP_MAP_FIRSTPRIVATE_POINTER); diff --git a/gcc/cp/semantics.cc b/gcc/cp/semantics.cc index 0cb17a6a8ab..646f4883d66 100644 --- a/gcc/cp/semantics.cc +++ b/gcc/cp/semantics.cc @@ -5497,6 +5497,10 @@ handle_omp_array_sections (tree c, enum c_omp_region_type ort) if (processing_template_decl && maybe_zero_len) return false; + struct dim { tree low_bound, length; }; + auto_vec dims (num); + dims.safe_grow (num); + for (i = num, t = OMP_CLAUSE_DECL (c); i > 0; t = TREE_CHAIN (t)) { @@ -5604,6 +5608,9 @@ handle_omp_array_sections (tree c, enum c_omp_region_type ort) else size = size_binop (MULT_EXPR, size, l); } + + dim d = { low_bound, length }; + dims[i] = d; } if (!processing_template_decl) { @@ -5647,6 +5654,24 @@ handle_omp_array_sections (tree c, enum c_omp_region_type ort) OMP_CLAUSE_DECL (c) = t; return false; } + + tree aref = t; + for (i = 0; i < dims.length (); i++) + { + if (dims[i].length && integer_onep (dims[i].length)) + { + tree lb = dims[i].low_bound; + aref = convert_from_reference (aref); + aref = build_array_ref (OMP_CLAUSE_LOCATION (c), aref, lb); + } + else + { + if (TREE_CODE (TREE_TYPE (aref)) == POINTER_TYPE) + t = aref; + break; + } + } + OMP_CLAUSE_DECL (c) = first; if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_HAS_DEVICE_ADDR) return false; @@ -5681,7 +5706,8 @@ handle_omp_array_sections (tree c, enum c_omp_region_type ort) bool reference_always_pointer = true; tree c2 = build_omp_clause (OMP_CLAUSE_LOCATION (c), OMP_CLAUSE_MAP); - if (TREE_CODE (t) == COMPONENT_REF) + if (TREE_CODE (t) == COMPONENT_REF || TREE_CODE (t) == ARRAY_REF + || (TREE_CODE (t) == INDIRECT_REF && !REFERENCE_REF_P (t))) { OMP_CLAUSE_SET_MAP_KIND (c2, GOMP_MAP_ATTACH_DETACH); diff --git a/gcc/gimplify.cc b/gcc/gimplify.cc index f570daa015a..77b95cd8000 100644 --- a/gcc/gimplify.cc +++ b/gcc/gimplify.cc @@ -9626,7 +9626,10 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p, || (component_ref_p && (INDIRECT_REF_P (decl) || TREE_CODE (decl) == MEM_REF - || TREE_CODE (decl) == ARRAY_REF))) + || TREE_CODE (decl) == ARRAY_REF)) + || (TREE_CODE (decl) == ARRAY_REF + && TREE_CODE (TREE_TYPE (decl)) == POINTER_TYPE + && OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ATTACH_DETACH)) && 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 diff --git a/gcc/testsuite/c-c++-common/gomp/target-enter-data-1.c b/gcc/testsuite/c-c++-common/gomp/target-enter-data-1.c index ce766d29e2d..3a1b488fa1f 100644 --- a/gcc/testsuite/c-c++-common/gomp/target-enter-data-1.c +++ b/gcc/testsuite/c-c++-common/gomp/target-enter-data-1.c @@ -21,4 +21,5 @@ void func (struct foo *f, int n, int m) #pragma omp target enter data map (to: f->bars[n].vectors[:f->bars[n].num_vectors]) } -/* { dg-final { scan-tree-dump-times "map\\(to:\\*_\[0-9\]+ \\\[len: _\[0-9\]+\\\]\\) map\\(attach:\[^-\]+->vectors \\\[bias: \[^\]\]+\\\]\\)" 3 "gimple" } } */ +/* { dg-final { scan-tree-dump-times "map\\(to:\\*_\[0-9\]+ \\\[len: _\[0-9\]+\\\]\\) map\\(attach:\\*_\[0-9\]+ \\\[bias: \[^\]\]+\\\]\\)" 1 "gimple" } } */ +/* { dg-final { scan-tree-dump-times "map\\(to:\\*_\[0-9\]+ \\\[len: _\[0-9\]+\\\]\\) map\\(attach:\[^-\]+->vectors \\\[bias: \[^\]\]+\\\]\\)" 2 "gimple" } } */ diff --git a/libgomp/testsuite/libgomp.c-c++-common/ptr-attach-2.c b/libgomp/testsuite/libgomp.c-c++-common/ptr-attach-2.c new file mode 100644 index 00000000000..889a4a253ae --- /dev/null +++ b/libgomp/testsuite/libgomp.c-c++-common/ptr-attach-2.c @@ -0,0 +1,60 @@ +#include + +struct blk { int x, y; }; +struct L +{ + #define N 10 + struct { + int num_blocks[N]; + struct blk * blocks[N]; + } m; +}; + +void foo (struct L *l) +{ + for (int i = 0; i < N; i++) + { + l->m.blocks[i] = (struct blk *) malloc (sizeof (struct blk) * N); + l->m.num_blocks[i] = N; + } + + #pragma omp target enter data map(to:l[:1]) + for (int i = 0; i < N; i++) + { + #pragma omp target enter data map(to:l->m.blocks[i][:l->m.num_blocks[i]]) + } + + #pragma omp target + { + for (int i = 0; i < N; i++) + for (int j = 0; j < N; j++) + { + l->m.blocks[i][j].x = i + j; + l->m.blocks[i][j].y = i * j; + } + } + + for (int i = 0; i < N; i++) + { + #pragma omp target exit data map(from:l->m.blocks[i][:l->m.num_blocks[i]]) + } + #pragma omp target exit data map(from:l[:1]) + + + for (int i = 0; i < N; i++) + for (int j = 0; j < N; j++) + { + if (l->m.blocks[i][j].x != i + j) + abort (); + if (l->m.blocks[i][j].y != i * j) + abort (); + } + +} + +int main (void) +{ + struct L l; + foo (&l); + return 0; +}