From patchwork Sat Mar 12 12:42:52 2022 Content-Type: text/plain; charset="utf-8" MIME-Version: 1.0 Content-Transfer-Encoding: 8bit X-Patchwork-Submitter: Thomas Schwinge X-Patchwork-Id: 51911 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 0B31E385E007 for ; Sat, 12 Mar 2022 12:43:24 +0000 (GMT) X-Original-To: gcc-patches@gcc.gnu.org Delivered-To: gcc-patches@gcc.gnu.org Received: from esa1.mentor.iphmx.com (esa1.mentor.iphmx.com [68.232.129.153]) by sourceware.org (Postfix) with ESMTPS id ACC793858C78 for ; Sat, 12 Mar 2022 12:43:05 +0000 (GMT) DMARC-Filter: OpenDMARC Filter v1.4.1 sourceware.org ACC793858C78 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,175,1643702400"; d="scan'208,223";a="75653590" Received: from orw-gwy-01-in.mentorg.com ([192.94.38.165]) by esa1.mentor.iphmx.com with ESMTP; 12 Mar 2022 04:43:04 -0800 IronPort-SDR: dAnQPYDQrfoAXcENsxk1c6POaXhmRfxvhYdv7oTggoMZ2EpUPWfEZiaq6wi5l1tRgPJhCc/fwI 87236lHLkATFzr2cYmsV1Pfrz9ANumkLHZPN/kNaV6od2UeiD7VXQvfVVr3FxIBeLYIRYW8p1m iD8jBe6TcTfurlhF4hdKnvaic0HXi6pt75EtrDmkmlxJFQoCbTrrhNix/6KIEZ/rN/TJgcsT/J 8uft25okt7qOUQJo45o/AQd4fQMGf3CNtoza6IOHS+mgczBLVF9pWZyPhNYHeX0vmCVQfxjKtD 15U= From: Thomas Schwinge To: Subject: OpenACC 'kernels' decomposition: Mark variables used in 'present' clauses as addressable [PR100280, PR104086] In-Reply-To: <87ilsjmjld.fsf@euler.schwinge.homeip.net> References: <877dqodhi1.fsf@euler.schwinge.homeip.net> <87ilsjmjld.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: Sat, 12 Mar 2022 13:42:52 +0100 Message-ID: <87fsnnmjeb.fsf@euler.schwinge.homeip.net> MIME-Version: 1.0 X-Originating-IP: [137.202.0.90] X-ClientProxiedBy: svr-ies-mbx-15.mgc.mentorg.com (139.181.222.15) To svr-ies-mbx-01.mgc.mentorg.com (139.181.222.1) X-Spam-Status: No, score=-11.9 required=5.0 tests=BAYES_00, GIT_PATCH_0, HEADER_FROM_DIFFERENT_DOMAINS, KAM_DMARC_STATUS, KAM_LOTSOFHASH, 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 , Arseny Solokha Errors-To: gcc-patches-bounces+patchwork=sourceware.org@gcc.gnu.org Sender: "Gcc-patches" Hi! On 2022-03-12T13:38:38+0100, I wrote: > On 2020-11-13T23:22:30+0100, I wrote: >> On 2019-02-01T00:59:30+0100, I wrote: >>> I've just pushed the attached nine patches to openacc-gcc-8-branch: >>> OpenACC 'kernels' construct changes: splitting of the construct into >>> several regions. >> >> Now, slightly more polished, I've pushed to master branch a variant of >> most of these patches combined in commit >> e898ce7997733c29dcab9c3c62ca102c7f9fa6eb "Decompose OpenACC 'kernels' >> constructs into parts, a sequence of compute constructs", see attached. >> >>> There's more work to be done there, and we're aware of a number of TODO >>> items, but nevertheless: it's a good first step. >> >> That's still the case... :-) > >> --- /dev/null >> +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/declare-vla-kernels-decompose-ice-1.c >> @@ -0,0 +1,8 @@ >> +/* { dg-additional-options "-fopenacc-kernels=decompose" } */ >> +/* Hopefully, this is the same issue as '../../../gcc/testsuite/c-c++-common/goacc/kernels-decompose-ice-1.c'. (Related, but not the same.) >> + { dg-ice "TODO" } >> + TODO { dg-prune-output "during GIMPLE pass: omplower" } >> + TODO { dg-do link } */ >> + >> +#undef KERNELS_DECOMPOSE_ICE_HACK >> +#include "declare-vla.c" > > Arseny had later reduced that, and filed . > To document the status quo, pushed to master branch > commit 9781ae3a254a8c17ef4ffa70f21ed1728ff3c707 > "Add 'c-c++-common/goacc/kernels-decompose-pr104086-1.c' [PR104086]" >> --- /dev/null >> +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/declare-vla-kernels-decompose.c >> @@ -0,0 +1,6 @@ >> +/* { dg-additional-options "-fopenacc-kernels=decompose" } */ >> + >> +/* See also 'declare-vla-kernels-decompose-ice-1.c'. */ >> + >> +#define KERNELS_DECOMPOSE_ICE_HACK >> +#include "declare-vla.c" >> --- a/libgomp/testsuite/libgomp.oacc-c-c++-common/declare-vla.c >> +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/declare-vla.c >> @@ -38,6 +38,12 @@ f_data (void) >> for (i = 0; i < N; i++) >> A[i] = -i; >> >> + /* See 'declare-vla-kernels-decompose.c'. */ >> +#ifdef KERNELS_DECOMPOSE_ICE_HACK >> + (volatile int *) &i; >> + (volatile int *) &N; >> +#endif >> + >> # pragma acc kernels >> for (i = 0; i < N; i++) >> A[i] = i; Pushed to master branch commit 337ed336d7dd83526891bdb436f0bfe9e351f69d "OpenACC 'kernels' decomposition: Mark variables used in 'present' clauses as addressable [PR100280, PR104086]", 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 337ed336d7dd83526891bdb436f0bfe9e351f69d Mon Sep 17 00:00:00 2001 From: Thomas Schwinge Date: Thu, 17 Feb 2022 14:18:57 +0100 Subject: [PATCH] OpenACC 'kernels' decomposition: Mark variables used in 'present' clauses as addressable [PR100280, PR104086] ... like in recent commit 9b32c1669aad5459dd053424f9967011348add83 "OpenACC 'kernels' decomposition: Mark variables used in synthesized data clauses as addressable [PR100280]". Otherwise, we may run into 'gcc/omp-low.cc:lower_omp_target': 13125 else if (is_gimple_reg (var)) 13126 { 13127 gcc_assert (offloaded); PR middle-end/100280 PR middle-end/104086 gcc/ * omp-oacc-kernels-decompose.cc (omp_oacc_kernels_decompose_1): Mark variables used in 'present' clauses as addressable. * omp-low.cc (scan_sharing_clauses) : Gracefully handle duplicate 'OMP_CLAUSE_MAP_DECL_MAKE_ADDRESSABLE'. gcc/testsuite/ * c-c++-common/goacc/kernels-decompose-pr104086-1.c: Adjust, extend. libgomp/ * testsuite/libgomp.oacc-c-c++-common/declare-vla-kernels-decompose-ice-1.c: Merge this... * testsuite/libgomp.oacc-c-c++-common/declare-vla-kernels-decompose.c: ..., and this... * testsuite/libgomp.oacc-c-c++-common/declare-vla.c: ... into this, and adjust. * testsuite/libgomp.oacc-c-c++-common/kernels-decompose-1.c: Extend. --- gcc/omp-low.cc | 27 +++++--- gcc/omp-oacc-kernels-decompose.cc | 32 +++++++++ .../goacc/kernels-decompose-pr104086-1.c | 37 +++++++++-- .../declare-vla-kernels-decompose-ice-1.c | 22 ------- .../declare-vla-kernels-decompose.c | 29 -------- .../libgomp.oacc-c-c++-common/declare-vla.c | 38 ++++++----- .../kernels-decompose-1.c | 66 ++++++++++++++++++- 7 files changed, 168 insertions(+), 83 deletions(-) delete mode 100644 libgomp/testsuite/libgomp.oacc-c-c++-common/declare-vla-kernels-decompose-ice-1.c delete mode 100644 libgomp/testsuite/libgomp.oacc-c-c++-common/declare-vla-kernels-decompose.c diff --git a/gcc/omp-low.cc b/gcc/omp-low.cc index d932d74cb03..cfc63d6a104 100644 --- a/gcc/omp-low.cc +++ b/gcc/omp-low.cc @@ -1501,11 +1501,14 @@ scan_sharing_clauses (tree clauses, omp_context *ctx) { gcc_checking_assert (DECL_P (decl)); - gcc_checking_assert (!TREE_ADDRESSABLE (decl)); - if (!make_addressable_vars) - make_addressable_vars = BITMAP_ALLOC (NULL); - bitmap_set_bit (make_addressable_vars, DECL_UID (decl)); - TREE_ADDRESSABLE (decl) = 1; + bool decl_addressable = TREE_ADDRESSABLE (decl); + if (!decl_addressable) + { + if (!make_addressable_vars) + make_addressable_vars = BITMAP_ALLOC (NULL); + bitmap_set_bit (make_addressable_vars, DECL_UID (decl)); + TREE_ADDRESSABLE (decl) = 1; + } if (dump_enabled_p ()) { @@ -1517,10 +1520,16 @@ scan_sharing_clauses (tree clauses, omp_context *ctx) # 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 (!decl_addressable) + dump_printf_loc (MSG_NOTE, d_u_loc, + "variable %<%T%>" + " made addressable\n", + decl); + else + dump_printf_loc (MSG_NOTE, d_u_loc, + "variable %<%T%>" + " already made addressable\n", + decl); #if __GNUC__ >= 10 # pragma GCC diagnostic pop #endif diff --git a/gcc/omp-oacc-kernels-decompose.cc b/gcc/omp-oacc-kernels-decompose.cc index ecbd3071e5d..40b04539894 100644 --- a/gcc/omp-oacc-kernels-decompose.cc +++ b/gcc/omp-oacc-kernels-decompose.cc @@ -1468,6 +1468,38 @@ omp_oacc_kernels_decompose_1 (gimple *kernels_stmt) /* Now that this data is mapped, turn the data clause on the inner OpenACC 'kernels' into a 'present' clause. */ OMP_CLAUSE_SET_MAP_KIND (c, GOMP_MAP_FORCE_PRESENT); + + /* See , + . */ + if (DECL_P (decl) + && !TREE_ADDRESSABLE (decl)) + { + /* Request that OMP lowering make 'decl' addressable. */ + OMP_CLAUSE_MAP_DECL_MAKE_ADDRESSABLE (new_clause) = 1; + + if (dump_enabled_p ()) + { + location_t loc = OMP_CLAUSE_LOCATION (new_clause); + 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, + "OpenACC % decomposition:" + " variable %<%T%> in %qs clause" + " requested to be made addressable\n", + decl, + user_omp_clause_code_name (new_clause, true)); +#if __GNUC__ >= 10 +# pragma GCC diagnostic pop +#endif + } + } } break; diff --git a/gcc/testsuite/c-c++-common/goacc/kernels-decompose-pr104086-1.c b/gcc/testsuite/c-c++-common/goacc/kernels-decompose-pr104086-1.c index eab10cf6c72..83fb75e28b2 100644 --- a/gcc/testsuite/c-c++-common/goacc/kernels-decompose-pr104086-1.c +++ b/gcc/testsuite/c-c++-common/goacc/kernels-decompose-pr104086-1.c @@ -1,8 +1,5 @@ -/* Reduced from 'libgomp.oacc-c-c++-common/declare-vla.c'. */ - -/* { dg-additional-options "-fchecking" } - { dg-ice TODO } - { dg-prune-output {during GIMPLE pass: omplower} } */ +/* Reduced from 'libgomp.oacc-c-c++-common/declare-vla.c', and then + extended. */ /* { dg-additional-options "--param openacc-kernels=decompose" } */ @@ -14,12 +11,38 @@ void foo (void) { #pragma acc data /* { dg-line l_data1 } */ - /* { dg-bogus {note: variable 'i' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} {TODO 'data'} { xfail *-*-* } l_data1 } */ + /* { dg-bogus {note: variable 'i' declared in block is candidate for adjusting OpenACC privatization level} {TODO 'data'} { xfail *-*-* } l_data1 } */ { int i; -#pragma acc kernels +#pragma acc kernels /* { dg-line l_compute1 } */ + /* { dg-note {OpenACC 'kernels' decomposition: variable 'i' in 'copy' clause requested to be made addressable} {} { target *-*-* } l_compute1 } + { dg-note {variable 'i' made addressable} {} { target *-*-* } l_compute1 } */ /* { dg-note {beginning 'gang-single' part in OpenACC 'kernels' region} {} { target *-*-* } .+1 } */ i = 0; + +#pragma acc kernels /* { dg-line l_compute2 } */ + /* { dg-note {OpenACC 'kernels' decomposition: variable 'i' in 'copy' clause requested to be made addressable} {} { target *-*-* } l_compute2 } + { dg-note {variable 'i' already made addressable} {} { target *-*-* } l_compute2 } */ + /* { dg-note {beginning 'gang-single' part in OpenACC 'kernels' region} {} { target *-*-* } .+1 } */ + i = -1; } } + +void +foo2 (void) +{ + int i[1]; + +#pragma acc kernels /* { dg-line l2_compute1 } */ + /* { dg-note {OpenACC 'kernels' decomposition: variable 'i' in 'copy' clause requested to be made addressable} {} { target *-*-* } l2_compute1 } + { dg-note {variable 'i' made addressable} {} { target *-*-* } l2_compute1 } */ + /* { dg-note {beginning 'gang-single' part in OpenACC 'kernels' region} {} { target *-*-* } .+1 } */ + i[0] = 0; + +#pragma acc kernels /* { dg-line l2_compute2 } */ + /* { dg-note {OpenACC 'kernels' decomposition: variable 'i' in 'copy' clause requested to be made addressable} {} { target *-*-* } l2_compute2 } + { dg-note {variable 'i' already made addressable} {} { target *-*-* } l2_compute2 } */ + /* { dg-note {beginning 'gang-single' part in OpenACC 'kernels' region} {} { target *-*-* } .+1 } */ + i[0] = -1; +} diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/declare-vla-kernels-decompose-ice-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/declare-vla-kernels-decompose-ice-1.c deleted file mode 100644 index 3e5b6bab233..00000000000 --- a/libgomp/testsuite/libgomp.oacc-c-c++-common/declare-vla-kernels-decompose-ice-1.c +++ /dev/null @@ -1,22 +0,0 @@ -/* { dg-additional-options "--param=openacc-kernels=decompose" } */ -/* ICE similar to PR100280, but not the same. - { dg-ice "TODO" } - TODO { dg-prune-output "during GIMPLE pass: omplower" } - TODO { dg-do link } */ - -/* { dg-additional-options "-fopt-info-omp-all" } - { dg-additional-options "-foffload=-fopt-info-all-omp" } */ - -/* { dg-additional-options "--param=openacc-privatization=noisy" } - { dg-additional-options "-foffload=--param=openacc-privatization=noisy" } - Prune a few: uninteresting, and potentially varying depending on GCC configuration (data types): - { dg-prune-output {note: variable 'D\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} } */ - -#undef KERNELS_DECOMPOSE_ICE_HACK -#include "declare-vla.c" - -/* { dg-note {beginning 'parloops' part in OpenACC 'kernels' region} {} { target *-*-* } 27 } */ - -/* { dg-note {beginning 'parloops' part in OpenACC 'kernels' region} {} { target *-*-* } 61 } */ - -/* { dg-bogus {note: variable [^\n\r]+ candidate for adjusting OpenACC privatization level} {TODO 'data'} { xfail *-*-* } 42 } */ diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/declare-vla-kernels-decompose.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/declare-vla-kernels-decompose.c deleted file mode 100644 index 142aceec9cd..00000000000 --- a/libgomp/testsuite/libgomp.oacc-c-c++-common/declare-vla-kernels-decompose.c +++ /dev/null @@ -1,29 +0,0 @@ -/* { dg-additional-options "--param=openacc-kernels=decompose" } */ - -/* See also 'declare-vla-kernels-decompose-ice-1.c'. */ - -/* { dg-additional-options "-fopt-info-omp-all" } - { dg-additional-options "-foffload=-fopt-info-all-omp" } */ - -/* { dg-additional-options "--param=openacc-privatization=noisy" } - { dg-additional-options "-foffload=--param=openacc-privatization=noisy" } - Prune a few: uninteresting, and potentially varying depending on GCC configuration (data types): - { dg-prune-output {note: variable 'D\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} } */ - -#define KERNELS_DECOMPOSE_ICE_HACK -#include "declare-vla.c" - -/* { dg-note {beginning 'parloops' part in OpenACC 'kernels' region} {} { target *-*-* } 27 } */ - -/* { dg-note {beginning 'parloops' part in OpenACC 'kernels' region} {} { target *-*-* } 61 } */ - -/* { dg-bogus {note: variable [^\n\r]+ candidate for adjusting OpenACC privatization level} {TODO 'data'} { xfail *-*-* } 42 } */ - -/* { dg-note {variable 'i\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } 58 } - { dg-note {variable 'N\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } 58 } */ - -/* { dg-optimized {assigned OpenACC seq loop parallelism} {} { target { ! __OPTIMIZE__ } } 24 } - { dg-optimized {assigned OpenACC gang loop parallelism} {} { target { __OPTIMIZE__ } } 24 } */ - -/* { dg-optimized {assigned OpenACC seq loop parallelism} {} { target { ! __OPTIMIZE__ } } 58 } - { dg-optimized {assigned OpenACC gang loop parallelism} {} { target { __OPTIMIZE__ } } 58 } */ diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/declare-vla.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/declare-vla.c index 4ce2e6d1f18..f6fc3ffefa4 100644 --- a/libgomp/testsuite/libgomp.oacc-c-c++-common/declare-vla.c +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/declare-vla.c @@ -1,5 +1,7 @@ /* Verify OpenACC 'declare' with VLAs. */ +/* { dg-additional-options "--param=openacc-kernels=decompose" } */ + /* { dg-additional-options "-fopt-info-omp-all" } { dg-additional-options "-foffload=-fopt-info-all-omp" } */ @@ -8,6 +10,15 @@ Prune a few: uninteresting, and potentially varying depending on GCC configuration (data types): { dg-prune-output {note: variable 'D\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} } */ +/* It's only with Tcl 8.5 (released in 2007) that "the variable 'varName' + passed to 'incr' may be unset, and in that case, it will be set to [...]", + so to maintain compatibility with earlier Tcl releases, we manually + initialize counter variables: + { dg-line l_dummy[variable c_compute 0] } + { dg-message "dummy" "" { target iN-VAl-Id } l_dummy } to avoid + "WARNING: dg-line var l_dummy defined, but not used". */ + + #include @@ -21,9 +32,10 @@ f (void) for (i = 0; i < N; i++) A[i] = -i; -#pragma acc kernels - /* { dg-optimized {assigned OpenACC seq loop parallelism} {} { target { ! __OPTIMIZE__ } } .-1 } - { dg-optimized {assigned OpenACC gang loop parallelism} {} { target { __OPTIMIZE__ } } .-2 } */ +#pragma acc kernels /* { dg-line l_compute[incr 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 } */ + /* { dg-note {beginning 'parloops' part in OpenACC 'kernels' region} {} { target *-*-* } .+1 } */ for (i = 0; i < N; i++) A[i] = i; @@ -49,15 +61,14 @@ f_data (void) for (i = 0; i < N; i++) A[i] = -i; - /* See 'declare-vla-kernels-decompose.c'. */ -#ifdef KERNELS_DECOMPOSE_ICE_HACK - (volatile int *) &i; - (volatile int *) &N; -#endif - -# pragma acc kernels - /* { dg-optimized {assigned OpenACC seq loop parallelism} {} { target { ! __OPTIMIZE__ } } .-1 } - { dg-optimized {assigned OpenACC gang loop parallelism} {} { target { __OPTIMIZE__ } } .-2 } */ +# pragma acc kernels /* { dg-line l_compute[incr c_compute] } */ + /* { dg-note {OpenACC 'kernels' decomposition: variable 'i' in 'copy' clause requested to be made addressable} {} { target *-*-* } l_compute$c_compute } + { dg-note {variable 'i' made addressable} {} { target *-*-* } l_compute$c_compute } */ + /* { dg-note {OpenACC 'kernels' decomposition: variable 'N' in 'copy' clause requested to be made addressable} {} { target *-*-* } l_compute$c_compute } + { dg-note {variable 'N' made addressable} {} { 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 } */ + /* { dg-note {beginning 'parloops' part in OpenACC 'kernels' region} {} { target *-*-* } .+1 } */ for (i = 0; i < N; i++) A[i] = i; @@ -78,6 +89,3 @@ main () return 0; } - - -/* { dg-note dummy "" { target n-on-e } } to disable 'prune_notes'. */ 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 40786c750d1..eb424776b6b 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 @@ -24,7 +24,9 @@ static int g1; static int g2; -int main() +/* PR100280, etc. */ + +static void f1 () { int a = 0; /*TODO Without making 'a' addressable, for GCN offloading we will not see the expected value copied out. (But it does work for nvptx offloading, strange...) */ @@ -153,5 +155,67 @@ int main() assert (g2 == N * (N + 1) / 2); assert (f1 == 2432902008176640000ULL); +#undef N +} + + +/* PR104086 */ + +static void f2 () +{ +#pragma acc data + /* { dg-bogus {note: variable [^\n\r]+ candidate for adjusting OpenACC privatization level} {TODO 'data'} { xfail *-*-* } .-1 } */ + { + int i; + +#pragma acc kernels /* { dg-line l_compute[incr c_compute] } */ + /* { dg-note {OpenACC 'kernels' decomposition: variable 'i' in 'copy' clause requested to be made addressable} {} { target *-*-* } l_compute$c_compute } + { dg-note {variable 'i' made addressable} {} { target *-*-* } l_compute$c_compute } */ + /* { dg-note {beginning 'gang-single' part in OpenACC 'kernels' region} {} { target *-*-* } .+1 } */ + i = 1; + + assert (i == 1); + +#pragma acc kernels /* { dg-line l_compute[incr c_compute] } */ + /* { dg-note {OpenACC 'kernels' decomposition: variable 'i' in 'copy' clause requested to be made addressable} {} { target *-*-* } l_compute$c_compute } + { dg-note {variable 'i' already made addressable} {} { target *-*-* } l_compute$c_compute } */ + /* { dg-note {beginning 'gang-single' part in OpenACC 'kernels' region} {} { target *-*-* } .+1 } */ + i = -1; + + assert (i == -1); + } + + + int ia[1]; + +#pragma acc kernels /* { dg-line l_compute[incr c_compute] } */ + /* { dg-note {OpenACC 'kernels' decomposition: variable 'ia' in 'copy' clause requested to be made addressable} {} { target *-*-* } l_compute$c_compute } + { dg-note {variable 'ia' made addressable} {} { target *-*-* } l_compute$c_compute } */ + /* { dg-note {beginning 'gang-single' part in OpenACC 'kernels' region} {} { target *-*-* } .+1 } */ + ia[0] = -2; + + assert (ia[0] == -2); + +#pragma acc kernels /* { dg-line l_compute[incr c_compute] } */ + /* { dg-note {OpenACC 'kernels' decomposition: variable 'ia' in 'copy' clause requested to be made addressable} {} { target *-*-* } l_compute$c_compute } + { dg-note {variable 'ia' already 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 *-*-* } l_compute$c_compute } */ + /* { dg-note {beginning 'parloops' part in OpenACC 'kernels' region} {} { target *-*-* } .+1 } */ + for (int i = 0; i < 100; ++i) + ++ia[0]; + + assert (ia[0] == -2 + 100); +} + + +int main() +{ + f1 (); + + f2 (); + return 0; } -- 2.34.1