From patchwork Mon Feb 8 16:17:56 2021 Content-Type: text/plain; charset="utf-8" MIME-Version: 1.0 Content-Transfer-Encoding: 7bit X-Patchwork-Submitter: Chung-Lin Tang X-Patchwork-Id: 1437749 Return-Path: X-Original-To: incoming@patchwork.ozlabs.org Delivered-To: patchwork-incoming@bilbo.ozlabs.org Authentication-Results: ozlabs.org; spf=pass (sender SPF authorized) smtp.mailfrom=gcc.gnu.org (client-ip=2620:52:3:1:0:246e:9693:128c; helo=sourceware.org; envelope-from=gcc-patches-bounces@gcc.gnu.org; receiver=) Received: from sourceware.org (server2.sourceware.org [IPv6:2620:52:3:1:0:246e:9693:128c]) (using TLSv1.3 with cipher TLS_AES_256_GCM_SHA384 (256/256 bits) key-exchange X25519 server-signature RSA-PSS (4096 bits) server-digest SHA256) (No client certificate requested) by ozlabs.org (Postfix) with ESMTPS id 4DZB5g3QVtz9sB4 for ; Tue, 9 Feb 2021 03:18:14 +1100 (AEDT) Received: from server2.sourceware.org (localhost [IPv6:::1]) by sourceware.org (Postfix) with ESMTP id 25F883898536; Mon, 8 Feb 2021 16:18:11 +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 6D0CA3898536 for ; Mon, 8 Feb 2021 16:18:06 +0000 (GMT) DMARC-Filter: OpenDMARC Filter v1.3.2 sourceware.org 6D0CA3898536 Authentication-Results: sourceware.org; dmarc=none (p=none dis=none) header.from=codesourcery.com Authentication-Results: sourceware.org; spf=pass smtp.mailfrom=ChungLin_Tang@mentor.com IronPort-SDR: mP3XbMohjhVmlTG14ME3DhLG9XFIhMJ4T4JcnKKjKYRWhDMH6/I6lluUHHN1rIMsXjuZXz4y4S czDCL+5owFkyhcjnyQmEpt/EM2eRYgtTAT++vR4m9P5Qb5BikHJF6z/G43chS+Guf8ehSjQ/es tir5AIcS4oSfvYqap3iRA9kQES4rzX+PM/1szdQgnFESgd4Pr1i9eUxyryigFIC+MToD5iltlV ppIgNMf/0iK0hBOw12GRMZHzfN7iKWlgInTmsy+XkxjA4wrJQaBvm5JkQBAbkIPoqf/W8rIcDo rqg= X-IronPort-AV: E=Sophos;i="5.81,162,1610438400"; d="scan'208,223";a="57868890" Received: from orw-gwy-02-in.mentorg.com ([192.94.38.167]) by esa2.mentor.iphmx.com with ESMTP; 08 Feb 2021 08:18:04 -0800 IronPort-SDR: OcmanH+RAmRQjPcmBe+s5q5FJCgdmAlmRtjECo7bU3a/1FT5c6/Jr8iWmefvtzObPsde+57+ya r62wbiScyqDiuJ9o7sdDHy9tS6Jr7nYxtT4v5yPcmEhVRukHHPmYd15VL5sLRI333+E2ldSCFo dH6PeYOVe5EwuCzUx7+t3WWzXsA8J8Cchg09mkMir/geHJQF9/XLm+ChhWVxW/iGX9s/3Y1jKR B+kPGPo5zBVDIKAjzIBSR7LYVzRcO+R6yM9q+pmdGG8QTuP3W0QdGIXkmtd34uZ9CZGmgNN7xt 2LM= From: Chung-Lin Tang Subject: [PATCH, OG10, committed] Support A->B expressions in map clause To: gcc-patches , Catherine Moore , Julian Brown Message-ID: <12e8b66f-c321-3a50-e6ea-fdc701e1f278@codesourcery.com> Date: Tue, 9 Feb 2021 00:17:56 +0800 User-Agent: Mozilla/5.0 (Macintosh; Intel Mac OS X 10.13; rv:78.0) Gecko/20100101 Thunderbird/78.7.0 MIME-Version: 1.0 Content-Language: en-US X-ClientProxiedBy: svr-orw-mbx-04.mgc.mentorg.com (147.34.90.204) To svr-orw-mbx-04.mgc.mentorg.com (147.34.90.204) X-Spam-Status: No, score=-10.1 required=5.0 tests=BAYES_00, GIT_PATCH_0, HEADER_FROM_DIFFERENT_DOMAINS, KAM_DMARC_STATUS, RCVD_IN_DNSWL_NONE, SPF_HELO_PASS, SPF_PASS, TXREP autolearn=ham autolearn_force=no version=3.4.2 X-Spam-Checker-Version: SpamAssassin 3.4.2 (2018-09-13) 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@gcc.gnu.org Sender: "Gcc-patches" This patch tries to allow map(A->ptr) to be properly handled the same way as map(B.ptr) expressions. map(struct:*A) clauses are now produced during gimplify. Julian, I'm CCing you since IIRC you seemed to be the author of this area of code. Would appreciate if you gave a look if you have time, though I've already went ahead and pushed to OG10 after testing results looked okay. Thanks, Chung-Lin gcc/ChangeLog: * gimplify.c ("tree-hash-traits.h"): Add include. (gimplify_scan_omp_clauses): Change struct_map_to_clause to type hash_map *. Adjust struct map handling to handle cases of *A and A->B expressions. (gimplify_adjust_omp_clauses): Move GOMP_MAP_STRUCT removal code for exit data directives code to earlier position. gcc/testsuite/ChangeLog: * g++.dg/gomp/target-3.C: Adjust testcase gimple scanning. * g++.dg/gomp/target-this-2.C: Likewise. * g++.dg/gomp/target-this-3.C: Likewise. * g++.dg/gomp/target-this-4.C: Likewise. libgomp/ChangeLog: * testsuite/libgomp.c++/target-23.C: New testcase. From bf8605f14ec33ea31233a3567f3184fee667b695 Mon Sep 17 00:00:00 2001 From: Chung-Lin Tang Date: Mon, 8 Feb 2021 07:53:55 -0800 Subject: [PATCH] Enable gimplify GOMP_MAP_STRUCT handling of (COMPONENT_REF (INDIRECT_REF ...)) map clauses. This patch tries to allow map(A->ptr) to be properly handled the same way as map(B.ptr) expressions. map(struct:*A) clauses are now produced during gimplify. This patch, as of time of commit, is only pushed to devel/omp/gcc-10, not yet submitted as mainline patch to upstream. 2021-02-08 Chung-Lin Tang gcc/ChangeLog: * gimplify.c ("tree-hash-traits.h"): Add include. (gimplify_scan_omp_clauses): Change struct_map_to_clause to type hash_map *. Adjust struct map handling to handle cases of *A and A->B expressions. (gimplify_adjust_omp_clauses): Move GOMP_MAP_STRUCT removal code for exit data directives code to earlier position. gcc/testsuite/ChangeLog: * g++.dg/gomp/target-3.C: Adjust testcase gimple scanning. * g++.dg/gomp/target-this-2.C: Likewise. * g++.dg/gomp/target-this-3.C: Likewise. * g++.dg/gomp/target-this-4.C: Likewise. libgomp/ChangeLog: * testsuite/libgomp.c++/target-23.C: New testcase. --- gcc/gimplify.c | 51 +++++++++++++++++++++++-------- gcc/testsuite/g++.dg/gomp/target-3.C | 2 +- gcc/testsuite/g++.dg/gomp/target-this-2.C | 2 +- gcc/testsuite/g++.dg/gomp/target-this-3.C | 2 +- gcc/testsuite/g++.dg/gomp/target-this-4.C | 4 +-- libgomp/testsuite/libgomp.c++/target-23.C | 34 +++++++++++++++++++++ 6 files changed, 78 insertions(+), 17 deletions(-) create mode 100644 libgomp/testsuite/libgomp.c++/target-23.C diff --git a/gcc/gimplify.c b/gcc/gimplify.c index b90ba5b..ba19017 100644 --- a/gcc/gimplify.c +++ b/gcc/gimplify.c @@ -53,6 +53,7 @@ along with GCC; see the file COPYING3. If not see #include "langhooks.h" #include "tree-cfg.h" #include "tree-ssa.h" +#include "tree-hash-traits.h" #include "omp-general.h" #include "omp-low.h" #include "gimple-low.h" @@ -8514,7 +8515,7 @@ 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_map_to_clause = NULL; hash_set *struct_deref_set = NULL; tree *prev_list_p = NULL, *orig_list_p = list_p; int handled_depend_iterators = -1; @@ -9082,12 +9083,15 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p, && TREE_CODE (decl) == INDIRECT_REF && TREE_CODE (TREE_OPERAND (decl, 0)) == COMPONENT_REF && (TREE_CODE (TREE_TYPE (TREE_OPERAND (decl, 0))) - == REFERENCE_TYPE)) + == REFERENCE_TYPE) + && (OMP_CLAUSE_MAP_KIND (c) + != GOMP_MAP_POINTER_TO_ZERO_LENGTH_ARRAY_SECTION)) { pd = &TREE_OPERAND (decl, 0); decl = TREE_OPERAND (decl, 0); } bool indir_p = false; + bool component_ref_p = false; tree orig_decl = decl; tree decl_ref = NULL_TREE; if ((region_type & (ORT_ACC | ORT_TARGET | ORT_TARGET_DATA)) != 0 @@ -9098,6 +9102,7 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p, 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)) @@ -9117,8 +9122,11 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p, } } } - else if (TREE_CODE (decl) == COMPONENT_REF) + 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 @@ -9188,7 +9196,8 @@ 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 (DECL_P (decl) + if ((DECL_P (decl) + || (component_ref_p && INDIRECT_REF_P (decl))) && 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 @@ -9245,7 +9254,10 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p, gcc_assert (base == decl); splay_tree_node n - = splay_tree_lookup (ctx->variables, (splay_tree_key)decl); + = (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) @@ -9271,7 +9283,11 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p, OMP_CLAUSE_SET_MAP_KIND (c, k); has_attachments = true; } - if (n == NULL || (n->value & GOVD_MAP) == 0) + 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))) { tree l = build_omp_clause (OMP_CLAUSE_LOCATION (c), OMP_CLAUSE_MAP); @@ -9290,7 +9306,8 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p, ? DECL_SIZE_UNIT (OMP_CLAUSE_DECL (l)) : 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 + = new hash_map; struct_map_to_clause->put (decl, l); if (ptr || attach_detach) { @@ -9324,7 +9341,8 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p, flags |= GOVD_SEEN; if (has_attachments) flags |= GOVD_MAP_HAS_ATTACHMENTS; - goto do_add_decl; + if (DECL_P (decl)) + goto do_add_decl; } else if (struct_map_to_clause) { @@ -9433,6 +9451,13 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p, } else if (*sc != c) { + if (gimplify_expr (pd, pre_p, NULL, is_gimple_lvalue, + fb_lvalue) + == GS_ERROR) + { + remove = true; + break; + } *list_p = OMP_CLAUSE_CHAIN (c); OMP_CLAUSE_CHAIN (c) = *sc; *sc = c; @@ -10811,6 +10836,12 @@ gimplify_adjust_omp_clauses (gimple_seq *pre_p, gimple_seq body, tree *list_p, } } } + if (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_STRUCT + && (code == OMP_TARGET_EXIT_DATA || code == OACC_EXIT_DATA)) + { + remove = true; + break; + } if (!DECL_P (decl)) { if ((ctx->region_type & ORT_TARGET) != 0 @@ -10857,10 +10888,6 @@ gimplify_adjust_omp_clauses (gimple_seq *pre_p, gimple_seq body, tree *list_p, = OMP_CLAUSE_CHAIN (OMP_CLAUSE_CHAIN (c)); } } - else if (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_STRUCT - && (code == OMP_TARGET_EXIT_DATA - || code == OACC_EXIT_DATA)) - remove = true; else if (DECL_SIZE (decl) && TREE_CODE (DECL_SIZE (decl)) != INTEGER_CST && OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_POINTER diff --git a/gcc/testsuite/g++.dg/gomp/target-3.C b/gcc/testsuite/g++.dg/gomp/target-3.C index fe2e38b..f4d40ec 100644 --- a/gcc/testsuite/g++.dg/gomp/target-3.C +++ b/gcc/testsuite/g++.dg/gomp/target-3.C @@ -33,4 +33,4 @@ T::bar (int x) template struct T<0>; -/* { dg-final { scan-tree-dump-times "map\\(alloc:this->b \\\[len: \[0-9\]+\\\]\\) map\\(alloc:this->a \\\[len: \[0-9\]+\\\]\\)" 4 "gimple" } } */ +/* { 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" } } */ diff --git a/gcc/testsuite/g++.dg/gomp/target-this-2.C b/gcc/testsuite/g++.dg/gomp/target-this-2.C index a5e8321..679c85a 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\(tofrom:\*_[0-9]+ \[len: [0-9]+\]\) map\(always_pointer:__closure->__this \[pointer assign, bias: 0\]\)} "gimple" } } */ +/* { 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" } } */ diff --git a/gcc/testsuite/g++.dg/gomp/target-this-3.C b/gcc/testsuite/g++.dg/gomp/target-this-3.C index 208ea07..a450b37 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.* map\(tofrom:\*this \[len: [0-9]+\]\) map\(firstprivate:this \[pointer assign, bias: 0\]\) map\(alloc:\*this->refptr \[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\]\) firstprivate\(n\)} "gimple" } } */ +/* { dg-final { scan-tree-dump {#pragma omp target num_teams.* 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\]\) firstprivate\(n\)} "gimple" } } */ /* { dg-final { scan-tree-dump {#pragma omp target num_teams.* 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:MEM.* \[len: 0\]\) firstprivate\(n\)} "gimple" } } */ diff --git a/gcc/testsuite/g++.dg/gomp/target-this-4.C b/gcc/testsuite/g++.dg/gomp/target-this-4.C index f42cf38..af23cbb 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.* 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\(attach_zero_length_array_section:_[0-9]+->ptr \[bias: 0\]\) map\(from:mapped \[len: 1\]\) map\(alloc:MEM.* \[len: 0\]\) firstprivate\(n\) map\(alloc:MEM.* \[len: 0\]\) map\(firstprivate:this \[pointer assign, bias: 0\]\)} "gimple" } } */ +/* { dg-final { scan-tree-dump {#pragma omp target num_teams.* 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:MEM.* \[len: 0\]\) firstprivate\(n\) map\(alloc:MEM.* \[len: 0\]\) map\(firstprivate:this \[pointer assign, bias: 0\]\)} "gimple" } } */ -/* { dg-final { scan-tree-dump {#pragma omp target num_teams.* 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\(alloc:\*_[0-9]+->refptr \[pointer assign, zero-length array section, bias: 0\]\) map\(attach:_3->refptr \[bias: 0\]\) map\(from:mapped \[len: [0-9]+\]\) map\(alloc:\*_[0-9]+ \[len: 0\]\) firstprivate\(n\) map\(alloc:MEM.* \[len: 0\]\) map\(firstprivate:this \[pointer assign, bias: 0\]\)} "gimple" } } */ +/* { dg-final { scan-tree-dump {#pragma omp target num_teams.* 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\]\) firstprivate\(n\) map\(alloc:MEM.* \[len: 0\]\) map\(firstprivate:this \[pointer assign, bias: 0\]\)} "gimple" } } */ diff --git a/libgomp/testsuite/libgomp.c++/target-23.C b/libgomp/testsuite/libgomp.c++/target-23.C new file mode 100644 index 0000000..d4f9ff3 --- /dev/null +++ b/libgomp/testsuite/libgomp.c++/target-23.C @@ -0,0 +1,34 @@ +extern "C" void abort (); + +struct S +{ + int *data; +}; + +int +main (void) +{ + #define SZ 10 + S *s = new S (); + s->data = new int[SZ]; + + for (int i = 0; i < SZ; i++) + s->data[i] = 0; + + #pragma omp target enter data map(to: s) + #pragma omp target enter data map(to: s->data[:SZ]) + #pragma omp target + { + for (int i = 0; i < SZ; i++) + s->data[i] = i; + } + #pragma omp target exit data map(from: s->data[:SZ]) + #pragma omp target exit data map(from: s) + + for (int i = 0; i < SZ; i++) + if (s->data[i] != i) + abort (); + + return 0; +} +