libgomp/testsuite/libgomp.oacc-c-c++-common/loop-red-g-1.c

   1 /* This code uses nvptx inline assembly guarded with acc_on_device, which is
   2    not optimized away at -O0, and then confuses the target assembler.
   3    { dg-skip-if "" { *-*-* } { "-O0" } { "" } } */
   4
   5 #include <stdio.h>
   6
   7 #define N (32*32*32+17)
   8 int main ()
   9 {
  10   int ix;
  11   int ondev = 0;
  12   int t = 0, h = 0;
  13
  14 #pragma acc parallel num_gangs(32) copy(ondev)
  15   {
  16 #pragma acc loop gang  reduction (+:t)
  17     for (unsigned ix = 0; ix < N; ix++)
  18       {
  19         int val = ix;
  20
  21         if (__builtin_acc_on_device (5))
  22           {
  23             int g = 0, w = 0, v = 0;
  24
  25             __asm__ volatile ("mov.u32 %0,%%ctaid.x;" : "=r" (g));
  26             __asm__ volatile ("mov.u32 %0,%%tid.y;" : "=r" (w));
  27             __asm__ volatile ("mov.u32 %0,%%tid.x;" : "=r" (v));
  28             val = (g << 16) | (w << 8) | v;
  29             ondev = 1;
  30           }
  31         t += val;
  32       }
  33   }
  34
  35   for (ix = 0; ix < N; ix++)
  36     {
  37       int val = ix;
  38       if(ondev)
  39         {
  40           int g = ix / ((N + 31) / 32);
  41           int w = 0;
  42           int v = 0;
  43
  44           val = (g << 16) | (w << 8) | v;
  45         }
  46       h += val;
  47     }
  48   if (t != h)
  49     {
  50       printf ("t=%x expected %x\n", t, h);
  51       return 1;
  52     }
  53
  54   return 0;
  55 }