This is the mail archive of the
gcc-patches@gcc.gnu.org
mailing list for the GCC project.
Re: [2/2] OpenACC routine support
- From: Jakub Jelinek <jakub at redhat dot com>
- To: Nathan Sidwell <nathan at acm dot org>
- Cc: GCC Patches <gcc-patches at gcc dot gnu dot org>
- Date: Mon, 2 Nov 2015 20:41:07 +0100
- Subject: Re: [2/2] OpenACC routine support
- Authentication-results: sourceware.org; auth=none
- References: <5637B1CF dot 5060408 at acm dot org> <5637B827 dot 1020909 at acm dot org>
- Reply-to: Jakub Jelinek <jakub at redhat dot com>
On Mon, Nov 02, 2015 at 02:23:19PM -0500, Nathan Sidwell wrote:
> +#pragma acc routine gang
> +void __attribute__ ((noinline)) gang (int ary[N])
> +{
> +#pragma acc loop gang
> + 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;
Does this work even with -O0? I mean, the assembler is invalid
for any target other than PTX, so you are relying on aggressively folding
this away.
Jakub