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