libgomp/testsuite/libgomp.oacc-c-c++-common/loop-gwv-2.c

   1 /* { dg-additional-options "-fopt-info-note-omp" }
   2    { dg-additional-options "--param=openacc-privatization=noisy" }
   3    { dg-additional-options "-foffload=-fopt-info-note-omp" }
   4    { dg-additional-options "-foffload=--param=openacc-privatization=noisy" }
   5    for testing/documenting aspects of that functionality.  */
   6
   7 #include <stdio.h>
   8 #include <openacc.h>
   9 #include <string.h>
  10 #include <gomp-constants.h>
  11 #include <stdlib.h>
  12
  13 #if 0
  14 #define DEBUG(DIM, IDX, VAL) \
  15   fprintf (stderr, "%sdist[%d] = %d\n", (DIM), (IDX), (VAL))
  16 #else
  17 #define DEBUG(DIM, IDX, VAL)
  18 #endif
  19
  20 #define N (32*32*32)
  21
  22 int
  23 check (const char *dim, int *dist, int dimsize)
  24 {
  25   int ix;
  26   int exit = 0;
  27
  28   for (ix = 0; ix < dimsize; ix++)
  29     {
  30       DEBUG(dim, ix, dist[ix]);
  31       if (dist[ix] < (N) / (dimsize + 0.5)
  32           || dist[ix] > (N) / (dimsize - 0.5))
  33         {
  34           fprintf (stderr, "did not distribute to %ss (%d not between %d "
  35                    "and %d)\n", dim, dist[ix], (int) ((N) / (dimsize + 0.5)),
  36                    (int) ((N) / (dimsize - 0.5)));
  37           exit |= 1;
  38         }
  39     }
  40
  41   return exit;
  42 }
  43
  44 int main ()
  45 {
  46   int ary[N];
  47   int ix;
  48   int exit = 0;
  49   int gangsize = 0, workersize = 0, vectorsize = 0;
  50   int *gangdist, *workerdist, *vectordist;
  51
  52   for (ix = 0; ix < N;ix++)
  53     ary[ix] = -1;
  54
  55 #pragma acc parallel num_gangs(32) num_workers(32) vector_length(32) \
  56             copy(ary) copyout(gangsize, workersize, vectorsize)
  57   /* { dg-note {variable 'ix' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } .-2 } */
  58   {
  59 #pragma acc loop gang worker vector
  60     /* { dg-note {variable 'ix' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } .-1 } */
  61     /* { dg-note {variable 'g' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } .-2 } */
  62     /* { dg-note {variable 'w' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } .-3 } */
  63     /* { dg-note {variable 'v' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } .-4 } */
  64     for (unsigned ix = 0; ix < N; ix++)
  65       {
  66         int g, w, v;
  67
  68         g = __builtin_goacc_parlevel_id (GOMP_DIM_GANG);
  69         w = __builtin_goacc_parlevel_id (GOMP_DIM_WORKER);
  70         v = __builtin_goacc_parlevel_id (GOMP_DIM_VECTOR);
  71
  72         ary[ix] = (g << 16) | (w << 8) | v;
  73       }
  74
  75     gangsize = __builtin_goacc_parlevel_size (GOMP_DIM_GANG);
  76     workersize = __builtin_goacc_parlevel_size (GOMP_DIM_WORKER);
  77     vectorsize = __builtin_goacc_parlevel_size (GOMP_DIM_VECTOR);
  78   }
  79
  80   gangdist = (int *) __builtin_alloca (gangsize * sizeof (int));
  81   workerdist = (int *) __builtin_alloca (workersize * sizeof (int));
  82   vectordist = (int *) __builtin_alloca (vectorsize * sizeof (int));
  83   memset (gangdist, 0, gangsize * sizeof (int));
  84   memset (workerdist, 0, workersize * sizeof (int));
  85   memset (vectordist, 0, vectorsize * sizeof (int));
  86
  87   /* Test that work is shared approximately equally amongst each active
  88      gang/worker/vector.  */
  89   for (ix = 0; ix < N; ix++)
  90     {
  91       int g = (ary[ix] >> 16) & 255;
  92       int w = (ary[ix] >> 8) & 255;
  93       int v = ary[ix] & 255;
  94
  95       gangdist[g]++;
  96       workerdist[w]++;
  97       vectordist[v]++;
  98     }
  99
 100   exit = check ("gang", gangdist, gangsize);
 101   exit |= check ("worker", workerdist, workersize);
 102   exit |= check ("vector", vectordist, vectorsize);
 103
 104   return exit;
 105 }