libgomp/testsuite/libgomp.oacc-c-c++-common/private-atomic-1-gang.c

   1 /* Tests for gang-private variables, 'atomic' access */
   2
   3 /* { dg-additional-options "-fopt-info-note-omp" }
   4    { dg-additional-options "--param=openacc-privatization=noisy" }
   5    { dg-additional-options "-foffload=-fopt-info-note-omp" }
   6    { dg-additional-options "-foffload=--param=openacc-privatization=noisy" }
   7    for testing/documenting aspects of that functionality.  */
   8
   9 /* It's only with Tcl 8.5 (released in 2007) that "the variable 'varName'
  10    passed to 'incr' may be unset, and in that case, it will be set to [...]",
  11    so to maintain compatibility with earlier Tcl releases, we manually
  12    initialize counter variables:
  13    { dg-line l_dummy[variable c_compute 0 c_loop 0] }
  14    { dg-message "dummy" "" { target iN-VAl-Id } l_dummy } to avoid
  15    "WARNING: dg-line var l_dummy defined, but not used".  */
  16
  17 #include <assert.h>
  18 #include <openacc.h>
  19
  20 int main (void)
  21 {
  22   int ret;
  23
  24
  25   ret = 0;
  26   #pragma acc parallel num_gangs(1444) num_workers(32) reduction(+: ret) /* { dg-line l_compute[incr c_compute] } */
  27   /* { dg-note {variable 'w' declared in block is candidate for adjusting OpenACC privatization level} "" { target *-*-* } l_compute$c_compute }
  28      { dg-note {variable 'w' ought to be adjusted for OpenACC privatization level: 'gang'} "" { target *-*-* } l_compute$c_compute }
  29      { dg-note {variable 'w' adjusted for OpenACC privatization level: 'gang'} "" { target { ! openacc_host_selected } } l_compute$c_compute } */
  30   /* { dg-note {variable 'i' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_compute$c_compute } */
  31   {
  32     int w = -22;
  33
  34     #pragma acc loop worker /* { dg-line l_loop[incr c_loop] } */
  35     /* { dg-note {variable 'i' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop$c_loop } */
  36     for (int i = 0; i < 2232; i++)
  37       {
  38         #pragma acc atomic update
  39         w++;
  40       }
  41
  42     ret = (w == -22 + 2232);
  43   }
  44   if (acc_get_device_type () == acc_device_host)
  45     assert (ret == 1);
  46   else
  47     assert (ret == 1444);
  48
  49
  50   ret = 0;
  51   #pragma acc parallel num_gangs(1414) vector_length(32) reduction(+: ret) /* { dg-line l_compute[incr c_compute] } */
  52   /* { dg-note {variable 'v' declared in block is candidate for adjusting OpenACC privatization level} "" { target *-*-* } l_compute$c_compute }
  53      { dg-note {variable 'v' ought to be adjusted for OpenACC privatization level: 'gang'} "" { target *-*-* } l_compute$c_compute }
  54      { dg-note {variable 'v' adjusted for OpenACC privatization level: 'gang'} "" { target { ! openacc_host_selected } } l_compute$c_compute } */
  55   /* { dg-note {variable 'i' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_compute$c_compute } */
  56   {
  57     int v = 10;
  58
  59     #pragma acc loop vector /* { dg-line l_loop[incr c_loop] } */
  60     /* { dg-note {variable 'i' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop$c_loop } */
  61     for (int i = 0; i < 3201; i++)
  62       {
  63         #pragma acc atomic update
  64         v++;
  65       }
  66
  67     ret = (v == 10 + 3201);
  68   }
  69   if (acc_get_device_type () == acc_device_host)
  70     assert (ret == 1);
  71   else
  72     assert (ret == 1414);
  73
  74
  75   ret = 0;
  76 #pragma acc parallel num_gangs(314) reduction(+: ret) /* { dg-line l_compute[incr c_compute] } */
  77   /* { dg-note {variable 'v' declared in block is candidate for adjusting OpenACC privatization level} "" { target *-*-* } l_compute$c_compute }
  78      { dg-note {variable 'v' ought to be adjusted for OpenACC privatization level: 'gang'} "" { target *-*-* } l_compute$c_compute }
  79      { dg-note {variable 'v' adjusted for OpenACC privatization level: 'gang'} "" { target { ! openacc_host_selected } } l_compute$c_compute } */
  80   {
  81     int v = -222;
  82
  83 #pragma acc atomic update
  84     ++v;
  85 #pragma acc atomic update
  86     ++v;
  87 #pragma acc atomic update
  88     ++v;
  89
  90     ret += (v == -222 + 3);
  91   }
  92   if (acc_get_device_type () == acc_device_host)
  93     assert (ret == 1);
  94   else
  95     assert (ret == 314);
  96
  97
  98   return 0;
  99 }