From patchwork Wed Jan 19 22:29:18 2022 Content-Type: text/plain; charset="utf-8" MIME-Version: 1.0 Content-Transfer-Encoding: 8bit X-Patchwork-Submitter: Thomas Schwinge X-Patchwork-Id: 50252 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 D482B3857439 for ; Wed, 19 Jan 2022 22:29:58 +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 371DD3858D35 for ; Wed, 19 Jan 2022 22:29:39 +0000 (GMT) DMARC-Filter: OpenDMARC Filter v1.4.1 sourceware.org 371DD3858D35 Authentication-Results: sourceware.org; dmarc=none (p=none dis=none) header.from=codesourcery.com Authentication-Results: sourceware.org; spf=pass smtp.mailfrom=mentor.com IronPort-SDR: F94HFaEnbYFEEIpQvRGQDg2GOwvGKLgGH5h8LhBGJR8QTF0v80QzNZswDKedRiW3i2EbnBie73 L4yENsNylA1Uu7i58d4gEtNgJWGRo8y9Zan7ccS5lHXpVN5CRBLEKEPtZ6Gwn1hekQZ5xM5xpC fK9Grg+LisvSnh/KSaD7cCxJ646GZR+/MwYwip/PkOvhEFG8P6ilaS24NKjVV1rvTbn5Yvm75x xg9M/N09cOj6P7xxgEVAaV9sjUb0fpD7Wl6lqXg3VNK08CobsMFyRxMcIA8OxIaaV92pjvfPsz lC8szc/jIJOEpmJacff5nT3R X-IronPort-AV: E=Sophos;i="5.88,300,1635235200"; d="scan'208,223";a="70978264" Received: from orw-gwy-02-in.mentorg.com ([192.94.38.167]) by esa2.mentor.iphmx.com with ESMTP; 19 Jan 2022 14:29:37 -0800 IronPort-SDR: vGnUNQq54oD/aS9kjnWKN0YQzkt4xyALpKFla9UiPGx5VOUOABdo6oYnk/S5kLevN4GrEg3/lr e155VEIKDr/oC40XiGVcbWwm4Vk90xt/gFqqGWVsz+lCxZyjgk3ouHcsAszpgQbaxcHTuOuSeX POgvASeO7Wb4wx7B04C1bmdkRyOkFn/HoBgO+0BP4PSHeA5FoGARz2rhx2DryK9fOGVr3ShXv7 kwPGjidY9ot+jMKeXS7OFRTb2ifrHC4E/g5YdR1oEQBnnjT0uHcJ7aQpYYHOjXLGw7Bx9yVHnW 19Y= From: Thomas Schwinge To: Subject: Catch 'GIMPLE_DEBUG' misbehavior in OpenACC 'kernels' decomposition [PR100400, PR103836, PR104061] (was: Decompose OpenACC 'kernels' constructs into parts, a sequence of compute constructs) In-Reply-To: <877dqodhi1.fsf@euler.schwinge.homeip.net> References: <877dqodhi1.fsf@euler.schwinge.homeip.net> User-Agent: Notmuch/0.29.1+93~g67ed7df (https://notmuchmail.org) Emacs/27.1 (i686-pc-linux-gnu) Date: Wed, 19 Jan 2022 23:29:18 +0100 Message-ID: <871r13we9t.fsf@dirichlet.schwinge.homeip.net> MIME-Version: 1.0 X-Originating-IP: [137.202.0.90] X-ClientProxiedBy: svr-ies-mbx-01.mgc.mentorg.com (139.181.222.1) To svr-ies-mbx-01.mgc.mentorg.com (139.181.222.1) X-Spam-Status: No, score=-12.1 required=5.0 tests=BAYES_00, GIT_PATCH_0, HEADER_FROM_DIFFERENT_DOMAINS, KAM_DMARC_STATUS, SPF_HELO_PASS, SPF_PASS, TXREP 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 2020-11-13T23:22:30+0100, I wrote: > I've pushed to master branch [...] commit > e898ce7997733c29dcab9c3c62ca102c7f9fa6eb "Decompose OpenACC 'kernels' > constructs into parts, a sequence of compute constructs", see attached. > > On 2019-02-01T00:59:30+0100, I wrote: >> 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... :-) (The pass is still disabled by default, by the way.) We've found that 'gcc/omp-oacc-kernels-decompose.cc' is currently not at all considerate of 'GIMPLE_DEBUG' statements -- and it's not always straight forward how to handle these (not rocket science either; but needs proper understanding and testing). Actually fixing it is a separate task, but it seems prudent to at least catch it, and document via a few test cases. OK to push "Catch 'GIMPLE_DEBUG' misbehavior in OpenACC 'kernels' decomposition [PR100400, PR103836, PR104061]", 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 568808ef7ccc97ebeae90bc7cb1aba6bd7659b24 Mon Sep 17 00:00:00 2001 From: Thomas Schwinge Date: Wed, 19 Jan 2022 14:04:42 +0100 Subject: [PATCH] Catch 'GIMPLE_DEBUG' misbehavior in OpenACC 'kernels' decomposition [PR100400, PR103836, PR104061] Actually fixing it is a separate task, but it seems prudent to at least catch it, and document via a few test cases. gcc/ PR middle-end/100400 PR middle-end/103836 PR middle-end/104061 * omp-oacc-kernels-decompose.cc (decompose_kernels_region_body): Catch 'GIMPLE_DEBUG'. gcc/testsuite/ PR middle-end/100400 PR middle-end/103836 PR middle-end/104061 * c-c++-common/goacc/kernels-decompose-pr100400-1-1.c: New. * c-c++-common/goacc/kernels-decompose-pr100400-1-2.c: New. * c-c++-common/goacc/kernels-decompose-pr100400-1-3.c: New. * c-c++-common/goacc/kernels-decompose-pr100400-1-4.c: New. * c-c++-common/goacc/kernels-decompose-pr103836-1-1.c: New. * c-c++-common/goacc/kernels-decompose-pr103836-1-2.c: New. * c-c++-common/goacc/kernels-decompose-pr103836-1-3.c: New. * c-c++-common/goacc/kernels-decompose-pr103836-1-4.c: New. * c-c++-common/goacc/kernels-decompose-pr104061-1-1.c: New. * c-c++-common/goacc/kernels-decompose-pr104061-1-2.c: New. * c-c++-common/goacc/kernels-decompose-pr104061-1-3.c: New. * c-c++-common/goacc/kernels-decompose-pr104061-1-4.c: New. --- gcc/omp-oacc-kernels-decompose.cc | 10 +++++ .../goacc/kernels-decompose-pr100400-1-1.c | 33 ++++++++++++++ .../goacc/kernels-decompose-pr100400-1-2.c | 40 +++++++++++++++++ .../goacc/kernels-decompose-pr100400-1-3.c | 42 ++++++++++++++++++ .../goacc/kernels-decompose-pr100400-1-4.c | 40 +++++++++++++++++ .../goacc/kernels-decompose-pr103836-1-1.c | 26 +++++++++++ .../goacc/kernels-decompose-pr103836-1-2.c | 29 +++++++++++++ .../goacc/kernels-decompose-pr103836-1-3.c | 30 +++++++++++++ .../goacc/kernels-decompose-pr103836-1-4.c | 30 +++++++++++++ .../goacc/kernels-decompose-pr104061-1-1.c | 30 +++++++++++++ .../goacc/kernels-decompose-pr104061-1-2.c | 33 ++++++++++++++ .../goacc/kernels-decompose-pr104061-1-3.c | 43 +++++++++++++++++++ .../goacc/kernels-decompose-pr104061-1-4.c | 41 ++++++++++++++++++ 13 files changed, 427 insertions(+) create mode 100644 gcc/testsuite/c-c++-common/goacc/kernels-decompose-pr100400-1-1.c create mode 100644 gcc/testsuite/c-c++-common/goacc/kernels-decompose-pr100400-1-2.c create mode 100644 gcc/testsuite/c-c++-common/goacc/kernels-decompose-pr100400-1-3.c create mode 100644 gcc/testsuite/c-c++-common/goacc/kernels-decompose-pr100400-1-4.c create mode 100644 gcc/testsuite/c-c++-common/goacc/kernels-decompose-pr103836-1-1.c create mode 100644 gcc/testsuite/c-c++-common/goacc/kernels-decompose-pr103836-1-2.c create mode 100644 gcc/testsuite/c-c++-common/goacc/kernels-decompose-pr103836-1-3.c create mode 100644 gcc/testsuite/c-c++-common/goacc/kernels-decompose-pr103836-1-4.c create mode 100644 gcc/testsuite/c-c++-common/goacc/kernels-decompose-pr104061-1-1.c create mode 100644 gcc/testsuite/c-c++-common/goacc/kernels-decompose-pr104061-1-2.c create mode 100644 gcc/testsuite/c-c++-common/goacc/kernels-decompose-pr104061-1-3.c create mode 100644 gcc/testsuite/c-c++-common/goacc/kernels-decompose-pr104061-1-4.c diff --git a/gcc/omp-oacc-kernels-decompose.cc b/gcc/omp-oacc-kernels-decompose.cc index 21872db3ed3..98eafdbe3a1 100644 --- a/gcc/omp-oacc-kernels-decompose.cc +++ b/gcc/omp-oacc-kernels-decompose.cc @@ -1255,6 +1255,16 @@ decompose_kernels_region_body (gimple *kernels_region, tree kernels_clauses) gsi_next (&gsi_n); gimple *stmt = gsi_stmt (gsi); + if (gimple_code (stmt) == GIMPLE_DEBUG) + { + if (flag_compare_debug_opt || flag_compare_debug) + /* Let the usual '-fcompare-debug' analysis bail out, as + necessary. */ + ; + else + sorry_at (loc, "%qs not yet supported", + gimple_code_name[gimple_code (stmt)]); + } gimple *omp_for = top_level_omp_for_in_stmt (stmt); bool is_unconditional_oacc_for_loop = false; if (omp_for != NULL) diff --git a/gcc/testsuite/c-c++-common/goacc/kernels-decompose-pr100400-1-1.c b/gcc/testsuite/c-c++-common/goacc/kernels-decompose-pr100400-1-1.c new file mode 100644 index 00000000000..f63800514c4 --- /dev/null +++ b/gcc/testsuite/c-c++-common/goacc/kernels-decompose-pr100400-1-1.c @@ -0,0 +1,33 @@ +/* { dg-additional-options "--param openacc-kernels=decompose" } */ + +/* { dg-additional-options "-g0" } */ +/* { dg-additional-options "-O1" } */ + +/* { dg-additional-options "-fopt-info-all-omp" } */ + +/* { dg-additional-options "--param=openacc-privatization=noisy" } */ + +int *p; + +void +foo (void) +{ +#pragma acc kernels + /* { dg-note {variable 'c' declared in block is candidate for adjusting OpenACC privatization level} {} { target *-*-* } .-1 } */ + /* { dg-note {variable 'c\.0' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } .-2 } */ + { + int c; + + /* { dg-note {beginning 'gang-single' part in OpenACC 'kernels' region} {} { target *-*-* } .+1 } */ + p = &c; + + /* { dg-note {parallelized loop nest in OpenACC 'kernels' region} {} { target *-*-* } .+1 } */ +#pragma acc loop independent + /* { dg-note {variable 'c\.0' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } .-1 } */ + /* { dg-note {variable 'c' in 'private' clause is candidate for adjusting OpenACC privatization level} {} { target *-*-* } .-2 } + { dg-note {variable 'c' ought to be adjusted for OpenACC privatization level: 'vector'} {} { target *-*-* } .-3 } */ + /* { dg-optimized {assigned OpenACC gang vector loop parallelism} {} { target *-*-* } .-4 } */ + for (c = 0; c < 1; ++c) + ; + } +} diff --git a/gcc/testsuite/c-c++-common/goacc/kernels-decompose-pr100400-1-2.c b/gcc/testsuite/c-c++-common/goacc/kernels-decompose-pr100400-1-2.c new file mode 100644 index 00000000000..1eee3b07a75 --- /dev/null +++ b/gcc/testsuite/c-c++-common/goacc/kernels-decompose-pr100400-1-2.c @@ -0,0 +1,40 @@ +/* { dg-additional-options "--param openacc-kernels=decompose" } */ + +/* { dg-additional-options "-fchecking" } + { dg-ice TODO { c++ } } + { dg-prune-output "during GIMPLE pass: omp_oacc_kernels_decompose" } */ + +/* { dg-additional-options "-g" } */ +/* { dg-additional-options "-O1" } so that we may get some 'GIMPLE_DEBUG's. */ + +/* { dg-additional-options "-fopt-info-all-omp" } */ + +/* { dg-additional-options "--param=openacc-privatization=noisy" } */ + +int *p; + +void +foo (void) +{ + /* { dg-bogus {sorry, unimplemented: 'gimple_debug' not yet supported} TODO { xfail *-*-* } .+1 } */ +#pragma acc kernels + /* { dg-note {variable 'c' declared in block is candidate for adjusting OpenACC privatization level} {} { xfail *-*-* } .-1 } */ + /* { dg-note {variable 'c\.0' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} {} { xfail *-*-* } .-2 } */ + { + /* { dg-bogus {note: beginning 'gang-single' part in OpenACC 'kernels' region} {w/ debug} { xfail c++ } .-1 } + { dg-bogus {note: beginning 'gang-single' part in OpenACC 'kernels' region} {w/ debug} { xfail c } .+1 } */ + int c; + + /* { dg-note {beginning 'gang-single' part in OpenACC 'kernels' region} {} { xfail *-*-* } .+1 } */ + p = &c; + + /* { dg-note {parallelized loop nest in OpenACC 'kernels' region} {} { xfail c++ } .+1 } */ +#pragma acc loop independent + /* { dg-note {variable 'c\.0' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} {} { xfail *-*-* } .-1 } */ + /* { dg-note {variable 'c' in 'private' clause is candidate for adjusting OpenACC privatization level} {} { xfail *-*-* } .-2 } + { dg-note {variable 'c' ought to be adjusted for OpenACC privatization level: 'vector'} {} { xfail *-*-* } .-3 } */ + /* { dg-optimized {assigned OpenACC gang vector loop parallelism} {} { xfail *-*-* } .-4 } */ + for (c = 0; c < 1; ++c) + ; + } +} diff --git a/gcc/testsuite/c-c++-common/goacc/kernels-decompose-pr100400-1-3.c b/gcc/testsuite/c-c++-common/goacc/kernels-decompose-pr100400-1-3.c new file mode 100644 index 00000000000..dce4e399fbe --- /dev/null +++ b/gcc/testsuite/c-c++-common/goacc/kernels-decompose-pr100400-1-3.c @@ -0,0 +1,42 @@ +/* { dg-additional-options "--param openacc-kernels=decompose" } */ + +/* { dg-additional-options "-fchecking" } + { dg-ice TODO { c++ } } + { dg-prune-output "during GIMPLE pass: omp_oacc_kernels_decompose" } */ + +/* { dg-additional-options "-fcompare-debug" } -- w/o debug compiled first. + { dg-bogus {error: during '-fcompare-debug' recompilation} TODO { xfail c++ } 0 } + { dg-bogus {error: [^\n\r]+: '-fcompare-debug' failure \(length\)} TODO { xfail c++ } 0 } */ +/* { dg-additional-options "-O1" } so that we may get some 'GIMPLE_DEBUG's. */ + +/* { dg-additional-options "-fopt-info-all-omp" } */ + +/* { dg-additional-options "--param=openacc-privatization=noisy" } */ + +int *p; + +void +foo (void) +{ + /* { dg-bogus {sorry, unimplemented: 'gimple_debug' not yet supported} {} { target *-*-* } .+1 } suppressed via '-fcompare-debug'. */ +#pragma acc kernels + /* { dg-note {variable 'c' declared in block is candidate for adjusting OpenACC privatization level} {} { target *-*-* } .-1 } */ + /* { dg-note {variable 'c\.0' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } .-2 } */ + { + /* { dg-bogus {note: beginning 'gang-single' part in OpenACC 'kernels' region} {w/ debug} { xfail c++ } .-1 } + { dg-bogus {note: beginning 'gang-single' part in OpenACC 'kernels' region} {w/ debug} { xfail c } .+1 } */ + int c; + + /* { dg-note {beginning 'gang-single' part in OpenACC 'kernels' region} {} { target *-*-* } .+1 } */ + p = &c; + + /* { dg-note {parallelized loop nest in OpenACC 'kernels' region} {} { target *-*-* } .+1 } */ +#pragma acc loop independent + /* { dg-note {variable 'c\.0' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } .-1 } */ + /* { dg-note {variable 'c' in 'private' clause is candidate for adjusting OpenACC privatization level} {} { target *-*-* } .-2 } + { dg-note {variable 'c' ought to be adjusted for OpenACC privatization level: 'vector'} {} { target *-*-* } .-3 } */ + /* { dg-optimized {assigned OpenACC gang vector loop parallelism} {} { target *-*-* } .-4 } */ + for (c = 0; c < 1; ++c) + ; + } +} diff --git a/gcc/testsuite/c-c++-common/goacc/kernels-decompose-pr100400-1-4.c b/gcc/testsuite/c-c++-common/goacc/kernels-decompose-pr100400-1-4.c new file mode 100644 index 00000000000..7ca4440d075 --- /dev/null +++ b/gcc/testsuite/c-c++-common/goacc/kernels-decompose-pr100400-1-4.c @@ -0,0 +1,40 @@ +/* { dg-additional-options "--param openacc-kernels=decompose" } */ + +/* { dg-additional-options "-fchecking" } + { dg-ice TODO { c++ } } + { dg-prune-output "during GIMPLE pass: omp_oacc_kernels_decompose" } */ + +/* { dg-additional-options "-g -fcompare-debug" } -- w/ debug compiled first. */ +/* { dg-additional-options "-O1" } so that we may get some 'GIMPLE_DEBUG's. */ + +/* { dg-additional-options "-fopt-info-all-omp" } */ + +/* { dg-additional-options "--param=openacc-privatization=noisy" } */ + +int *p; + +void +foo (void) +{ + /* { dg-bogus {sorry, unimplemented: 'gimple_debug' not yet supported} {} { target *-*-* } .+1 } suppressed via '-fcompare-debug'. */ +#pragma acc kernels + /* { dg-note {variable 'c' declared in block is candidate for adjusting OpenACC privatization level} {} { xfail c++ } .-1 } */ + /* { dg-note {variable 'c\.0' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} {} { xfail c++ } .-2 } */ + { + /* { dg-bogus {note: beginning 'gang-single' part in OpenACC 'kernels' region} {w/ debug} { xfail c++ } .-1 } + { dg-bogus {note: beginning 'gang-single' part in OpenACC 'kernels' region} {w/ debug} { xfail c } .+1 } */ + int c; + + /* { dg-note {beginning 'gang-single' part in OpenACC 'kernels' region} {} { xfail c++ } .+1 } */ + p = &c; + + /* { dg-note {parallelized loop nest in OpenACC 'kernels' region} {} { xfail c++ } .+1 } */ +#pragma acc loop independent + /* { dg-note {variable 'c\.0' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} {} { xfail c++ } .-1 } */ + /* { dg-note {variable 'c' in 'private' clause is candidate for adjusting OpenACC privatization level} {} { xfail c++ } .-2 } + { dg-note {variable 'c' ought to be adjusted for OpenACC privatization level: 'vector'} {} { xfail c++ } .-3 } */ + /* { dg-optimized {assigned OpenACC gang vector loop parallelism} {} { xfail c++ } .-4 } */ + for (c = 0; c < 1; ++c) + ; + } +} diff --git a/gcc/testsuite/c-c++-common/goacc/kernels-decompose-pr103836-1-1.c b/gcc/testsuite/c-c++-common/goacc/kernels-decompose-pr103836-1-1.c new file mode 100644 index 00000000000..46ca0c99d2f --- /dev/null +++ b/gcc/testsuite/c-c++-common/goacc/kernels-decompose-pr103836-1-1.c @@ -0,0 +1,26 @@ +/* { dg-additional-options "--param openacc-kernels=decompose" } */ + +/* { dg-additional-options "-g0" } */ +/* { dg-additional-options "-O1" } */ + +/* { dg-additional-options "-fopt-info-all-omp" } */ + +/* { dg-additional-options "--param=openacc-privatization=noisy" } */ + +extern int i; + +void +f_acc_kernels (void) +{ +#pragma acc kernels + /* { dg-note {variable 'i\.0' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } .-1 } */ + { + /* { dg-note {forwarded loop nest in OpenACC 'kernels' region to 'parloops' for analysis} {} { target *-*-* } .+1 } */ +#pragma acc loop + /* { dg-note {variable 'i\.0' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } .-1 } */ + /* { dg-note {variable 'i' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } .-2 } */ + /* { dg-optimized {assigned OpenACC seq loop parallelism} {} { target *-*-* } .-3 } */ + for (i = 0; i < 2; ++i) + ; + } +} diff --git a/gcc/testsuite/c-c++-common/goacc/kernels-decompose-pr103836-1-2.c b/gcc/testsuite/c-c++-common/goacc/kernels-decompose-pr103836-1-2.c new file mode 100644 index 00000000000..e0f24cee2db --- /dev/null +++ b/gcc/testsuite/c-c++-common/goacc/kernels-decompose-pr103836-1-2.c @@ -0,0 +1,29 @@ +/* { dg-additional-options "--param openacc-kernels=decompose" } */ + +/* { dg-additional-options "-g" } */ +/* { dg-additional-options "-O1" } so that we may get some 'GIMPLE_DEBUG's. */ + +/* { dg-additional-options "-fopt-info-all-omp" } */ + +/* { dg-additional-options "--param=openacc-privatization=noisy" } */ + +extern int i; + +void +f_acc_kernels (void) +{ + /* { dg-bogus {sorry, unimplemented: 'gimple_debug' not yet supported} TODO { xfail c++ } .+1 } */ +#pragma acc kernels + /* { dg-note {variable 'i\.0' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} {} { xfail c++ } .-1 } */ + { + /* { dg-bogus {note: beginning 'gang-single' part in OpenACC 'kernels' region} {w/ debug} { xfail c++ } .-1 } */ + + /* { dg-note {forwarded loop nest in OpenACC 'kernels' region to 'parloops' for analysis} {} { target *-*-* } .+1 } */ +#pragma acc loop + /* { dg-note {variable 'i\.0' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} {} { xfail c++ } .-1 } */ + /* { dg-note {variable 'i' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} {} { xfail c++ } .-2 } */ + /* { dg-optimized {assigned OpenACC seq loop parallelism} {} { xfail c++ } .-3 } */ + for (i = 0; i < 2; ++i) + ; + } +} diff --git a/gcc/testsuite/c-c++-common/goacc/kernels-decompose-pr103836-1-3.c b/gcc/testsuite/c-c++-common/goacc/kernels-decompose-pr103836-1-3.c new file mode 100644 index 00000000000..cbf1b7c3e25 --- /dev/null +++ b/gcc/testsuite/c-c++-common/goacc/kernels-decompose-pr103836-1-3.c @@ -0,0 +1,30 @@ +/* { dg-additional-options "--param openacc-kernels=decompose" } */ + +/* { dg-additional-options "-fcompare-debug" } -- w/o debug compiled first. + { dg-bogus {error: [^\n\r]+: '-fcompare-debug' failure \(length\)} TODO { xfail c++ } 0 } */ +/* { dg-additional-options "-O1" } so that we may get some 'GIMPLE_DEBUG's. */ + +/* { dg-additional-options "-fopt-info-all-omp" } */ + +/* { dg-additional-options "--param=openacc-privatization=noisy" } */ + +extern int i; + +void +f_acc_kernels (void) +{ + /* { dg-bogus {sorry, unimplemented: 'gimple_debug' not yet supported} {} { target *-*-* } .+1 } */ +#pragma acc kernels + /* { dg-note {variable 'i\.0' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } .-1 } */ + { + /* { dg-bogus {note: beginning 'gang-single' part in OpenACC 'kernels' region} {w/ debug} { xfail c++ } .-1 } */ + + /* { dg-note {forwarded loop nest in OpenACC 'kernels' region to 'parloops' for analysis} {} { target *-*-* } .+1 } */ +#pragma acc loop + /* { dg-note {variable 'i\.0' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } .-1 } */ + /* { dg-note {variable 'i' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } .-2 } */ + /* { dg-optimized {assigned OpenACC seq loop parallelism} {} { target *-*-* } .-3 } */ + for (i = 0; i < 2; ++i) + ; + } +} diff --git a/gcc/testsuite/c-c++-common/goacc/kernels-decompose-pr103836-1-4.c b/gcc/testsuite/c-c++-common/goacc/kernels-decompose-pr103836-1-4.c new file mode 100644 index 00000000000..21bbe37723f --- /dev/null +++ b/gcc/testsuite/c-c++-common/goacc/kernels-decompose-pr103836-1-4.c @@ -0,0 +1,30 @@ +/* { dg-additional-options "--param openacc-kernels=decompose" } */ + +/* { dg-additional-options "-g -fcompare-debug" } -- w/ debug compiled first. + { dg-bogus {error: [^\n\r]+: '-fcompare-debug' failure \(length\)} TODO { xfail c++ } 0 } */ +/* { dg-additional-options "-O1" } so that we may get some 'GIMPLE_DEBUG's. */ + +/* { dg-additional-options "-fopt-info-all-omp" } */ + +/* { dg-additional-options "--param=openacc-privatization=noisy" } */ + +extern int i; + +void +f_acc_kernels (void) +{ + /* { dg-bogus {sorry, unimplemented: 'gimple_debug' not yet supported} {} { target *-*-* } .+1 } */ +#pragma acc kernels + /* { dg-note {variable 'i\.0' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } .-1 } */ + { + /* { dg-bogus {note: beginning 'gang-single' part in OpenACC 'kernels' region} {w/ debug} { xfail c++ } .-1 } */ + + /* { dg-note {forwarded loop nest in OpenACC 'kernels' region to 'parloops' for analysis} {} { target *-*-* } .+1 } */ +#pragma acc loop + /* { dg-note {variable 'i\.0' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } .-1 } */ + /* { dg-note {variable 'i' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } .-2 } */ + /* { dg-optimized {assigned OpenACC seq loop parallelism} {} { target *-*-* } .-3 } */ + for (i = 0; i < 2; ++i) + ; + } +} diff --git a/gcc/testsuite/c-c++-common/goacc/kernels-decompose-pr104061-1-1.c b/gcc/testsuite/c-c++-common/goacc/kernels-decompose-pr104061-1-1.c new file mode 100644 index 00000000000..a58fce33426 --- /dev/null +++ b/gcc/testsuite/c-c++-common/goacc/kernels-decompose-pr104061-1-1.c @@ -0,0 +1,30 @@ +/* { dg-additional-options "--param openacc-kernels=decompose" } */ + +/* { dg-additional-options "-g0" } */ +/* { dg-additional-options "-O1" } */ + +/* { dg-additional-options "-fopt-info-all-omp" } */ + +/* { dg-additional-options "--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} } */ + +int arr_0; + +void +foo (void) +{ +#pragma acc kernels + /* { dg-note {variable 'arr_0\.0' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } .-1 } */ + { + int k; + + /* { dg-note {forwarded loop nest in OpenACC 'kernels' region to 'parloops' for analysis} {} { target *-*-* } .+1 } */ +#pragma acc loop + /* { dg-note {variable 'k' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } .-1 } */ + /* { dg-note {variable 'k' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } .-2 } */ + /* { dg-optimized {assigned OpenACC seq loop parallelism} {} { target *-*-* } .-3 } */ + for (k = 0; k < 2; k++) + arr_0 += k; + } +} 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 new file mode 100644 index 00000000000..d66dee6f8a7 --- /dev/null +++ b/gcc/testsuite/c-c++-common/goacc/kernels-decompose-pr104061-1-2.c @@ -0,0 +1,33 @@ +/* { dg-additional-options "--param openacc-kernels=decompose" } */ + +/* { dg-additional-options "-g" } */ +/* { dg-additional-options "-O1" } so that we may get some 'GIMPLE_DEBUG's. */ + +/* { dg-additional-options "-fopt-info-all-omp" } */ + +/* { dg-additional-options "--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} } */ + +int arr_0; + +void +foo (void) +{ + /* { dg-bogus {sorry, unimplemented: 'gimple_debug' not yet supported} TODO { xfail *-*-* } .+1 } */ +#pragma acc kernels + /* { dg-note {variable 'arr_0\.0' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} {} { xfail *-*-* } .-1 } */ + { + /* { dg-bogus {note: beginning 'gang-single' part in OpenACC 'kernels' region} {w/ debug} { xfail c++ } .-1 } + { dg-bogus {note: beginning 'gang-single' part in OpenACC 'kernels' region} {w/ debug} { xfail c } .+1 } */ + int k; + + /* { dg-note {forwarded loop nest in OpenACC 'kernels' region to 'parloops' for analysis} {} { target *-*-* } .+1 } */ +#pragma acc loop + /* { dg-note {variable 'k' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} {} { xfail *-*-* } .-1 } */ + /* { dg-note {variable 'k' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} {} { xfail *-*-* } .-2 } */ + /* { dg-optimized {assigned OpenACC seq loop parallelism} {} { xfail *-*-* } .-3 } */ + for (k = 0; k < 2; k++) + arr_0 += k; + } +} 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 new file mode 100644 index 00000000000..20c84e2f3db --- /dev/null +++ b/gcc/testsuite/c-c++-common/goacc/kernels-decompose-pr104061-1-3.c @@ -0,0 +1,43 @@ +/* { dg-additional-options "--param openacc-kernels=decompose" } */ + +/* { dg-additional-options "-fchecking" } + { dg-ice TODO } + { dg-prune-output {D\.[0-9]+ = arr_0\.0 \+ k;} } + { dg-prune-output {during GIMPLE pass: lower} } */ + +/* { dg-additional-options "-fcompare-debug" } -- w/o debug compiled first. + { dg-bogus {error: during '-fcompare-debug' recompilation} TODO { xfail *-*-* } 0 } + { dg-bogus {error: [^\n\r]+: '-fcompare-debug' failure \(length\)} TODO { xfail *-*-* } 0 } */ +/* { dg-additional-options "-O1" } so that we may get some 'GIMPLE_DEBUG's. */ + +/* { dg-additional-options "-fopt-info-all-omp" } */ + +/* { dg-additional-options "--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} } */ + +int arr_0; + +void +foo (void) +{ + /* { dg-bogus {sorry, unimplemented: 'gimple_debug' not yet supported} {} { target *-*-* } .+1 } suppressed via '-fcompare-debug'. */ +#pragma acc kernels + /* { dg-bogus {note: variable 'k' declared in block is candidate for adjusting OpenACC privatization level} {w/ debug} { xfail *-*-* } .-1 } */ + /* { dg-note {variable 'arr_0\.0' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } .-2 } */ + { + /* { dg-bogus {note: beginning 'gang-single' part in OpenACC 'kernels' region} {w/ debug} { xfail c++ } .-1 } + { dg-bogus {note: beginning 'gang-single' part in OpenACC 'kernels' region} {w/ debug} { xfail c } .+1 } */ + int k; + + /* { dg-note {forwarded loop nest in OpenACC 'kernels' region to 'parloops' for analysis} {} { target *-*-* } .+1 } */ +#pragma acc loop + /* { dg-note {variable 'k' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } .-1 } */ + /* { dg-note {variable 'k' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } .-2 } */ + /* { dg-bogus {note: variable 'k' in 'private' clause is candidate for adjusting OpenACC privatization level} {w/ debug} { xfail *-*-* } .-3 } */ + /* { dg-optimized {assigned OpenACC seq loop parallelism} {} { target *-*-* } .-4 } */ + for (k = 0; k < 2; k++) + arr_0 += k; + /* { dg-bogus {error: invalid operands in binary operation} {w/ debug} { xfail *-*-* } .-1 } */ + } +} 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 new file mode 100644 index 00000000000..6b6effe1791 --- /dev/null +++ b/gcc/testsuite/c-c++-common/goacc/kernels-decompose-pr104061-1-4.c @@ -0,0 +1,41 @@ +/* { dg-additional-options "--param openacc-kernels=decompose" } */ + +/* { dg-additional-options "-fchecking" } + { dg-ice TODO } + { dg-prune-output {D\.[0-9]+ = arr_0\.0 \+ k;} } + { dg-prune-output {during GIMPLE pass: lower} } */ + +/* { dg-additional-options "-g -fcompare-debug" } -- w/ debug compiled first. */ +/* { dg-additional-options "-O1" } so that we may get some 'GIMPLE_DEBUG's. */ + +/* { dg-additional-options "-fopt-info-all-omp" } */ + +/* { dg-additional-options "--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} } */ + +int arr_0; + +void +foo (void) +{ + /* { dg-bogus {sorry, unimplemented: 'gimple_debug' not yet supported} {} { target *-*-* } .+1 } suppressed via '-fcompare-debug'. */ +#pragma acc kernels + /* { dg-bogus {note: variable 'k' declared in block is candidate for adjusting OpenACC privatization level} {w/ debug} { xfail *-*-* } .-1 } */ + /* { dg-note {variable 'arr_0\.0' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } .-2 } */ + { + /* { dg-bogus {note: beginning 'gang-single' part in OpenACC 'kernels' region} {w/ debug} { xfail c++ } .-1 } + { dg-bogus {note: beginning 'gang-single' part in OpenACC 'kernels' region} {w/ debug} { xfail c } .+1 } */ + int k; + + /* { dg-note {forwarded loop nest in OpenACC 'kernels' region to 'parloops' for analysis} {} { target *-*-* } .+1 } */ +#pragma acc loop + /* { dg-note {variable 'k' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} {} { xfail *-*-* } .-1 } */ + /* { dg-note {variable 'k' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} {} { xfail *-*-* } .-2 } */ + /* { dg-bogus {note: variable 'k' in 'private' clause is candidate for adjusting OpenACC privatization level} {w/ debug} { xfail *-*-* } .-3 } */ + /* { dg-optimized {assigned OpenACC seq loop parallelism} {} { xfail *-*-* } .-4 } */ + for (k = 0; k < 2; k++) + arr_0 += k; + /* { dg-bogus {error: invalid operands in binary operation} {w/ debug} { xfail *-*-* } .-1 } */ + } +} -- 2.25.1