From patchwork Fri Dec 8 14:28:59 2023 Content-Type: text/plain; charset="utf-8" MIME-Version: 1.0 Content-Transfer-Encoding: 8bit X-Patchwork-Submitter: Tobias Burnus X-Patchwork-Id: 175820 Return-Path: Delivered-To: ouuuleilei@gmail.com Received: by 2002:a59:bcd1:0:b0:403:3b70:6f57 with SMTP id r17csp5492230vqy; Fri, 8 Dec 2023 06:29:34 -0800 (PST) X-Google-Smtp-Source: AGHT+IEDHqAo8KZ4QwCnbouexW510zEHmZvEnJPU/rNouBCbNnwXIoc6miSJO5lZetGmboAWGf40 X-Received: by 2002:a67:ee84:0:b0:464:77e6:ea70 with SMTP id n4-20020a67ee84000000b0046477e6ea70mr190687vsp.33.1702045773342; Fri, 08 Dec 2023 06:29:33 -0800 (PST) ARC-Seal: i=2; a=rsa-sha256; t=1702045773; cv=pass; d=google.com; s=arc-20160816; b=WSD8ZpObF6wtDBLlxuctr/zlSqSO3EJm8a6fZzC61vPg4b6ChHDZbxrgK0vmFnxLKc dpqxFSG8NQurF061oQ6TOuOjuBV5AG3XdOb1z/cFpeoMhlYPkF4zs3Guf19X8rm/MBop QYu9XhL/WfAkpVoGWkH8NadWJxh9aPCYlcFG7cPeW9nar6uQrAKXNYYdRthWc/xMqhuC cnFJ4xUKl2EmOM/mlBMekRwVp85RSRdkkqZtIlX9c0k/bOdTAcnMIrdTySI4L0rrZRpk C6Xu/hFmtPR3whRvQuGuDtbINk3TkP/r/MvMnPwkKTEsmBizgOf8gjatL8idlYy1it2k kZBQ== ARC-Message-Signature: i=2; a=rsa-sha256; c=relaxed/relaxed; d=google.com; s=arc-20160816; h=errors-to:list-subscribe:list-help:list-post:list-archive :list-unsubscribe:list-id:precedence:subject:to:from :content-language:user-agent:mime-version:date:message-id :ironport-sdr:arc-filter:dmarc-filter:delivered-to; bh=vQy7U9bjWbBjWxECYoPOMvbddtFawa3SJgyLB4cBFOc=; fh=CT9ANdoiYgjVt7HQGxGdzRB+HPa196wDRuos5o/1sIg=; b=f2NHZ3pLHxsCZvdrRnCMqVus/ebO5btiDf1wqlpTlNOY8+C9eo1Io5m6KZrp1k+gVP k3X1m6kQPNOmqt1H5mZBpftjnrxrxZY7tB49TfBqXAicHywJ2E5ZlchziKLkk+4q8bRo t1muGle+Fwwg9snLshuoZbXE0x6qJb1XS5+k9Jf8NNhI+KDpEu85L+fRQdc10Efs3W4B m2mW25dlYZxNF4fMK5ppvmEZdfYXN1BcOWvKXC8dVFrO/8B15uHJAv1fuvrkRw0LCr1g qUYyfEVW6l2n8Og3pPQeHmrkAwa2oJ8mxGRfRq6dOFfFYuRsZisi7eFFFg2H6/mju4zf LHkQ== ARC-Authentication-Results: i=2; mx.google.com; arc=pass (i=1); spf=pass (google.com: domain of gcc-patches-bounces+ouuuleilei=gmail.com@gcc.gnu.org designates 2620:52:3:1:0:246e:9693:128c as permitted sender) smtp.mailfrom="gcc-patches-bounces+ouuuleilei=gmail.com@gcc.gnu.org" Received: from server2.sourceware.org (server2.sourceware.org. [2620:52:3:1:0:246e:9693:128c]) by mx.google.com with ESMTPS id n19-20020a056102049300b00464767ae605si448469vsa.690.2023.12.08.06.29.33 for (version=TLS1_3 cipher=TLS_AES_256_GCM_SHA384 bits=256/256); Fri, 08 Dec 2023 06:29:33 -0800 (PST) Received-SPF: pass (google.com: domain of gcc-patches-bounces+ouuuleilei=gmail.com@gcc.gnu.org designates 2620:52:3:1:0:246e:9693:128c as permitted sender) client-ip=2620:52:3:1:0:246e:9693:128c; Authentication-Results: mx.google.com; arc=pass (i=1); spf=pass (google.com: domain of gcc-patches-bounces+ouuuleilei=gmail.com@gcc.gnu.org designates 2620:52:3:1:0:246e:9693:128c as permitted sender) smtp.mailfrom="gcc-patches-bounces+ouuuleilei=gmail.com@gcc.gnu.org" Received: from server2.sourceware.org (localhost [IPv6:::1]) by sourceware.org (Postfix) with ESMTP id 19AFC385734D for ; Fri, 8 Dec 2023 14:29:33 +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 143F03858C60 for ; Fri, 8 Dec 2023 14:29:06 +0000 (GMT) DMARC-Filter: OpenDMARC Filter v1.4.2 sourceware.org 143F03858C60 Authentication-Results: sourceware.org; dmarc=none (p=none dis=none) header.from=codesourcery.com Authentication-Results: sourceware.org; spf=pass smtp.mailfrom=mentor.com ARC-Filter: OpenARC Filter v1.0.0 sourceware.org 143F03858C60 Authentication-Results: server2.sourceware.org; arc=none smtp.remote-ip=68.232.137.180 ARC-Seal: i=1; a=rsa-sha256; d=sourceware.org; s=key; t=1702045749; cv=none; b=w9LQxdn0AyaN5R6HTlci32HSOquq6D3SD98Ebph5FHZgHOiv7KSqpU+FY69rmPfOcwgrzHlRRSHmCqbIErMsJtKLPbxhG38XePRyUn3UIJu23ovcfeL+cuwpyRBnynEvnChLfPAXJsCx3emJ2Sd4LLfs8KDMIC0gbStXwOVM/60= ARC-Message-Signature: i=1; a=rsa-sha256; d=sourceware.org; s=key; t=1702045749; c=relaxed/simple; bh=Sa7LzTcBXfhPvkgciK2g9XiIuryectKw6l9c5AR8UpY=; h=Message-ID:Date:MIME-Version:From:To:Subject; b=koVFSv9uQ2I+5EQos4rH8enboIH87XFtfel4s6sUiWnDkcp/D+/3HBtgGBv+gDfUDv27PPvbhsImMMfWQ66gLgr5Bt0JxHziVMNR0P0IhSa13a+X8Ast0YIw7EP7vGj2VyDQGzTRnaP+OK2N+LyNjfr7aAz5XqR53oeyu8t7JXg= ARC-Authentication-Results: i=1; server2.sourceware.org X-CSE-ConnectionGUID: XFtk25HlSoi5LSaAyM+fSQ== X-CSE-MsgGUID: LRBFSXMSRpenJ5LtHv87ow== X-IronPort-AV: E=Sophos;i="6.04,261,1695715200"; d="diff'?scan'208";a="24814581" Received: from orw-gwy-01-in.mentorg.com ([192.94.38.165]) by esa3.mentor.iphmx.com with ESMTP; 08 Dec 2023 06:29:04 -0800 IronPort-SDR: HLymXiPIxzBqLg6z25bDiY85b6lVwmBF432MnThQNJp3aWK5XIq7+E5oRLQrbVSNdcgIp03ZbT Gtq4qC66QMFMWnXzjWy8PBBGCvFNL4x9h0kf/jFwvBML+LW74DP23Ly9HmJ3h5fiv79yZa9wRS O3PLsvTLlfwSH5KBMyuItKMLfKrMTtR4j8cIjUbaAOCETSKBE/mhRK3A96n7Dm1GOIGfX9giEo agJR5BIRcn8Cyxx6j8OgYEE8WK1tN/FZaiPiTfNHII66vX5C9EIF5yjv3LAPooAtoEybmp7GI1 IQY= Message-ID: Date: Fri, 8 Dec 2023 15:28:59 +0100 MIME-Version: 1.0 User-Agent: Mozilla Thunderbird Content-Language: en-US From: Tobias Burnus To: gcc-patches , Jakub Jelinek Subject: [patch] OpenMP: Handle same-directive mapped vars with pointer predefined firstprivate [PR110639] X-Originating-IP: [137.202.0.90] X-ClientProxiedBy: svr-ies-mbx-12.mgc.mentorg.com (139.181.222.12) To svr-ies-mbx-12.mgc.mentorg.com (139.181.222.12) X-Spam-Status: No, score=-11.3 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.6 X-Spam-Checker-Version: SpamAssassin 3.4.6 (2021-04-09) on server2.sourceware.org X-BeenThere: gcc-patches@gcc.gnu.org X-Mailman-Version: 2.1.30 Precedence: list List-Id: Gcc-patches mailing list List-Unsubscribe: , List-Archive: List-Post: List-Help: List-Subscribe: , Errors-To: gcc-patches-bounces+ouuuleilei=gmail.com@gcc.gnu.org X-getmail-retrieved-from-mailbox: INBOX X-GMAIL-THRID: 1784724349038375688 X-GMAIL-MSGID: 1784724349038375688 This patch fixes the issue: int a[100]; p = &a[0]; #pragma omp target map(a) p[0] = p[99] = 3; where 'p' is predetermined firstprivate, i.e. it is firstprivatized but its address gets updated to the device address of 'a' as there is associated storage for the value of 'p', i.e. its pointee. [This is a C/C++-only feature that cannot be replicated by using a single clause. ('target data map(a) use_device_ptr(p)' + 'target is_device_ptr(p)' would do so in two steps. - or 'p2 = omp_get_mapped_ptr(p, devnum)' + 'target is_device_ptr(p2)'.)] Before this only worked when that storage was mapped before and not on the same directive. The gimplify_scan_omp_clauses change was done when I saw some runtime fails; I think those were due to a bug in libgomp (now fixed) and not due to having two pointer privatisations in a now different order. Still, they at least prevent mapping 'this' multiple times when 'this' is not 'this' but __closure->this which is at least a missed optimization. And also for libgomp.c++/pr108286.C which has a normal 'this' and map(tofrom:*this [len: 16]). Build and tested without offloading and with nvptx offloading. Comments, remarks, suggestions? * * * (I wonder whether our current approach of removing explicit MAP if its DECL is unsued is the right one if there is any GOVD_MAP_0LEN_ARRAY around - or even any OMP_CLAUSE_MAP_MAYBE_ZERO_LENGTH_ARRAY_SECTION.) (See new libgomp.c-c++-common/target-implicit-map-6.c; BTW, I tried: before '(void) a;' but that only worked with C and not with C++.) * * * The other issue in the PR (still to be done) is for code like: int a[100]; p = &a[0]; #pragma omp target map(a[20:20]) // Map only a[20] to a[40], but p points to &a[0] p[20] = p[30] = 3; where 'p' points to the base address of 'a' but p[0] == a[0] it not actually mapped. As we currently do not keep track of base pointer, this won't work. I have not (yet) explored how to best implement this. * * * OpenMP Spec: The first feature is not new, but I have not checked the wording in 4.5 or 5.0; it might be that older versions only required it to work for storage mapped before the current taget directive. But at least TR12 is very explicit in permitting it and the (nonpublic) issue which lead to the 5.1 change also uses this. (See PR.) (The second feature is definitely new in OpenMP 5.1.) TR12 states in "14.8 target Construct" [379:8-10]: "[C/C++] If a list item in a map clause has a base pointer that is predetermined firstprivate (see Section 6.1.1) and on entry to the target region the list item is mapped, the firstprivate pointer is updated via corresponding base pointer initialization." (For OpenMP 5.1, read its Section 2.21.7.2.) Tobias ----------------- Siemens Electronic Design Automation GmbH; Anschrift: Arnulfstraße 201, 80634 München; Gesellschaft mit beschränkter Haftung; Geschäftsführer: Thomas Heurung, Frank Thürauf; Sitz der Gesellschaft: München; Registergericht München, HRB 106955 OpenMP: Handle same-directive mapped vars with pointer predefined firstprivate [PR110639] Predefined 'firstprivate' for pointer variables firstprivatizes the pointer but if it is associated with a mapped target, its address is updated to the corresponding target. (If not, the host value remains.) This commit extends this handling to also update the pointer address for storaged mapped on the same directive. The 'gimplify_scan_omp_clauses' change avoids adding an additional map(alloc:this) (+ptr assignment) when there is already a map(tofrom:*this) (+ptr assignment) This shows up for libgomp.c++/pr108286.C and also when 'this' is actually '__closure->this' (-> g++.dg/gomp/target-{this-{2,4},lambda-1}.C). PR middle-end/110639 gcc/ChangeLog: * gimplify.cc (struct gimplify_adjust_omp_clauses_data): Add append_list. (gimplify_adjust_omp_clauses_1, gimplify_adjust_omp_clauses): Add GOVD_MAP_0LEN_ARRAY clauses at the end. (gimplify_scan_omp_clauses): Mark also '*var' as found not only 'var'. libgomp/ChangeLog: * target.c (gomp_map_vars_internal): Handle also variables mapped in the same directive for GOVD_MAP_0LEN_ARRAY. * testsuite/libgomp.c++/pr108286.C: Add gimple tree-scan test. * testsuite/libgomp.c-c++-common/target-implicit-map-6.c: New test. gcc/testsuite/ChangeLog: * g++.dg/gomp/target-this-2.C: Remove 'this' pointer mapping alreay mapped via __closure->this. * g++.dg/gomp/target-this-4.C: Likewise. * g++.dg/gomp/target-lambda-1.C: Likewise. Move 'iptr' pointer mapping to the end in scan-tree-dump. gcc/gimplify.cc | 45 ++++- gcc/testsuite/g++.dg/gomp/target-lambda-1.C | 4 +- gcc/testsuite/g++.dg/gomp/target-this-2.C | 4 +- gcc/testsuite/g++.dg/gomp/target-this-4.C | 6 +- libgomp/target.c | 11 +- libgomp/testsuite/libgomp.c++/pr108286.C | 4 + .../libgomp.c-c++-common/target-implicit-map-6.c | 212 +++++++++++++++++++++ 7 files changed, 276 insertions(+), 10 deletions(-) diff --git a/gcc/gimplify.cc b/gcc/gimplify.cc index 342e43a7f25..2234fd6b7e1 100644 --- a/gcc/gimplify.cc +++ b/gcc/gimplify.cc @@ -11586,6 +11586,23 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p, else if (!DECL_P (decl)) { tree d = decl, *pd; + pd = &OMP_CLAUSE_DECL (c); + if (TREE_CODE (decl) == INDIRECT_REF) + { + tree d2 = TREE_OPERAND (decl, 0); + STRIP_NOPS (d2); + if (DECL_P (d2)) + { + if (gimplify_expr (pd, pre_p, NULL, is_gimple_lvalue, + fb_lvalue) == GS_ERROR) + { + remove = true; + break; + } + decl = d2; + goto handle_map_decl; + } + } if (TREE_CODE (d) == ARRAY_REF) { while (TREE_CODE (d) == ARRAY_REF) @@ -11594,7 +11611,6 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p, && TREE_CODE (TREE_TYPE (d)) == ARRAY_TYPE) decl = d; } - pd = &OMP_CLAUSE_DECL (c); if (d == decl && TREE_CODE (decl) == INDIRECT_REF && TREE_CODE (TREE_OPERAND (decl, 0)) == COMPONENT_REF @@ -11774,6 +11790,7 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p, } break; } + handle_map_decl: flags = GOVD_MAP | GOVD_EXPLICIT; if (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ALWAYS_TO || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ALWAYS_TOFROM) @@ -11806,7 +11823,7 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p, OMP_CLAUSE_SET_MAP_KIND (c, map_kind); } - goto do_add; + goto do_add_decl; case OMP_CLAUSE_AFFINITY: gimplify_omp_affinity (list_p, pre_p); @@ -12571,6 +12588,7 @@ omp_find_stores_stmt (gimple_stmt_iterator *gsi_p, struct gimplify_adjust_omp_clauses_data { tree *list_p; + tree append_list; gimple_seq *pre_p; }; @@ -12691,6 +12709,7 @@ gimplify_adjust_omp_clauses_1 (splay_tree_node n, void *data) && omp_shared_to_firstprivate_optimizable_decl_p (decl)) omp_mark_stores (gimplify_omp_ctxp->outer_context, decl); + bool len0_append_list_used = false; tree chain = *list_p; clause = build_omp_clause (input_location, code); OMP_CLAUSE_DECL (clause) = decl; @@ -12707,6 +12726,11 @@ gimplify_adjust_omp_clauses_1 (splay_tree_node n, void *data) OMP_CLAUSE_FIRSTPRIVATE_IMPLICIT (clause) = 1; else if (code == OMP_CLAUSE_MAP && (flags & GOVD_MAP_0LEN_ARRAY) != 0) { + /* For GOVD_MAP_0LEN_ARRAY, add the clauses to append_list such + that those come after any data mapping. */ + len0_append_list_used = true; + struct gimplify_adjust_omp_clauses_data *adjdata + = (struct gimplify_adjust_omp_clauses_data *) data; tree nc = build_omp_clause (input_location, OMP_CLAUSE_MAP); OMP_CLAUSE_DECL (nc) = decl; if (TREE_CODE (TREE_TYPE (decl)) == REFERENCE_TYPE @@ -12721,8 +12745,9 @@ gimplify_adjust_omp_clauses_1 (splay_tree_node n, void *data) OMP_CLAUSE_SET_MAP_KIND (clause, GOMP_MAP_ALLOC); OMP_CLAUSE_MAP_MAYBE_ZERO_LENGTH_ARRAY_SECTION (clause) = 1; OMP_CLAUSE_SET_MAP_KIND (nc, GOMP_MAP_FIRSTPRIVATE_POINTER); - OMP_CLAUSE_CHAIN (nc) = chain; + OMP_CLAUSE_CHAIN (nc) = adjdata->append_list; OMP_CLAUSE_CHAIN (clause) = nc; + adjdata->append_list = clause; struct gimplify_omp_ctx *ctx = gimplify_omp_ctxp; gimplify_omp_ctxp = ctx->outer_context; gimplify_expr (&TREE_OPERAND (OMP_CLAUSE_DECL (clause), 0), @@ -12833,7 +12858,8 @@ gimplify_adjust_omp_clauses_1 (splay_tree_node n, void *data) (ctx->region_type & ORT_ACC) != 0); gimplify_omp_ctxp = ctx; } - *list_p = clause; + if (!len0_append_list_used) + *list_p = clause; struct gimplify_omp_ctx *ctx = gimplify_omp_ctxp; gimplify_omp_ctxp = ctx->outer_context; /* Don't call omp_finish_clause on implicitly added OMP_CLAUSE_PRIVATE @@ -12842,7 +12868,7 @@ gimplify_adjust_omp_clauses_1 (splay_tree_node n, void *data) if (code != OMP_CLAUSE_PRIVATE || ctx->region_type != ORT_SIMD) lang_hooks.decls.omp_finish_clause (clause, pre_p, (ctx->region_type & ORT_ACC) != 0); - if (gimplify_omp_ctxp) + if (gimplify_omp_ctxp && !len0_append_list_used) for (; clause != chain; clause = OMP_CLAUSE_CHAIN (clause)) if (OMP_CLAUSE_CODE (clause) == OMP_CLAUSE_MAP && DECL_P (OMP_CLAUSE_SIZE (clause))) @@ -13445,6 +13471,7 @@ gimplify_adjust_omp_clauses (gimple_seq *pre_p, gimple_seq body, tree *list_p, /* Add in any implicit data sharing. */ struct gimplify_adjust_omp_clauses_data data; + data.append_list = NULL; if ((gimplify_omp_ctxp->region_type & ORT_ACC) == 0) { /* OpenMP. Implicit clauses are added at the start of the clause list, @@ -13472,6 +13499,14 @@ gimplify_adjust_omp_clauses (gimple_seq *pre_p, gimple_seq body, tree *list_p, "iterator"); break; } + if (data.append_list != NULL_TREE && *data.list_p != NULL_TREE) + { + for (c = *data.list_p; c && OMP_CLAUSE_CHAIN (c); c = OMP_CLAUSE_CHAIN (c)) + ; + OMP_CLAUSE_CHAIN (c) = data.append_list; + } + else if (data.append_list != NULL_TREE) + *data.list_p = data.append_list; gimplify_omp_ctxp = ctx->outer_context; delete_omp_context (ctx); diff --git a/gcc/testsuite/g++.dg/gomp/target-lambda-1.C b/gcc/testsuite/g++.dg/gomp/target-lambda-1.C index 5ce8ceadb19..b4f1593af02 100644 --- a/gcc/testsuite/g++.dg/gomp/target-lambda-1.C +++ b/gcc/testsuite/g++.dg/gomp/target-lambda-1.C @@ -87,7 +87,9 @@ int main (void) return 0; } -/* { dg-final { scan-tree-dump {#pragma omp target num_teams.* firstprivate\(b\) map\(alloc:MEM.* \[len: 0\]\) map\(firstprivate:iptr \[pointer assign, bias: 0\]\) map\(alloc:MEM.* \[len: 0\]\) map\(firstprivate:this \[pointer assign, bias: 0\]\) 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\(from:mapped \[len: [0-9]+\]\) map\(alloc:\*_[0-9]+ \[len: 0\]\) map\(alloc:\*_[0-9]+ \[len: 0\]\) map\(attach_zero_length_array_section:__closure->__iptr \[bias: 0\]\) map\(attach_zero_length_array_section:_[0-9]+->ptr \[bias: 0\]\)} "gimple" } } */ +/* Note that 'this' = '__closure->__this' such that no pointer-assign for 'this' should appear. */ + +/* { dg-final { scan-tree-dump {#pragma omp target num_teams.* firstprivate\(b\) 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\(from:mapped \[len: [0-9]+\]\) map\(alloc:\*_[0-9]+ \[len: 0\]\) map\(alloc:\*_[0-9]+ \[len: 0\]\) map\(attach_zero_length_array_section:__closure->__iptr \[bias: 0\]\) map\(attach_zero_length_array_section:_[0-9]+->ptr \[bias: 0\]\) map\(alloc:MEM.* \[len: 0\]\) map\(firstprivate:iptr \[pointer assign, bias: 0\]\)[\r\n]} "gimple" } } */ /* { dg-final { scan-tree-dump {#pragma omp target num_teams.* firstprivate\(end\) firstprivate\(begin\) map\(to:loop \[len: [0-9]+\]\) map\(alloc:\*_[0-9]+ \[len: 0\]\) map\(attach_zero_length_array_section:loop\.__data1 \[bias: 0\]\)} "gimple" } } */ diff --git a/gcc/testsuite/g++.dg/gomp/target-this-2.C b/gcc/testsuite/g++.dg/gomp/target-this-2.C index cc08e7e8693..eecab5a25e8 100644 --- a/gcc/testsuite/g++.dg/gomp/target-this-2.C +++ b/gcc/testsuite/g++.dg/gomp/target-this-2.C @@ -46,4 +46,6 @@ int main (void) return 0; } -/* { dg-final { scan-tree-dump {map\(alloc:MEM\[\(char \*\)_[0-9]+\] \[len: [0-9]+\]\) map\(firstprivate:this \[pointer assign, bias: 0\]\) firstprivate\(m\) 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\(from:v \[len: [0-9]+\]\)} "gimple" } } */ +/* Note that 'this' = '__closure->__this' such that no pointer-assign for 'this' should appear. */ + +/* { dg-final { scan-tree-dump {firstprivate\(n\) firstprivate\(m\) 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\(from:v \[len: [0-9]+\]\)[\r\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 9ade3cc0b2b..845f89f6997 100644 --- a/gcc/testsuite/g++.dg/gomp/target-this-4.C +++ b/gcc/testsuite/g++.dg/gomp/target-this-4.C @@ -102,6 +102,8 @@ int main (void) return 0; } -/* { dg-final { scan-tree-dump {#pragma omp target num_teams.* firstprivate\(n\) map\(alloc:MEM.* \[len: 0\]\) map\(firstprivate:this \[pointer assign, bias: 0\]\) 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\(from:mapped \[len: 1\]\) map\(alloc:\*_[0-9]+ \[len: 0\]\) map\(attach_zero_length_array_section:_[0-9]+->ptr \[bias: 0\]\)} "gimple" } } */ +/* Note that 'this' = '__closure->__this' such that no pointer-assign for 'this' should appear. */ -/* { dg-final { scan-tree-dump {#pragma omp target num_teams.* firstprivate\(n\) map\(alloc:MEM.* \[len: 0\]\) map\(firstprivate:this \[pointer assign, bias: 0\]\) 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\(from:mapped \[len: [0-9]+\]\) map\(alloc:\*_[0-9]+ \[len: 0\]\) map\(alloc:\*_[0-9]+ \[pointer assign, zero-length array section, bias: 0\]\) map\(attach:_[0-9]+->refptr \[bias: 0\]\)} "gimple" } } */ +/* { dg-final { scan-tree-dump {#pragma omp target num_teams.* firstprivate\(n\) 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\(from:mapped \[len: 1\]\) map\(alloc:\*_[0-9]+ \[len: 0\]\) map\(attach_zero_length_array_section:_[0-9]+->ptr \[bias: 0\]\)[\r\n]} "gimple" } } */ + +/* { dg-final { scan-tree-dump {#pragma omp target num_teams.* firstprivate\(n\) 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\(from:mapped \[len: [0-9]+\]\) map\(alloc:\*_[0-9]+ \[len: 0\]\) map\(alloc:\*_[0-9]+ \[pointer assign, zero-length array section, bias: 0\]\) map\(attach:_[0-9]+->refptr \[bias: 0\]\)[\r\n]} "gimple" } } */ diff --git a/libgomp/target.c b/libgomp/target.c index f30c20255d3..c1f26e6860b 100644 --- a/libgomp/target.c +++ b/libgomp/target.c @@ -1149,7 +1149,8 @@ gomp_map_vars_internal (struct gomp_device_descr *devicep, splay_tree_key n; if ((kind & typemask) == GOMP_MAP_ZERO_LEN_ARRAY_SECTION) { - n = gomp_map_0len_lookup (mem_map, &cur_node); + /* Defer lookup when mapped item found. */ + n = not_found_cnt ? NULL : gomp_map_0len_lookup (mem_map, &cur_node); if (!n) { tgt->list[i].key = NULL; @@ -1417,7 +1418,15 @@ gomp_map_vars_internal (struct gomp_device_descr *devicep, } continue; case GOMP_MAP_FIRSTPRIVATE_INT: + continue; case GOMP_MAP_ZERO_LEN_ARRAY_SECTION: + cur_node.host_start = (uintptr_t) hostaddrs[i]; + cur_node.host_end = cur_node.host_start + sizes[i]; + n = gomp_map_0len_lookup (mem_map, &cur_node); + if (n) + gomp_map_vars_existing (devicep, aq, n, &cur_node, + &tgt->list[i], kind & typemask, false, + implicit, NULL, refcount_set); continue; case GOMP_MAP_USE_DEVICE_PTR_IF_PRESENT: /* The OpenACC 'host_data' construct only allows 'use_device' diff --git a/libgomp/testsuite/libgomp.c++/pr108286.C b/libgomp/testsuite/libgomp.c++/pr108286.C index ee88c2f9fd0..3ae2f5dd16e 100644 --- a/libgomp/testsuite/libgomp.c++/pr108286.C +++ b/libgomp/testsuite/libgomp.c++/pr108286.C @@ -1,5 +1,6 @@ // PR c++/108286 // { dg-do run } +// { dg-additional-options "-fdump-tree-gimple" } struct S { int @@ -27,3 +28,6 @@ main () if (s.foo () != 42) __builtin_abort (); } + +/* Ensure that 'this' is mapped but only once and not additionally via 'this[:0]'. */ +/* { dg-final { scan-tree-dump "#pragma omp target num_teams\\(-2\\) thread_limit\\(0\\) map\\(tofrom:\\*this \\\[len: \[0-9\]+\\\]\\) map\\(firstprivate:this \\\[pointer assign, bias: 0\\\]\\) nowait map\\(tofrom:res \\\[len: \[0-9\]+\\\]\\) map\\(tofrom:\\*_\[0-9\]+ \\\[len: _\[0-9\]+\\\]\\) map\\(attach:this->ptr \\\[bias: 0\\\]\\)\[\r\n\]" "gimple" } } */ diff --git a/libgomp/testsuite/libgomp.c-c++-common/target-implicit-map-6.c b/libgomp/testsuite/libgomp.c-c++-common/target-implicit-map-6.c new file mode 100644 index 00000000000..6a2d3bfd0e3 --- /dev/null +++ b/libgomp/testsuite/libgomp.c-c++-common/target-implicit-map-6.c @@ -0,0 +1,212 @@ +/* Prefined firstprivate privatizes the pointer + and then updates the value to point to the corresponding + device variable, if existing. + + See PR middle-end/110639 + and TR12 in "14.8 target Construct" [379:8-10] + or OpenMP 5.1 in "2.21.7.2 Pointer Initialization for Device Data Environments". */ + +#include +#include + +int my_false = 0; + +int +f (int x, int y) +{ + return x + y; +} + +void +no_other_clause () +{ + int data = 3; + int *p = &data; + #pragma omp target enter data map(data) + #pragma omp target + *p = 5; + #pragma omp target exit data map(data) + if (*p != 5) + abort (); +} + +void +test1 (int devnum) +{ + int start = 0, n = 100; + int a[100]; + int *p = &a[0]; + + for (int i = start; i < start+n; i++) + a[i] = 10*i; + + #pragma omp target map(a) device(device_num : devnum) + { + if (my_false) /* Ensure that 'map(a)' is not optimized away. */ + a[8] = 1; + for (int i = start; i < start+n; i++) + p[i] = f(p[i], i); + p = NULL; + } + + if (p != &a[0]) + abort (); + for (int i = start; i < start+n; i++) + if (a[i] != f(10 *i, i)) + abort (); +} + + + +void +test2 (int devnum) +{ + int start = 0, n = 100; + int a[100]; + int *p = &a[0]; + + for (int i = start; i < start+n; i++) + a[i] = 10*i; + + #pragma omp target enter data map(a) device(device_num : devnum) + #pragma omp target device(device_num : devnum) + { + (void) a; /* Ensure that 'map(a)' is not optimized away. */ + for (int i = start; i < start+n; i++) + p[i] = f(p[i], i); + p = NULL; + } + #pragma omp target exit data map(a) device(device_num : devnum) + + if (p != &a[0]) + abort (); + for (int i = start; i < start+n; i++) + if (a[i] != f(10 *i, i)) + abort (); +} + +void +test3 (int devnum) +{ + int start = 8, n = 10; + int a[100]; + int *p = &a[start]; + + for (int i = start; i < start+n; i++) + a[i] = 10*i; + + /* p points to a[start] */ + #pragma omp target map(a[start:n]) device(device_num : devnum) + { + if (my_false) /* Ensure that 'map(a)' is not optimized away. */ + a[8] = 1; + for (int i = 0; i < n; i++) + p[i] = f(p[i], i + start); + p = NULL; + } + + if (p != &a[start]) + abort (); + for (int i = start; i < start+n; i++) + if (a[i] != f(10 *i, i)) + abort (); +} + +void +test4 (int devnum) +{ + int start = 8, n = 10; + int a[100]; + int *p = &a[start]; + + for (int i = start; i < start+n; i++) + a[i] = 10*i; + + /* p points to a[start] */ + #pragma omp target enter data map(a[start:n]) device(device_num : devnum) + #pragma omp target device(device_num : devnum) + { + for (int i = 0; i < n; i++) + p[i] = f(p[i], i + start); + p = NULL; + } + #pragma omp target exit data map(a[start:n]) device(device_num : devnum) + + if (p != &a[start]) + abort (); + for (int i = start; i < start+n; i++) + if (a[i] != f(10 *i, i)) + abort (); +} + +void +test5 (int devnum) +{ + int start = 8, n = 10; + int a[100]; + int *p = &a[start + 5]; + + for (int i = start; i < start+n; i++) + a[i] = 10*i; + + /* p points to a[start + 5] */ + #pragma omp target map(a[start:n]) device(device_num : devnum) + { + if (my_false) /* Ensure that 'map(a)' is not optimized away. */ + a[8] = 1; + for (int i = 0; i < n; i++) + p[i - 5] = f(p[i - 5], i + start); + p = NULL; + } + + if (p != &a[start + 5]) + abort (); + for (int i = start; i < start+n; i++) + if (a[i] != f(10 *i, i)) + abort (); +} + +void +test6 (int devnum) +{ + int start = 8, n = 10; + int a[100]; + int *p = &a[start + 5]; + + for (int i = start; i < start+n; i++) + a[i] = 10*i; + + /* p points to a[start + 5] */ + #pragma omp target enter data map(a[start:n]) device(device_num : devnum) + #pragma omp target device(device_num : devnum) + { + for (int i = 0; i < n; i++) + p[i - 5] = f(p[i - 5], i + start); + p = NULL; + } + #pragma omp target exit data map(a[start:n]) device(device_num : devnum) + + if (p != &a[start + 5]) + abort (); + for (int i = start; i < start+n; i++) + if (a[i] != f(10 *i, i)) + abort (); +} + +int +main () +{ + int n = omp_get_num_devices (); + no_other_clause (); + for (int i = omp_initial_device; i <= n; i++) + { + /* First with 'a' mapped on target; then 'a' on target enter data. */ + test1 (i); /* p = &a[0], map(a) */ + test2 (i); + test3 (i); /* p = &a[start], map(a[start:n]) */ + test4 (i); + test5 (i); /* p = &a[start + 5], map(a[start:n]) */ + test6 (i); + } + return 0; +}