This is the mail archive of the
gcc-patches@gcc.gnu.org
mailing list for the GCC project.
Re: Also test -O0 for OpenACC C, C++ offloading test cases
- From: Thomas Schwinge <thomas at codesourcery dot com>
- To: Bernd Schmidt <bschmidt at redhat dot com>
- Cc: <gcc-patches at gcc dot gnu dot org>, Jakub Jelinek <jakub at redhat dot com>, Nathan Sidwell <nathan at codesourcery dot com>
- Date: Wed, 23 Mar 2016 07:47:59 +0100
- Subject: Re: Also test -O0 for OpenACC C, C++ offloading test cases
- Authentication-results: sourceware.org; auth=none
- References: <87mvpqrets dot fsf at kepler dot schwinge dot homeip dot net> <56F1CC9B dot 1060003 at redhat dot com>
Hi!
On Tue, 22 Mar 2016 23:52:11 +0100, Bernd Schmidt <bschmidt@redhat.com> wrote:
> 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 works with
> > + optimizations enabled.
> > + { dg-skip-if "" { *-*-* } { "-O0" } { "" } } */
>
> What exactly is going on with these? Do these tests fail with -O0, and
> 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 = 0; ix < N; ix++)
{
if (__builtin_acc_on_device (5))
{
int g = 0, w = 0, v = 0;
__asm__ volatile ("mov.u32 %0,%%ctaid.x;" : "=r" (g));
__asm__ volatile ("mov.u32 %0,%%tid.y;" : "=r" (w));
__asm__ volatile ("mov.u32 %0,%%tid.x;" : "=r" (v));
ary[ix] = (g << 16) | (w << 8) | v;
}
else
ary[ix] = 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ÃÃe
Thomas