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, workersize, vectorsize;
  50   int *gangdist, *workerdist, *vectordist;
  51
  52   for (ix = 0; ix < N;ix++)
  53     ary[ix] = -1;
  54
  55 #define NG 32
  56 #define NW 32
  57 #define VL 32
  58 #pragma acc parallel num_gangs(NG) num_workers(NW) vector_length(VL) \
  59             copy(ary)
  60   /* { dg-note {variable 'ix' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } .-2 } */
  61   {
  62 #pragma acc loop gang worker vector
  63     /* { dg-note {variable 'ix' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } .-1 } */
  64     /* { dg-note {variable 'g' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } .-2 } */
  65     /* { dg-note {variable 'w' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } .-3 } */
  66     /* { dg-note {variable 'v' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } .-4 } */
  67     for (unsigned ix = 0; ix < N; ix++)
  68       {
  69         int g, w, v;
  70
  71         g = __builtin_goacc_parlevel_id (GOMP_DIM_GANG);
  72         w = __builtin_goacc_parlevel_id (GOMP_DIM_WORKER);
  73         v = __builtin_goacc_parlevel_id (GOMP_DIM_VECTOR);
  74
  75         ary[ix] = (g << 16) | (w << 8) | v;
  76       }
  77   }
  78   gangsize = NG;
  79   workersize = NW;
  80   vectorsize = VL;
  81 #if defined ACC_DEVICE_TYPE_host
  82   gangsize = 1;
  83   workersize = 1;
  84   vectorsize = 1;
  85 #elif defined ACC_DEVICE_TYPE_radeon
  86   /* AMD GCN has an upper limit of 'num_workers(16)'.  */
  87   if (workersize > 16)
  88     workersize = 16;
  89   /* AMD GCN uses the autovectorizer for the vector dimension: the use
  90      of a function call in vector-partitioned code in this test is not
  91      currently supported.  */
  92   vectorsize = 1;
  93 #endif
  94
  95   gangdist = (int *) __builtin_alloca (gangsize * sizeof (int));
  96   workerdist = (int *) __builtin_alloca (workersize * sizeof (int));
  97   vectordist = (int *) __builtin_alloca (vectorsize * sizeof (int));
  98   memset (gangdist, 0, gangsize * sizeof (int));
  99   memset (workerdist, 0, workersize * sizeof (int));
 100   memset (vectordist, 0, vectorsize * sizeof (int));
 101
 102   /* Test that work is shared approximately equally amongst each active
 103      gang/worker/vector.  */
 104   for (ix = 0; ix < N; ix++)
 105     {
 106       int g = (ary[ix] >> 16) & 255;
 107       int w = (ary[ix] >> 8) & 255;
 108       int v = ary[ix] & 255;
 109
 110       if (g >= gangsize
 111           || w >= workersize
 112           || v >= vectorsize)
 113         __builtin_abort ();
 114
 115       gangdist[g]++;
 116       workerdist[w]++;
 117       vectordist[v]++;
 118     }
 119
 120   exit = check ("gang", gangdist, gangsize);
 121   exit |= check ("worker", workerdist, workersize);
 122   exit |= check ("vector", vectordist, vectorsize);
 123
 124   return exit;
 125 }