From patchwork Fri Mar 4 13:41:47 2022 Content-Type: text/plain; charset="utf-8" MIME-Version: 1.0 Content-Transfer-Encoding: 8bit X-Patchwork-Submitter: Thomas Schwinge X-Patchwork-Id: 51573 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 B5F813857C6D for ; Fri, 4 Mar 2022 13:44:14 +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 572303857C56; Fri, 4 Mar 2022 13:41:57 +0000 (GMT) DMARC-Filter: OpenDMARC Filter v1.4.1 sourceware.org 572303857C56 Authentication-Results: sourceware.org; dmarc=none (p=none dis=none) header.from=codesourcery.com Authentication-Results: sourceware.org; spf=pass smtp.mailfrom=mentor.com X-IronPort-AV: E=Sophos;i="5.90,155,1643702400"; d="scan'208,223";a="72566435" Received: from orw-gwy-02-in.mentorg.com ([192.94.38.167]) by esa3.mentor.iphmx.com with ESMTP; 04 Mar 2022 05:41:57 -0800 IronPort-SDR: zlvN7o20NRjinwOjRIe+GAo3Fv/HQD6qvE/D8NUgm02CyeYGuewSJkMBytxsb9dswgaOrsiyWD Wbhh1urj/NjCMnJieeO0+y5C3v/kFIcChOmzEtPKHoZyUMPRIjp9FpkHL8lJGeyEsRmX87d9t4 twyBc4g1UoPlnO7sR2ei8clQuXaoy4L362AfSduqb5igkyifO9XfSYUbZZJP2UonzSs+k5kcm1 o7vEeEjQayI0enrvv4dF+gu57Yv8d0Rv5o1EbaTiCsj6vxrthWDNod2VRn9Y/fCrpzrQI+Xacu AjQ= From: Thomas Schwinge To: Subject: OpenACC 'kernels' decomposition: Move 'TREE_ADDRESSABLE' setting into OMP lowering [PR100280] In-Reply-To: <87iluovu47.fsf@euler.schwinge.homeip.net> References: <20190508145157.08beb4df@squid.athome> <87iluovu47.fsf@euler.schwinge.homeip.net> User-Agent: Notmuch/0.29.3+94~g74c3f1b (https://notmuchmail.org) Emacs/27.1 (x86_64-pc-linux-gnu) Date: Fri, 4 Mar 2022 14:41:47 +0100 Message-ID: <878rtpomw4.fsf@euler.schwinge.homeip.net> MIME-Version: 1.0 X-Originating-IP: [137.202.0.90] X-ClientProxiedBy: SVR-IES-MBX-07.mgc.mentorg.com (139.181.222.7) To svr-ies-mbx-01.mgc.mentorg.com (139.181.222.1) X-Spam-Status: No, score=-12.0 required=5.0 tests=BAYES_00, GIT_PATCH_0, HEADER_FROM_DIFFERENT_DOMAINS, KAM_DMARC_STATUS, KAM_SHORT, 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: , Cc: Jakub Jelinek , Julian Brown , fortran@gcc.gnu.org, asolokha@gmx.com Errors-To: gcc-patches-bounces+patchwork=sourceware.org@gcc.gnu.org Sender: "Gcc-patches" Hi! On 2022-01-13T10:54:16+0100, I wrote: > On 2019-05-08T14:51:57+0100, Julian Brown wrote: >> - The "addressable" bit is set during the kernels conversion pass for >> variables that have "create" (alloc) clauses created for them in the >> synthesised outer data region (instead of in the front-end, etc., >> where it can't be done accurately). Such variables actually have >> their address taken during transformations made in a later pass >> (omp-low, I think), but there's a phase-ordering problem that means >> the flag should be set earlier. > > The actual issue is a bit different, but yes, there is a problem. > The related ICE has also been reported as > "ICE in lower_omp_target, at omp-low.c:12287". (And I'm confused why we > didn't run into that with the OpenACC 'kernels' decomposition > originally.) I've pushed to master branch > commit 9b32c1669aad5459dd053424f9967011348add83 > "OpenACC 'kernels' decomposition: Mark variables used in synthesized data > clauses as addressable [PR100280]", see attached. > --- a/gcc/omp-oacc-kernels-decompose.cc > +++ b/gcc/omp-oacc-kernels-decompose.cc > @@ -793,7 +793,8 @@ make_data_region_try_statement (location_t loc, gimple *body) > > /* If INNER_BIND_VARS holds variables, build an OpenACC data region with > location LOC containing BODY and having 'create (var)' clauses for each > - variable. If INNER_CLEANUP is present, add a try-finally statement with > + variable (as a side effect, such variables also get TREE_ADDRESSABLE set). > + If INNER_CLEANUP is present, add a try-finally statement with > this cleanup code in the finally block. Return the new data region, or > the original BODY if no data region was needed. */ > > @@ -842,6 +843,9 @@ maybe_build_inner_data_region (location_t loc, gimple *body, > inner_data_clauses = new_clause; > > prev_mapped_var = v; > + > + /* See . */ > + TREE_ADDRESSABLE (v) = 1; > } > } Pushed to master branch commit de6e81ea961219d0726db67776d11ce75a4cae1b "OpenACC 'kernels' decomposition: Move 'TREE_ADDRESSABLE' setting into OMP lowering [PR100280]", see attached. Grüße Thomas ----------------- 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 From de6e81ea961219d0726db67776d11ce75a4cae1b Mon Sep 17 00:00:00 2001 From: Thomas Schwinge Date: Tue, 15 Feb 2022 23:03:49 +0100 Subject: [PATCH] OpenACC 'kernels' decomposition: Move 'TREE_ADDRESSABLE' setting into OMP lowering [PR100280] ... in preparation for later changes. No functional change. Follow-up to commit 9b32c1669aad5459dd053424f9967011348add83 "OpenACC 'kernels' decomposition: Mark variables used in synthesized data clauses as addressable [PR100280]". PR middle-end/100280 gcc/ * tree.h (OMP_CLAUSE_MAP_DECL_MAKE_ADDRESSABLE): New. * tree-core.h: Document it. * omp-low.cc (scan_sharing_clauses) : Handle 'OMP_CLAUSE_MAP_DECL_MAKE_ADDRESSABLE'. * omp-oacc-kernels-decompose.cc (maybe_build_inner_data_region): Set 'OMP_CLAUSE_MAP_DECL_MAKE_ADDRESSABLE' instead of 'TREE_ADDRESSABLE'. gcc/testsuite/ * c-c++-common/goacc/classify-kernels-unparallelized.c: Adjust. * c-c++-common/goacc/classify-kernels.c: Likewise. * c-c++-common/goacc/kernels-decompose-2.c: Likewise. * c-c++-common/goacc/kernels-decompose-pr100280-1.c: Likewise. * c-c++-common/goacc/kernels-decompose-pr104061-1-2.c: Likewise. * c-c++-common/goacc/kernels-decompose-pr104061-1-3.c: Likewise. * c-c++-common/goacc/kernels-decompose-pr104061-1-4.c: Likewise. * c-c++-common/goacc/kernels-decompose-pr104132-1.c: Likewise. * c-c++-common/goacc/kernels-decompose-pr104133-1.c: Likewise. libgomp/ * testsuite/libgomp.oacc-c-c++-common/f-asyncwait-1.c: Adjust. * testsuite/libgomp.oacc-c-c++-common/kernels-decompose-1.c: Likewise. --- gcc/omp-low.cc | 31 +++++++++++++++++++ gcc/omp-oacc-kernels-decompose.cc | 5 +-- .../goacc/classify-kernels-unparallelized.c | 3 +- .../c-c++-common/goacc/classify-kernels.c | 3 +- .../c-c++-common/goacc/kernels-decompose-2.c | 6 ++-- .../goacc/kernels-decompose-pr100280-1.c | 3 +- .../goacc/kernels-decompose-pr104061-1-2.c | 3 +- .../goacc/kernels-decompose-pr104061-1-3.c | 3 +- .../goacc/kernels-decompose-pr104061-1-4.c | 3 +- .../goacc/kernels-decompose-pr104132-1.c | 3 +- .../goacc/kernels-decompose-pr104133-1.c | 3 +- gcc/tree-core.h | 3 ++ gcc/tree.h | 5 +++ .../libgomp.oacc-c-c++-common/f-asyncwait-1.c | 9 ++++-- .../kernels-decompose-1.c | 3 +- 15 files changed, 70 insertions(+), 16 deletions(-) diff --git a/gcc/omp-low.cc b/gcc/omp-low.cc index 2294456b27d..6654bfd426e 100644 --- a/gcc/omp-low.cc +++ b/gcc/omp-low.cc @@ -1495,6 +1495,37 @@ scan_sharing_clauses (tree clauses, omp_context *ctx) if (ctx->outer) scan_omp_op (&OMP_CLAUSE_SIZE (c), ctx->outer); decl = OMP_CLAUSE_DECL (c); + /* If requested, make 'decl' addressable. */ + if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP + && OMP_CLAUSE_MAP_DECL_MAKE_ADDRESSABLE (c)) + { + gcc_checking_assert (DECL_P (decl)); + + gcc_checking_assert (!TREE_ADDRESSABLE (decl)); + TREE_ADDRESSABLE (decl) = 1; + + if (dump_enabled_p ()) + { + location_t loc = OMP_CLAUSE_LOCATION (c); + const dump_user_location_t d_u_loc + = dump_user_location_t::from_location_t (loc); + /* PR100695 "Format decoder, quoting in 'dump_printf' etc." */ +#if __GNUC__ >= 10 +# pragma GCC diagnostic push +# pragma GCC diagnostic ignored "-Wformat" +#endif + dump_printf_loc (MSG_NOTE, d_u_loc, + "variable %<%T%>" + " made addressable\n", + decl); +#if __GNUC__ >= 10 +# pragma GCC diagnostic pop +#endif + } + + /* Done. */ + OMP_CLAUSE_MAP_DECL_MAKE_ADDRESSABLE (c) = 0; + } /* Global variables with "omp declare target" attribute don't need to be copied, the receiver side will use them directly. However, global variables with "omp declare target link" diff --git a/gcc/omp-oacc-kernels-decompose.cc b/gcc/omp-oacc-kernels-decompose.cc index 5093386f718..ecbd3071e5d 100644 --- a/gcc/omp-oacc-kernels-decompose.cc +++ b/gcc/omp-oacc-kernels-decompose.cc @@ -847,7 +847,8 @@ maybe_build_inner_data_region (location_t loc, gimple *body, /* See . */ if (!TREE_ADDRESSABLE (v)) { - TREE_ADDRESSABLE (v) = 1; + /* Request that OMP lowering make 'v' addressable. */ + OMP_CLAUSE_MAP_DECL_MAKE_ADDRESSABLE (new_clause) = 1; if (dump_enabled_p ()) { @@ -861,7 +862,7 @@ maybe_build_inner_data_region (location_t loc, gimple *body, dump_printf_loc (MSG_NOTE, d_u_loc, "OpenACC % decomposition:" " variable %<%T%> declared in block" - " made addressable\n", + " requested to be made addressable\n", v); #if __GNUC__ >= 10 # pragma GCC diagnostic pop diff --git a/gcc/testsuite/c-c++-common/goacc/classify-kernels-unparallelized.c b/gcc/testsuite/c-c++-common/goacc/classify-kernels-unparallelized.c index 2496462f777..61871d118a9 100644 --- a/gcc/testsuite/c-c++-common/goacc/classify-kernels-unparallelized.c +++ b/gcc/testsuite/c-c++-common/goacc/classify-kernels-unparallelized.c @@ -29,7 +29,8 @@ extern unsigned int f (unsigned int); void KERNELS () { #pragma acc kernels copyin (a[0:N], b[0:N]) copyout (c[0:N]) /* { dg-line l_compute1 } */ - /* { dg-note {OpenACC 'kernels' decomposition: variable 'i' declared in block made addressable} {} { target *-*-* } l_compute1 } */ + /* { dg-note {OpenACC 'kernels' decomposition: variable 'i' declared in block requested to be made addressable} {} { target *-*-* } l_compute1 } */ + /* { dg-note {variable 'i' made addressable} {} { target *-*-* } l_compute1 } */ /* { dg-note {variable 'i' declared in block is candidate for adjusting OpenACC privatization level} {} { target *-*-* } l_compute1 } */ /* { dg-note {variable '[abc]\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_compute1 } */ /* { dg-optimized {assigned OpenACC seq loop parallelism} {} { target *-*-* } l_compute1 } */ diff --git a/gcc/testsuite/c-c++-common/goacc/classify-kernels.c b/gcc/testsuite/c-c++-common/goacc/classify-kernels.c index 3ba0411a57a..1473337986f 100644 --- a/gcc/testsuite/c-c++-common/goacc/classify-kernels.c +++ b/gcc/testsuite/c-c++-common/goacc/classify-kernels.c @@ -25,7 +25,8 @@ extern unsigned int *__restrict c; void KERNELS () { #pragma acc kernels copyin (a[0:N], b[0:N]) copyout (c[0:N]) /* { dg-line l_compute1 } */ - /* { dg-note {OpenACC 'kernels' decomposition: variable 'i' declared in block made addressable} {} { target *-*-* } l_compute1 } */ + /* { dg-note {OpenACC 'kernels' decomposition: variable 'i' declared in block requested to be made addressable} {} { target *-*-* } l_compute1 } */ + /* { dg-note {variable 'i' made addressable} {} { target *-*-* } l_compute1 } */ /* { dg-note {variable 'i' declared in block is candidate for adjusting OpenACC privatization level} {} { target *-*-* } l_compute1 } */ /* { dg-note {variable '[abc]\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_compute1 } */ /* { dg-optimized {assigned OpenACC gang loop parallelism} {} { target *-*-* } l_compute1 } */ diff --git a/gcc/testsuite/c-c++-common/goacc/kernels-decompose-2.c b/gcc/testsuite/c-c++-common/goacc/kernels-decompose-2.c index 5fbd10221a2..bf158311dae 100644 --- a/gcc/testsuite/c-c++-common/goacc/kernels-decompose-2.c +++ b/gcc/testsuite/c-c++-common/goacc/kernels-decompose-2.c @@ -66,14 +66,16 @@ main () #pragma acc kernels /* { dg-line l_compute[incr c_compute] } */ /* { dg-note {beginning 'gang-single' part in OpenACC 'kernels' region} {} { target *-*-* } l_compute$c_compute } */ - /* { dg-note {OpenACC 'kernels' decomposition: variable 'i' declared in block made addressable} {} { target *-*-* } l_compute$c_compute } */ + /* { dg-note {OpenACC 'kernels' decomposition: variable 'i' declared in block requested to be made addressable} {} { target *-*-* } l_compute$c_compute } */ + /* { dg-note {variable 'i' made addressable} {} { target *-*-* } l_compute$c_compute } */ /* { dg-note {variable 'i' declared in block is candidate for adjusting OpenACC privatization level} {} { target *-*-* } l_compute$c_compute } */ { int i; } #pragma acc kernels /* { dg-line l_compute[incr c_compute] } */ - /* { dg-note {OpenACC 'kernels' decomposition: variable 'i' declared in block made addressable} {} { target *-*-* } l_compute$c_compute } */ + /* { dg-note {OpenACC 'kernels' decomposition: variable 'i' declared in block requested to be made addressable} {} { target *-*-* } l_compute$c_compute } */ + /* { dg-note {variable 'i' made addressable} {} { target *-*-* } l_compute$c_compute } */ /* { dg-note {variable 'i' declared in block is candidate for adjusting OpenACC privatization level} {} { target *-*-* } l_compute$c_compute } */ /* { dg-optimized {assigned OpenACC gang loop parallelism} {} { target *-*-* } l_compute$c_compute } */ /* { dg-note {beginning 'parloops' part in OpenACC 'kernels' region} {} { target *-*-* } .+1 } */ diff --git a/gcc/testsuite/c-c++-common/goacc/kernels-decompose-pr100280-1.c b/gcc/testsuite/c-c++-common/goacc/kernels-decompose-pr100280-1.c index 91e589777b0..1c1e22c00ac 100644 --- a/gcc/testsuite/c-c++-common/goacc/kernels-decompose-pr100280-1.c +++ b/gcc/testsuite/c-c++-common/goacc/kernels-decompose-pr100280-1.c @@ -10,7 +10,8 @@ void foo (void) /* { dg-line l_f_1 } */ { #pragma acc kernels /* { dg-line l_k_1 } */ - /* { dg-note {OpenACC 'kernels' decomposition: variable 'i' declared in block made addressable} {} { target *-*-* } l_k_1 } */ + /* { dg-note {OpenACC 'kernels' decomposition: variable 'i' declared in block requested to be made addressable} {} { target *-*-* } l_k_1 } */ + /* { dg-note {variable 'i' made addressable} {} { target *-*-* } l_k_1 } */ /* { dg-note {variable 'i' declared in block is candidate for adjusting OpenACC privatization level} {} { target *-*-* } l_k_1 } */ /* { dg-optimized {assigned OpenACC seq loop parallelism} {} { target *-*-* } l_k_1 } */ /* { dg-bogus {note: beginning 'parloops' part in OpenACC 'kernels' region} {TODO location} { xfail *-*-* } l_f_1 } diff --git a/gcc/testsuite/c-c++-common/goacc/kernels-decompose-pr104061-1-2.c b/gcc/testsuite/c-c++-common/goacc/kernels-decompose-pr104061-1-2.c index 8ca02cb2119..336cf2ad425 100644 --- a/gcc/testsuite/c-c++-common/goacc/kernels-decompose-pr104061-1-2.c +++ b/gcc/testsuite/c-c++-common/goacc/kernels-decompose-pr104061-1-2.c @@ -16,7 +16,8 @@ foo (void) { /* { dg-bogus {sorry, unimplemented: 'gimple_debug' not yet supported} TODO { xfail *-*-* } .+1 } */ #pragma acc kernels /* { dg-line l_compute1 } */ - /* { dg-bogus {note: OpenACC 'kernels' decomposition: variable 'k' declared in block made addressable} {w/ debug} { xfail *-*-* } l_compute1 } */ + /* { dg-bogus {note: OpenACC 'kernels' decomposition: variable 'k' declared in block requested to be made addressable} {w/ debug} { xfail *-*-* } l_compute1 } */ + /* { dg-bogus {note: variable 'k' made addressable} {w/ debug} { xfail *-*-* } l_compute1 } */ /* { dg-note {variable 'arr_0\.0' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} {} { xfail *-*-* } l_compute1 } */ { /* { dg-bogus {note: beginning 'gang-single' part in OpenACC 'kernels' region} {w/ debug} { xfail c++ } .-1 } diff --git a/gcc/testsuite/c-c++-common/goacc/kernels-decompose-pr104061-1-3.c b/gcc/testsuite/c-c++-common/goacc/kernels-decompose-pr104061-1-3.c index 05a196d8f36..f41dda86122 100644 --- a/gcc/testsuite/c-c++-common/goacc/kernels-decompose-pr104061-1-3.c +++ b/gcc/testsuite/c-c++-common/goacc/kernels-decompose-pr104061-1-3.c @@ -23,7 +23,8 @@ foo (void) { /* { dg-bogus {sorry, unimplemented: 'gimple_debug' not yet supported} {} { target *-*-* } .+1 } suppressed via '-fcompare-debug'. */ #pragma acc kernels /* { dg-line l_compute1 } */ - /* { dg-bogus {note: OpenACC 'kernels' decomposition: variable 'k' declared in block made addressable} {w/ debug} { xfail *-*-* } l_compute1 } */ + /* { dg-bogus {note: OpenACC 'kernels' decomposition: variable 'k' declared in block requested to be made addressable} {w/ debug} { xfail *-*-* } l_compute1 } */ + /* { dg-bogus {note: variable 'k' made addressable} {w/ debug} { xfail *-*-* } l_compute1 } */ /* { dg-bogus {note: variable 'k' declared in block is candidate for adjusting OpenACC privatization level} {w/ debug} { xfail *-*-* } l_compute1 } */ /* { dg-note {variable 'arr_0\.0' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_compute1 } */ { diff --git a/gcc/testsuite/c-c++-common/goacc/kernels-decompose-pr104061-1-4.c b/gcc/testsuite/c-c++-common/goacc/kernels-decompose-pr104061-1-4.c index 07cb592bcf9..cde95a7b7ac 100644 --- a/gcc/testsuite/c-c++-common/goacc/kernels-decompose-pr104061-1-4.c +++ b/gcc/testsuite/c-c++-common/goacc/kernels-decompose-pr104061-1-4.c @@ -21,7 +21,8 @@ foo (void) { /* { dg-bogus {sorry, unimplemented: 'gimple_debug' not yet supported} {} { target *-*-* } .+1 } suppressed via '-fcompare-debug'. */ #pragma acc kernels /* { dg-line l_compute1 } */ - /* { dg-bogus {note: OpenACC 'kernels' decomposition: variable 'k' declared in block made addressable} {w/ debug} { xfail *-*-* } l_compute1 } */ + /* { dg-bogus {note: OpenACC 'kernels' decomposition: variable 'k' declared in block requested to be made addressable} {w/ debug} { xfail *-*-* } l_compute1 } */ + /* { dg-bogus {note: variable 'k' made addressable} {w/ debug} { xfail *-*-* } l_compute1 } */ /* { dg-bogus {note: variable 'k' declared in block is candidate for adjusting OpenACC privatization level} {w/ debug} { xfail *-*-* } l_compute1 } */ /* { dg-note {variable 'arr_0\.0' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_compute1 } */ { diff --git a/gcc/testsuite/c-c++-common/goacc/kernels-decompose-pr104132-1.c b/gcc/testsuite/c-c++-common/goacc/kernels-decompose-pr104132-1.c index 5b0fe42efbf..4f38a83bb19 100644 --- a/gcc/testsuite/c-c++-common/goacc/kernels-decompose-pr104132-1.c +++ b/gcc/testsuite/c-c++-common/goacc/kernels-decompose-pr104132-1.c @@ -18,7 +18,8 @@ void foo (void) { #pragma acc kernels /* { dg-line l_compute1 } */ - /* { dg-note {OpenACC 'kernels' decomposition: variable 'k' declared in block made addressable} {} { target *-*-* } l_compute1 } */ + /* { dg-note {OpenACC 'kernels' decomposition: variable 'k' declared in block requested to be made addressable} {} { target *-*-* } l_compute1 } */ + /* { dg-note {variable 'k' made addressable} {} { target *-*-* } l_compute1 } */ /* { dg-note {variable 'k' declared in block is candidate for adjusting OpenACC privatization level} {} { target *-*-* } l_compute1 } */ { int k; diff --git a/gcc/testsuite/c-c++-common/goacc/kernels-decompose-pr104133-1.c b/gcc/testsuite/c-c++-common/goacc/kernels-decompose-pr104133-1.c index 4536b5c179e..0499665777d 100644 --- a/gcc/testsuite/c-c++-common/goacc/kernels-decompose-pr104133-1.c +++ b/gcc/testsuite/c-c++-common/goacc/kernels-decompose-pr104133-1.c @@ -18,7 +18,8 @@ void foo (void) { #pragma acc kernels /* { dg-line l_compute1 } */ - /* { dg-note {OpenACC 'kernels' decomposition: variable 'k' declared in block made addressable} {} { target *-*-* } l_compute1 } */ + /* { dg-note {OpenACC 'kernels' decomposition: variable 'k' declared in block requested to be made addressable} {} { target *-*-* } l_compute1 } */ + /* { dg-note {variable 'k' made addressable} {} { target *-*-* } l_compute1 } */ /* { dg-note {variable 'k' declared in block is candidate for adjusting OpenACC privatization level} {} { target *-*-* } l_compute1 } */ /* { dg-note {variable 'arr_0\.0' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_compute1 } */ /* { dg-note {variable 'arr_0\.1' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_compute1 } */ diff --git a/gcc/tree-core.h b/gcc/tree-core.h index 01a1ce499da..4530bd8c2c7 100644 --- a/gcc/tree-core.h +++ b/gcc/tree-core.h @@ -1155,6 +1155,9 @@ struct GTY(()) tree_base { PREDICT_EXPR_OUTCOME in PREDICT_EXPR + OMP_CLAUSE_MAP_DECL_MAKE_ADDRESSABLE in + OMP_CLAUSE + static_flag: TREE_STATIC in diff --git a/gcc/tree.h b/gcc/tree.h index 36ceed57064..da6f3b38ba5 100644 --- a/gcc/tree.h +++ b/gcc/tree.h @@ -1699,6 +1699,11 @@ class auto_suppress_location_wrappers #define OMP_CLAUSE_MAP_RUNTIME_IMPLICIT_P(NODE) \ (OMP_CLAUSE_SUBCODE_CHECK (NODE, OMP_CLAUSE_MAP)->base.deprecated_flag) +/* Flag that 'OMP_CLAUSE_DECL (NODE)' is to be made addressable during OMP + lowering. */ +#define OMP_CLAUSE_MAP_DECL_MAKE_ADDRESSABLE(NODE) \ + (OMP_CLAUSE_SUBCODE_CHECK (NODE, OMP_CLAUSE_MAP)->base.addressable_flag) + /* True on an OMP_CLAUSE_USE_DEVICE_PTR with an OpenACC 'if_present' clause. */ #define OMP_CLAUSE_USE_DEVICE_PTR_IF_PRESENT(NODE) \ diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/f-asyncwait-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/f-asyncwait-1.c index d360aad4736..9f7b8b4af2f 100644 --- a/libgomp/testsuite/libgomp.oacc-c-c++-common/f-asyncwait-1.c +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/f-asyncwait-1.c @@ -260,7 +260,8 @@ main (void) #pragma acc data copy (a[0:N]) copy (b[0:N]) copy (c[0:N]) copy (d[0:N]) { #pragma acc kernels async (1) /* { dg-line l_compute[incr c_compute] } */ - /* { dg-note {OpenACC 'kernels' decomposition: variable 'i' declared in block made addressable} "" { target *-*-* } l_compute$c_compute } */ + /* { dg-note {OpenACC 'kernels' decomposition: variable 'i' declared in block requested to be made addressable} "" { target *-*-* } l_compute$c_compute } */ + /* { dg-note {variable 'i' made addressable} {} { target *-*-* } l_compute$c_compute } */ /* { dg-note {variable 'i' declared in block is candidate for adjusting OpenACC privatization level} "" { target *-*-* } l_compute$c_compute } */ /* { dg-optimized "assigned OpenACC seq loop parallelism" "" { target { ! __OPTIMIZE__ } } l_compute$c_compute } { dg-optimized "assigned OpenACC gang loop parallelism" "" { target { __OPTIMIZE__ } } l_compute$c_compute } */ @@ -269,7 +270,8 @@ main (void) b[i] = (a[i] * a[i] * a[i]) / a[i]; #pragma acc kernels async (1) /* { dg-line l_compute[incr c_compute] } */ - /* { dg-note {OpenACC 'kernels' decomposition: variable 'i' declared in block made addressable} "" { target *-*-* } l_compute$c_compute } */ + /* { dg-note {OpenACC 'kernels' decomposition: variable 'i' declared in block requested to be made addressable} "" { target *-*-* } l_compute$c_compute } */ + /* { dg-note {variable 'i' made addressable} {} { target *-*-* } l_compute$c_compute } */ /* { dg-note {variable 'i' declared in block is candidate for adjusting OpenACC privatization level} "" { target *-*-* } l_compute$c_compute } */ /* { dg-optimized "assigned OpenACC seq loop parallelism" "" { target { ! __OPTIMIZE__ } } l_compute$c_compute } { dg-optimized "assigned OpenACC gang loop parallelism" "" { target { __OPTIMIZE__ } } l_compute$c_compute } */ @@ -313,7 +315,8 @@ main (void) #pragma acc data copy (a[0:N], b[0:N], c[0:N], d[0:N], e[0:N]) { #pragma acc kernels async (1) /* { dg-line l_compute[incr c_compute] } */ - /* { dg-note {OpenACC 'kernels' decomposition: variable 'i' declared in block made addressable} "" { target *-*-* } l_compute$c_compute } */ + /* { dg-note {OpenACC 'kernels' decomposition: variable 'i' declared in block requested to be made addressable} "" { target *-*-* } l_compute$c_compute } */ + /* { dg-note {variable 'i' made addressable} {} { target *-*-* } l_compute$c_compute } */ /* { dg-note {variable 'i' declared in block is candidate for adjusting OpenACC privatization level} "" { target *-*-* } l_compute$c_compute } */ /* { dg-optimized "assigned OpenACC seq loop parallelism" "" { target { ! __OPTIMIZE__ } } l_compute$c_compute } { dg-optimized "assigned OpenACC gang loop parallelism" "" { target { __OPTIMIZE__ } } l_compute$c_compute } */ diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-decompose-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-decompose-1.c index b99497ea6c5..85c39871f94 100644 --- a/libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-decompose-1.c +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-decompose-1.c @@ -32,7 +32,8 @@ int main() { /* { dg-note {beginning 'gang-single' part in OpenACC 'kernels' region} {} { target *-*-* } .+1 } */ int c = 234; - /* { dg-note {OpenACC 'kernels' decomposition: variable 'c' declared in block made addressable} "" { target *-*-* } l_compute$c_compute } */ + /* { dg-note {OpenACC 'kernels' decomposition: variable 'c' declared in block requested to be made addressable} "" { target *-*-* } l_compute$c_compute } */ + /* { dg-note {variable 'c' made addressable} {} { target *-*-* } l_compute$c_compute } */ /* { dg-note {variable 'c' declared in block is candidate for adjusting OpenACC privatization level} "" { target *-*-* } l_compute$c_compute } */ #pragma acc loop independent gang /* { dg-line l_loop_i[incr c_loop_i] } */ -- 2.34.1