libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-private-vars-loop-worker-3.c

   1 /* Test of worker-private variables declared on a loop directive, broadcasting
   2    to vector-partitioned mode.  Back-to-back worker loops.  */
   3
   4 /* { dg-additional-options "--param=openacc-kernels=decompose" } */
   5
   6 /* { dg-additional-options "-fopt-info-omp-all" }
   7    { dg-additional-options "-foffload=-fopt-info-omp-all" } */
   8
   9 /* { dg-additional-options "--param=openacc-privatization=noisy" }
  10    { dg-additional-options "-foffload=--param=openacc-privatization=noisy" }
  11    Prune a few: uninteresting:
  12    { dg-prune-output {note: variable 'D\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} } */
  13
  14 /* It's only with Tcl 8.5 (released in 2007) that "the variable 'varName'
  15    passed to 'incr' may be unset, and in that case, it will be set to [...]",
  16    so to maintain compatibility with earlier Tcl releases, we manually
  17    initialize counter variables:
  18    { dg-line l_dummy[variable c_compute 0 c_loop_i 0 c_loop_j 0 c_loop_k 0] }
  19    { dg-message "dummy" "" { target iN-VAl-Id } l_dummy } to avoid
  20    "WARNING: dg-line var l_dummy defined, but not used".  */
  21
  22 #include <assert.h>
  23
  24 int
  25 main (int argc, char* argv[])
  26 {
  27   int x = 5, i, arr[32 * 32 * 32];
  28
  29   for (i = 0; i < 32 * 32 * 32; i++)
  30     arr[i] = i;
  31
  32   #pragma acc kernels copy(arr) /* { dg-line l_compute[incr c_compute] } */
  33   /* [PR104784] For some reason, for C++, the OpenACC 'kernels' decomposition
  34      decides that a data region is needed for 'j', and subsequently requests it
  35      to be made addressable.
  36      { dg-note {OpenACC 'kernels' decomposition: variable 'j' declared in block requested to be made addressable} {} { target c++ } l_compute$c_compute }
  37      { dg-note {variable 'j' made addressable} {} { target c++ } l_compute$c_compute }
  38      { dg-note {variable 'j' declared in block is candidate for adjusting OpenACC privatization level} {} { target c++ } l_compute$c_compute } */
  39   {
  40     int j;
  41
  42     /* { dg-note {forwarded loop nest in OpenACC 'kernels' region to 'parloops' for analysis} {} { target *-*-* } .+1 } */
  43     #pragma acc loop gang(num:32) /* { dg-line l_loop_i[incr c_loop_i] } */
  44     /* { dg-note {variable 'j' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target c } l_loop_i$c_loop_i } */
  45     /* { dg-note {variable 'i' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_loop_i$c_loop_i } */
  46     for (i = 0; i < 32; i++)
  47       {
  48         #pragma acc loop worker(num:32) private(x) /* { dg-line l_loop_j[incr c_loop_j] } */
  49         /* { dg-note {variable 'x' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_loop_j$c_loop_j } */
  50         /* { dg-note {variable 'j' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_loop_j$c_loop_j } */
  51         /* { dg-note {variable 'k' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_loop_j$c_loop_j } */
  52         for (j = 0; j < 32; j++)
  53           {
  54             int k;
  55             x = i ^ j * 3;
  56
  57             #pragma acc loop vector(length:32) /* { dg-line l_loop_k[incr c_loop_k] } */
  58             /* { dg-note {variable 'k' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_loop_k$c_loop_k } */
  59             for (k = 0; k < 32; k++)
  60               arr[i * 1024 + j * 32 + k] += x * k;
  61           }
  62
  63         #pragma acc loop worker(num:32) private(x) /* { dg-line l_loop_j[incr c_loop_j] } */
  64         /* { dg-note {variable 'x' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_loop_j$c_loop_j } */
  65         /* { dg-note {variable 'j' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_loop_j$c_loop_j } */
  66         /* { dg-note {variable 'k' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_loop_j$c_loop_j } */
  67         for (j = 0; j < 32; j++)
  68           {
  69             int k;
  70             x = i | j * 5;
  71
  72             #pragma acc loop vector(length:32) /* { dg-line l_loop_k[incr c_loop_k] } */
  73             /* { dg-note {variable 'k' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_loop_k$c_loop_k } */
  74             for (k = 0; k < 32; k++)
  75               arr[i * 1024 + j * 32 + k] += x * k;
  76           }
  77       }
  78     /* { dg-optimized {assigned OpenACC seq loop parallelism} {} { target *-*-* } l_loop_i$c_loop_i } */
  79   }
  80
  81   for (i = 0; i < 32; i++)
  82     for (int j = 0; j < 32; j++)
  83       for (int k = 0; k < 32; k++)
  84         {
  85           int idx = i * 1024 + j * 32 + k;
  86           assert (arr[idx] == idx + (i ^ j * 3) * k + (i | j * 5) * k);
  87         }
  88
  89   return 0;
  90 }