libgomp/testsuite/libgomp.oacc-c-c++-common/routine-wv-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 worker
   8 void __attribute__ ((noinline)) worker (int ary[N])
   9 {
  10 #pragma acc loop worker 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 workersize, vectorsize;
  34
  35   for (ix = 0; ix < N;ix++)
  36     ary[ix] = -1;
  37
  38 #define NW 32
  39 #define VL 32
  40 #pragma acc parallel num_workers(NW) vector_length(VL) \
  41             copy(ary) copy(ondev)
  42   {
  43     ondev = acc_on_device (acc_device_not_host);
  44     worker (ary);
  45   }
  46   workersize = NW;
  47   vectorsize = VL;
  48 #ifdef ACC_DEVICE_TYPE_radeon
  49   /* AMD GCN has an upper limit of 'num_workers(16)'.  */
  50   if (workersize > 16)
  51     workersize = 16;
  52   /* AMD GCN uses the autovectorizer for the vector dimension: the use
  53      of a function call in vector-partitioned code in this test is not
  54      currently supported.  */
  55   vectorsize = 1;
  56 #endif
  57
  58   for (ix = 0; ix < N; ix++)
  59     {
  60       int expected = ix;
  61       if(ondev)
  62         {
  63           int g = 0;
  64           int w = (ix / vectorsize) % workersize;
  65           int v = ix % vectorsize;
  66
  67           expected = (g << 16) | (w << 8) | v;
  68         }
  69
  70       if (ary[ix] != expected)
  71         {
  72           exit = 1;
  73           printf ("ary[%d]=%x expected %x\n", ix, ary[ix], expected);
  74         }
  75     }
  76
  77   return exit;
  78 }