1  /* Test "big" private data.  */
       2  
       3  /* { dg-additional-options -fno-inline } for stable results regarding OpenACC 'routine'.  */
       4  
       5  /* { dg-additional-options -fopt-info-all-omp }
       6     { dg-additional-options --param=openacc-privatization=noisy }
       7     { dg-additional-options -foffload=-fopt-info-all-omp }
       8     { dg-additional-options -foffload=--param=openacc-privatization=noisy }
       9     for testing/documenting aspects of that functionality.  */
      10  
      11  /* { dg-additional-options -Wopenacc-parallelism } for testing/documenting
      12     aspects of that functionality.  */
      13  
      14  /* For GCN offloading compilation, we (expectedly) run into a
      15     'gang-private data-share memory exhausted' error: the default
      16     '-mgang-private-size' is too small.  Raise it so that 'uint32_t x[344]' plus
      17     some internal-use data fits in:
      18     { dg-additional-options -foffload-options=amdgcn-amdhsa=-mgang-private-size=1555 { target openacc_radeon_accel_selected } } */
      19  
      20  /* It's only with Tcl 8.5 (released in 2007) that "the variable 'varName'
      21     passed to 'incr' may be unset, and in that case, it will be set to [...]",
      22     so to maintain compatibility with earlier Tcl releases, we manually
      23     initialize counter variables:
      24     { dg-line l_dummy[variable c_compute 0 c_loop 0] }
      25     { dg-message dummy {} { target iN-VAl-Id } l_dummy } to avoid
      26     "WARNING: dg-line var l_dummy defined, but not used".  */
      27  
      28  #include <assert.h>
      29  #include <stdint.h>
      30  
      31  
      32  /* Based on 'private-variables.c:loop_g_5'.  */
      33  
      34  /* To demonstrate PR105421 "GCN offloading, raised '-mgang-private-size':
      35     'HSA_STATUS_ERROR_MEMORY_APERTURE_VIOLATION'", a 'struct' indirection, for
      36     example, has been necessary in combination with a separate routine.  */
      37  
      38  struct data
      39  {
      40    uint32_t *x;
      41    uint32_t *arr;
      42    uint32_t i;
      43  };
      44  
      45  #pragma acc routine worker
      46  static void
      47  loop_g_5_r(struct data *data)
      48  {
      49    uint32_t *x = data->x;
      50    uint32_t *arr = data->arr;
      51    uint32_t i = data->i;
      52  
      53  #pragma acc loop /* { dg-line l_loop[incr c_loop] } */
      54    /* { dg-note {variable 'j' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_loop$c_loop } */
      55    /* { dg-optimized {assigned OpenACC worker vector loop parallelism} {} { target *-*-* } l_loop$c_loop } */
      56    for (int j = 0; j < 320; j++)
      57      arr[i * 320 + j] += x[(i * 320 + j) % 344];
      58  }
      59  
      60  void loop_g_5()
      61  {
      62    uint32_t x[344], i, arr[320 * 320];
      63  
      64    for (i = 0; i < 320 * 320; i++)
      65      arr[i] = i;
      66  
      67    #pragma acc parallel copy(arr)
      68    {
      69      #pragma acc loop gang private(x) /* { dg-line l_loop[incr c_loop] } */
      70      /* { dg-note {variable 'x' in 'private' clause is candidate for adjusting OpenACC privatization level} {} { target *-*-* } l_loop$c_loop }
      71         { dg-note {variable 'x' ought to be adjusted for OpenACC privatization level: 'gang'} {} { target *-*-* } l_loop$c_loop }
      72         { dg-note {variable 'x' adjusted for OpenACC privatization level: 'gang'} {} { target { ! openacc_host_selected } } l_loop$c_loop } */
      73      /* { dg-note {variable 'i' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_loop$c_loop } */
      74      /* { dg-note {variable 'data' declared in block is candidate for adjusting OpenACC privatization level} {} { target *-*-* } l_loop$c_loop }
      75         { dg-note {variable 'data' ought to be adjusted for OpenACC privatization level: 'gang'} {} { target *-*-* } l_loop$c_loop }
      76         { dg-note {variable 'data' adjusted for OpenACC privatization level: 'gang'} {} { target { ! openacc_host_selected } } l_loop$c_loop } */
      77      /* { dg-note {variable 'j' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_loop$c_loop } */
      78      /* { dg-optimized {assigned OpenACC gang loop parallelism} {} { target *-*-* } l_loop$c_loop } */
      79      for (i = 0; i < 320; i++)
      80        {
      81          for (int j = 0; j < 344; j++)
      82  	  x[j] = j * (2 + i);
      83  
      84  	struct data data = { x, arr, i };
      85  	loop_g_5_r(&data); /* { dg-line l_compute[incr c_compute] } */
      86  	/* { dg-optimized {assigned OpenACC worker vector loop parallelism} {} { target *-*-* } l_compute$c_compute } */
      87        }
      88    }
      89  
      90    for (i = 0; i < 320 * 320; i++)
      91      assert(arr[i] == i + (i % 344) * (2 + (i / 320)));
      92  }
      93  
      94  
      95  int main ()
      96  {
      97    loop_g_5();
      98  
      99    return 0;
     100  }