libgomp/testsuite/libgomp.oacc-c-c++-common/routine-v-1.c

   1 #include <stdio.h>
   2 #include <openacc.h>
   3 #include <gomp-constants.h>
   4
   5 #define N (32*32*32+17)
   6
   7 #pragma acc routine vector
   8 void __attribute__ ((noinline)) vector (int ary[N])
   9 {
  10 #pragma acc loop vector
  11   for (unsigned ix = 0; ix < N; ix++)
  12     {
  13       if (acc_on_device (acc_device_not_host))
  14         {
  15           int g, w, v;
  16
  17           g = __builtin_goacc_parlevel_id (GOMP_DIM_GANG);
  18           w = __builtin_goacc_parlevel_id (GOMP_DIM_WORKER);
  19           v = __builtin_goacc_parlevel_id (GOMP_DIM_VECTOR);
  20           ary[ix] = (g << 16) | (w << 8) | v;
  21         }
  22       else
  23         ary[ix] = ix;
  24     }
  25 }
  26
  27 int main ()
  28 {
  29   int ary[N];
  30   int ix;
  31   int exit = 0;
  32   int ondev = 0;
  33   int vectorsize;
  34
  35   for (ix = 0; ix < N;ix++)
  36     ary[ix] = -1;
  37
  38 #define VL 32
  39 #pragma acc parallel vector_length(VL) \
  40             copy(ary) copy(ondev)
  41   {
  42     ondev = acc_on_device (acc_device_not_host);
  43     vector (ary);
  44   }
  45   vectorsize = VL;
  46 #ifdef ACC_DEVICE_TYPE_radeon
  47   /* AMD GCN uses the autovectorizer for the vector dimension: the use
  48      of a function call in vector-partitioned code in this test is not
  49      currently supported.  */
  50   vectorsize = 1;
  51 #endif
  52
  53   for (ix = 0; ix < N; ix++)
  54     {
  55       int expected = ix;
  56       if(ondev)
  57         {
  58           int g = 0;
  59           int w = 0;
  60           int v = ix % vectorsize;
  61
  62           expected = (g << 16) | (w << 8) | v;
  63         }
  64
  65       if (ary[ix] != expected)
  66         {
  67           exit = 1;
  68           printf ("ary[%d]=%x expected %x\n", ix, ary[ix], expected);
  69         }
  70     }
  71
  72   return exit;
  73 }