1  /* To avoid 'error: shared-memory region overflow':
       2     { dg-additional-options "-foffload-options=amdgcn-amdhsa=-mgang-private-size=64" { target openacc_radeon_accel_selected } }
       3  */
       4  
       5  #include <assert.h>
       6  #include <stdio.h>
       7  
       8  #if ACC_DEVICE_TYPE_nvidia
       9  /* To avoid 'libgomp: The Nvidia accelerator has insufficient resources'.  */
      10  #define NUM_WORKERS 28
      11  #else
      12  #define NUM_WORKERS 32
      13  #endif
      14  
      15  #define LOCAL(n) double n = input;
      16  #define LOCALS(n) LOCAL(n##1) LOCAL(n##2) LOCAL(n##3) LOCAL(n##4) \
      17  		  LOCAL(n##5) LOCAL(n##6) LOCAL(n##7) LOCAL(n##8)
      18  #define LOCALS2(n) LOCALS(n##a) LOCALS(n##b) LOCALS(n##c) LOCALS(n##d) \
      19  		   LOCALS(n##e) LOCALS(n##f) LOCALS(n##g) LOCALS(n##h)
      20  
      21  #define USE(n) n
      22  #define USES(n,OP) USE(n##1) OP USE(n##2) OP USE(n##3) OP USE (n##4) OP \
      23  		   USE(n##5) OP USE(n##6) OP USE(n##7) OP USE (n##8)
      24  #define USES2(n,OP) USES(n##a,OP) OP USES(n##b,OP) OP USES(n##c,OP) OP \
      25  		    USES(n##d,OP) OP USES(n##e,OP) OP USES(n##f,OP) OP \
      26  		    USES(n##g,OP) OP USES(n##h,OP)
      27  
      28  int main (void)
      29  {
      30    int ret;
      31    int input = 1;
      32  
      33    #pragma acc parallel num_gangs(1) num_workers(NUM_WORKERS) copyout(ret)
      34    {
      35      int w = 0;
      36      LOCALS2(h);
      37  
      38      #pragma acc loop worker reduction(+:w)
      39      for (int i = 0; i < 32; i++)
      40        {
      41  	int u = USES2(h,+);
      42  	w += u;
      43        }
      44  
      45      printf ("w=%d\n", w);
      46      /* { dg-output "w=2048(\n|\r\n|\r)" } */
      47  
      48      LOCALS2(i);
      49  
      50      #pragma acc loop worker reduction(+:w)
      51      for (int i = 0; i < 32; i++)
      52        {
      53  	int u = USES2(i,+);
      54  	w += u;
      55        }
      56  
      57      printf ("w=%d\n", w);
      58      /* { dg-output "w=4096(\n|\r\n|\r)" } */
      59  
      60      LOCALS2(j);
      61      LOCALS2(k);
      62  
      63      #pragma acc loop worker reduction(+:w)
      64      for (int i = 0; i < 32; i++)
      65        {
      66  	int u = USES2(j,+);
      67  	w += u;
      68        }
      69  
      70      printf ("w=%d\n", w);
      71      /* { dg-output "w=6144(\n|\r\n|\r)" } */
      72  
      73      #pragma acc loop worker reduction(+:w)
      74      for (int i = 0; i < 32; i++)
      75        {
      76  	int u = USES2(k,+);
      77  	w += u;
      78        }
      79  
      80      ret = (w == 64 * 32 * 4);
      81      printf ("w=%d\n", w);
      82      /* { dg-output "w=8192(\n|\r\n|\r)" } */
      83    }
      84  
      85    assert (ret);
      86  
      87    return 0;
      88  }