From mboxrd@z Thu Jan 1 00:00:00 1970 Return-Path: Received: (qmail 110831 invoked by alias); 23 Mar 2016 06:48:20 -0000 Mailing-List: contact gcc-patches-help@gcc.gnu.org; run by ezmlm Precedence: bulk List-Id: List-Archive: List-Post: List-Help: Sender: gcc-patches-owner@gcc.gnu.org Received: (qmail 110809 invoked by uid 89); 23 Mar 2016 06:48:19 -0000 Authentication-Results: sourceware.org; auth=none X-Virus-Found: No X-Spam-SWARE-Status: No, score=-1.9 required=5.0 tests=AWL,BAYES_00,RCVD_IN_DNSWL_NONE,SPF_PASS autolearn=ham version=3.3.2 spammy=Hx-languages-length:1835 X-HELO: relay1.mentorg.com Received: from relay1.mentorg.com (HELO relay1.mentorg.com) (192.94.38.131) by sourceware.org (qpsmtpd/0.93/v0.84-503-g423c35a) with (AES256-GCM-SHA384 encrypted) ESMTPS; Wed, 23 Mar 2016 06:48:16 +0000 Received: from nat-ies.mentorg.com ([192.94.31.2] helo=SVR-IES-FEM-01.mgc.mentorg.com) by relay1.mentorg.com with esmtp id 1aicaJ-0001zY-7Q from Thomas_Schwinge@mentor.com ; Tue, 22 Mar 2016 23:48:03 -0700 Received: from hertz.schwinge.homeip.net (137.202.0.76) by SVR-IES-FEM-01.mgc.mentorg.com (137.202.0.104) with Microsoft SMTP Server id 14.3.224.2; Wed, 23 Mar 2016 06:48:02 +0000 From: Thomas Schwinge To: Bernd Schmidt CC: , Jakub Jelinek , Nathan Sidwell Subject: Re: Also test -O0 for OpenACC C, C++ offloading test cases In-Reply-To: <56F1CC9B.1060003@redhat.com> References: <87mvpqrets.fsf@kepler.schwinge.homeip.net> <56F1CC9B.1060003@redhat.com> User-Agent: Notmuch/0.9-101-g81dad07 (http://notmuchmail.org) Emacs/24.4.1 (x86_64-pc-linux-gnu) Date: Wed, 23 Mar 2016 07:35:00 -0000 Message-ID: <87mvppbsgw.fsf@hertz.schwinge.homeip.net> MIME-Version: 1.0 Content-Type: text/plain; charset="utf-8" Content-Transfer-Encoding: quoted-printable X-SW-Source: 2016-03/txt/msg01281.txt.bz2 Hi! On Tue, 22 Mar 2016 23:52:11 +0100, Bernd Schmidt wro= te: > On 03/22/2016 11:23 AM, Thomas Schwinge wrote: > > --- libgomp/testsuite/libgomp.oacc-c-c++-common/routine-w-1.c > > +++ libgomp/testsuite/libgomp.oacc-c-c++-common/routine-w-1.c > > @@ -1,5 +1,6 @@ > > -/* { dg-do run } */ > > -/* { dg-additional-options "-O2" } */ > > +/* Dead code elimination for blocks guarded by acc_on_device () only w= orks with > > + optimizations enabled. > > + { dg-skip-if "" { *-*-* } { "-O0" } { "" } } */ >=20 > What exactly is going on with these? Do these tests fail with -O0, and=20 > is that likely to be a problem in practice? Want me to re-word that? :-| I thought it would be obvious from looking at the test case code; will not be a problem in practice. It's because of constructs used in the test cases, like the following, for example: #pragma acc loop worker for (unsigned ix =3D 0; ix < N; ix++) { if (__builtin_acc_on_device (5)) { int g =3D 0, w =3D 0, v =3D 0; =20=20=20=20 __asm__ volatile ("mov.u32 %0,%%ctaid.x;" : "=3Dr" (g)); __asm__ volatile ("mov.u32 %0,%%tid.y;" : "=3Dr" (w)); __asm__ volatile ("mov.u32 %0,%%tid.x;" : "=3Dr" (v)); ary[ix] =3D (g << 16) | (w << 8) | v; } else ary[ix] =3D ix; } Without optimizations, the target (x86_64) assembler will bail out seeing the device (nvptx) inline assembly code, even if it's dead code always because of the acc_on_device () conditional. Long ago, my suggestion has been to have GCC provide builtin functions for users to retrieve the number of gangs, workers, vectors, and the current thread's IDs of these; not sure why Nathan didn't implement that? (Should be easy to do -- want me to have a look at that, as a separate patch?) > Also, why remove the dg-do run? Because that's the default anyway. Gr=C3=BC=C3=9Fe Thomas