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

   1 /* Test of worker-private variables declared in a local scope, broadcasting
   2    to vector-partitioned mode.  Aggregate worker variable.  */
   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 typedef struct
  25 {
  26   int x, y;
  27 } vec2;
  28
  29 int
  30 main (int argc, char* argv[])
  31 {
  32   int i, arr[32 * 32 * 32];
  33
  34   for (i = 0; i < 32 * 32 * 32; i++)
  35     arr[i] = i;
  36
  37   #pragma acc kernels copy(arr) /* { dg-line l_compute[incr c_compute] } */
  38   /* [PR104784] For some reason, for C++, the OpenACC 'kernels' decomposition
  39      decides that a data region is needed for 'j', and subsequently requests it
  40      to be made addressable.
  41      { dg-note {OpenACC 'kernels' decomposition: variable 'j' declared in block requested to be made addressable} {} { target c++ } l_compute$c_compute }
  42      { dg-note {variable 'j' made addressable} {} { target c++ } l_compute$c_compute }
  43      { dg-note {variable 'j' declared in block is candidate for adjusting OpenACC privatization level} {} { target c++ } l_compute$c_compute } */
  44   {
  45     int j;
  46
  47     /* { dg-note {forwarded loop nest in OpenACC 'kernels' region to 'parloops' for analysis} {} { target *-*-* } .+1 } */
  48     #pragma acc loop gang(num:32) /* { dg-line l_loop_i[incr c_loop_i] } */
  49     /* { 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 } */
  50     /* { dg-note {variable 'i' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_loop_i$c_loop_i } */
  51     for (i = 0; i < 32; i++)
  52       {
  53         #pragma acc loop worker(num:32) /* { dg-line l_loop_j[incr c_loop_j] } */
  54         /* { dg-note {variable 'j' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_loop_j$c_loop_j } */
  55         /* { dg-note {variable 'k' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_loop_j$c_loop_j } */
  56         /* { dg-note {variable 'pt' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_loop_j$c_loop_j } */
  57         for (j = 0; j < 32; j++)
  58           {
  59             int k;
  60             vec2 pt;
  61
  62             pt.x = i ^ j * 3;
  63             pt.y = i | j * 5;
  64
  65             #pragma acc loop vector(length:32) /* { dg-line l_loop_k[incr c_loop_k] } */
  66             /* { dg-note {variable 'k' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_loop_k$c_loop_k } */
  67             for (k = 0; k < 32; k++)
  68               arr[i * 1024 + j * 32 + k] += pt.x * k;
  69
  70             #pragma acc loop vector(length:32) /* { dg-line l_loop_k[incr c_loop_k] } */
  71             /* { dg-note {variable 'k' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_loop_k$c_loop_k } */
  72             for (k = 0; k < 32; k++)
  73               arr[i * 1024 + j * 32 + k] += pt.y * k;
  74           }
  75       }
  76     /* { dg-optimized {assigned OpenACC seq loop parallelism} {} { target *-*-* } l_loop_i$c_loop_i } */
  77   }
  78
  79   for (i = 0; i < 32; i++)
  80     for (int j = 0; j < 32; j++)
  81       for (int k = 0; k < 32; k++)
  82         {
  83           int idx = i * 1024 + j * 32 + k;
  84           assert (arr[idx] == idx + (i ^ j * 3) * k + (i | j * 5) * k);
  85         }
  86
  87   return 0;
  88 }