From patchwork Sat Mar 12 12:38:38 2022 Content-Type: text/plain; charset="utf-8" MIME-Version: 1.0 Content-Transfer-Encoding: 8bit X-Patchwork-Submitter: Thomas Schwinge X-Patchwork-Id: 51910 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 D8CB33853812 for ; Sat, 12 Mar 2022 12:39:02 +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 42C083858C78 for ; Sat, 12 Mar 2022 12:38:46 +0000 (GMT) DMARC-Filter: OpenDMARC Filter v1.4.1 sourceware.org 42C083858C78 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="72900647" Received: from orw-gwy-01-in.mentorg.com ([192.94.38.165]) by esa3.mentor.iphmx.com with ESMTP; 12 Mar 2022 04:38:46 -0800 IronPort-SDR: iPP2RaTG5HW6aB/3ONgR+ZdPK2PjrOOx0ZFub3d1olOjQBjm4LMbn15FGM6+DC4POwRmyLOk+v 4SaYvm6rpf2nMDDHO7kvzmKk5GAHfXT/qewxJKM836oqRTNhFTK7hciY+BnayYR6YbebfBKJbP scsXNLQGCar98VuwW6HIUHMmkb93QuBFDKE69boUaUGnSo8/ad+4J8Wln5PHw8UAnzuQxqmd4K RnYXUQEBvKU+wtJGMS+yJ6iksh7weFeCwS83I/w6CwASYFC8hzFf0VnXgpQ/QM95/IzTb/vc2d G08= From: Thomas Schwinge To: Subject: Add 'c-c++-common/goacc/kernels-decompose-pr104086-1.c' [PR104086] In-Reply-To: <877dqodhi1.fsf@euler.schwinge.homeip.net> References: <877dqodhi1.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:38:38 +0100 Message-ID: <87ilsjmjld.fsf@euler.schwinge.homeip.net> MIME-Version: 1.0 X-Originating-IP: [137.202.0.90] X-ClientProxiedBy: svr-ies-mbx-12.mgc.mentorg.com (139.181.222.12) 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 , 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: > 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'. > + { 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]", see attached. Grüße Thomas > --- /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; ----------------- 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 9781ae3a254a8c17ef4ffa70f21ed1728ff3c707 Mon Sep 17 00:00:00 2001 From: Thomas Schwinge Date: Tue, 18 Jan 2022 17:22:14 +0100 Subject: [PATCH] Add 'c-c++-common/goacc/kernels-decompose-pr104086-1.c' [PR104086] ..., currently XFAILed with 'dg-ice', as it runs into 'gcc/omp-low.cc:lower_omp_target': 13125 else if (is_gimple_reg (var)) 13126 { 13127 gcc_assert (offloaded); This means, the recent PR100280 etc. changes are still not sufficient. gcc/testsuite/ PR middle-end/104086 * c-c++-common/goacc/kernels-decompose-pr104086-1.c: New file. --- .../goacc/kernels-decompose-pr104086-1.c | 25 +++++++++++++++++++ 1 file changed, 25 insertions(+) create mode 100644 gcc/testsuite/c-c++-common/goacc/kernels-decompose-pr104086-1.c 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 new file mode 100644 index 00000000000..eab10cf6c72 --- /dev/null +++ b/gcc/testsuite/c-c++-common/goacc/kernels-decompose-pr104086-1.c @@ -0,0 +1,25 @@ +/* Reduced from 'libgomp.oacc-c-c++-common/declare-vla.c'. */ + +/* { dg-additional-options "-fchecking" } + { dg-ice TODO } + { dg-prune-output {during GIMPLE pass: omplower} } */ + +/* { dg-additional-options "--param openacc-kernels=decompose" } */ + +/* { dg-additional-options "-fopt-info-all-omp" } */ + +/* { dg-additional-options "--param=openacc-privatization=noisy" } */ + +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 } */ + { + int i; + +#pragma acc kernels + /* { dg-note {beginning 'gang-single' part in OpenACC 'kernels' region} {} { target *-*-* } .+1 } */ + i = 0; + } +} -- 2.34.1