libgomp/testsuite/libgomp.oacc-c-c++-common/loop-auto-1.c

   1 /* { dg-additional-options "-fopenacc-dim=32" } */
   2
   3 #include <stdio.h>
   4 #include <openacc.h>
   5 #include <gomp-constants.h>
   6
   7 int check (const int *ary, int size, int gp, int wp, int vp)
   8 {
   9   int exit = 0;
  10   int ix;
  11   int gangs[32], workers[32], vectors[32];
  12
  13   for (ix = 0; ix < 32; ix++)
  14     gangs[ix] = workers[ix] = vectors[ix] = 0;
  15
  16   for (ix = 0; ix < size; ix++)
  17     {
  18       vectors[ary[ix] & 0xff]++;
  19       workers[(ary[ix] >> 8) & 0xff]++;
  20       gangs[(ary[ix] >> 16) & 0xff]++;
  21     }
  22
  23   for (ix = 0; ix < 32; ix++)
  24     {
  25       if (gp)
  26         {
  27           int expect = gangs[0];
  28           if (gangs[ix] != expect)
  29             {
  30               exit = 1;
  31               printf ("gang %d not used %d times\n", ix, expect);
  32             }
  33         }
  34       else if (ix && gangs[ix])
  35         {
  36           exit = 1;
  37           printf ("gang %d unexpectedly used\n", ix);
  38         }
  39
  40       if (wp)
  41         {
  42           int expect = workers[0];
  43           if (workers[ix] != expect)
  44             {
  45               exit = 1;
  46               printf ("worker %d not used %d times\n", ix, expect);
  47             }
  48         }
  49       else if (ix && workers[ix])
  50         {
  51           exit = 1;
  52           printf ("worker %d unexpectedly used\n", ix);
  53         }
  54
  55       if (vp)
  56         {
  57           int expect = vectors[0];
  58           if (vectors[ix] != expect)
  59             {
  60               exit = 1;
  61               printf ("vector %d not used %d times\n", ix, expect);
  62             }
  63         }
  64       else if (ix && vectors[ix])
  65         {
  66           exit = 1;
  67           printf ("vector %d unexpectedly used\n", ix);
  68         }
  69
  70     }
  71   return exit;
  72 }
  73
  74 #pragma acc routine seq
  75 static int __attribute__((noinline)) place ()
  76 {
  77   int r = 0;
  78
  79   int g = 0, w = 0, v = 0;
  80   g = __builtin_goacc_parlevel_id (GOMP_DIM_GANG);
  81   w = __builtin_goacc_parlevel_id (GOMP_DIM_WORKER);
  82   v = __builtin_goacc_parlevel_id (GOMP_DIM_VECTOR);
  83   r = (g << 16) | (w << 8) | v;
  84
  85   return r;
  86 }
  87
  88 static void clear (int *ary, int size)
  89 {
  90   int ix;
  91
  92   for (ix = 0; ix < size; ix++)
  93     ary[ix] = -1;
  94 }
  95
  96 int vector_1 (int *ary, int size)
  97 {
  98   clear (ary, size);
  99
 100 #pragma acc parallel num_workers (32) vector_length(32) copy(ary[0:size]) firstprivate (size)
 101   {
 102 #pragma acc loop gang
 103     for (int jx = 0; jx < 1; jx++)
 104 #pragma acc loop auto
 105       for (int ix = 0; ix < size; ix++)
 106         ary[ix] = place ();
 107   }
 108
 109   return check (ary, size, 0, 1, 1);
 110 }
 111
 112 int vector_2 (int *ary, int size)
 113 {
 114   clear (ary, size);
 115
 116 #pragma acc parallel num_workers (32) vector_length(32) copy(ary[0:size]) firstprivate (size)
 117   {
 118 #pragma acc loop worker
 119     for (int jx = 0; jx < size  / 64; jx++)
 120 #pragma acc loop auto
 121       for (int ix = 0; ix < 64; ix++)
 122         ary[ix + jx * 64] = place ();
 123   }
 124
 125   return check (ary, size, 0, 1, 1);
 126 }
 127
 128 int worker_1 (int *ary, int size)
 129 {
 130   clear (ary, size);
 131
 132 #pragma acc parallel num_workers (32) vector_length(32) copy(ary[0:size]) firstprivate (size)
 133   {
 134 #pragma acc loop gang
 135     for (int kx = 0; kx < 1; kx++)
 136 #pragma acc loop auto
 137       for (int jx = 0; jx <  size  / 64; jx++)
 138 #pragma acc loop vector
 139         for (int ix = 0; ix < 64; ix++)
 140           ary[ix + jx * 64] = place ();
 141   }
 142
 143   return check (ary, size, 0,  1, 1);
 144 }
 145
 146 int gang_1 (int *ary, int size)
 147 {
 148   clear (ary, size);
 149
 150 #pragma acc parallel num_gangs (32) num_workers (32) vector_length(32) copy(ary[0:size]) firstprivate (size)
 151   {
 152 #pragma acc loop auto
 153     for (int jx = 0; jx <  size  / 64; jx++)
 154 #pragma acc loop worker
 155       for (int ix = 0; ix < 64; ix++)
 156         ary[ix + jx * 64] = place ();
 157   }
 158
 159   return check (ary, size, 1, 1, 0);
 160 }
 161
 162 int gang_2 (int *ary, int size)
 163 {
 164   clear (ary, size);
 165
 166 #pragma acc parallel num_gangs (32) num_workers (32) vector_length(32) copy(ary[0:size]) firstprivate (size)
 167   {
 168 #pragma acc loop auto
 169     for (int kx = 0; kx < size / (32 * 32); kx++)
 170 #pragma acc loop auto
 171       for (int jx = 0; jx <  32; jx++)
 172 #pragma acc loop auto
 173         for (int ix = 0; ix < 32; ix++)
 174           ary[ix + jx * 32 + kx * 32 * 32] = place ();
 175   }
 176
 177   return check (ary, size, 1, 1, 1);
 178 }
 179
 180 int gang_3 (int *ary, int size)
 181 {
 182   clear (ary, size);
 183
 184 #pragma acc parallel num_workers (32) vector_length(32) copy(ary[0:size]) firstprivate (size)
 185   {
 186 #pragma acc loop auto
 187     for (int jx = 0; jx <  size  / 64; jx++)
 188 #pragma acc loop auto
 189       for (int ix = 0; ix < 64; ix++)
 190         ary[ix + jx * 64] = place ();
 191   }
 192
 193   return check (ary, size, 1, 1, 1);
 194 }
 195
 196 int gang_4 (int *ary, int size)
 197 {
 198   clear (ary, size);
 199
 200 #pragma acc parallel vector_length(32) copy(ary[0:size]) firstprivate (size)
 201   {
 202 #pragma acc loop auto
 203     for (int jx = 0; jx <  size; jx++)
 204       ary[jx] = place ();
 205   }
 206
 207   return check (ary, size, 1, 0, 1);
 208 }
 209
 210 #define N (32*32*32*2)
 211 int main ()
 212 {
 213   int ondev = 0;
 214
 215 #pragma acc parallel copy(ondev)
 216   {
 217     ondev = acc_on_device (acc_device_not_host);
 218   }
 219   if (!ondev)
 220     return 0;
 221
 222   int ary[N];
 223
 224   if (vector_1 (ary,  N))
 225     return 1;
 226   if (vector_2 (ary,  N))
 227     return 1;
 228
 229   if (worker_1 (ary,  N))
 230     return 1;
 231
 232   if (gang_1 (ary,  N))
 233     return 1;
 234   if (gang_2 (ary,  N))
 235     return 1;
 236   if (gang_3 (ary,  N))
 237     return 1;
 238   if (gang_4 (ary,  N))
 239     return 1;
 240
 241   return 0;
 242 }