1  /* { dg-additional-options "-Wopenacc-parallelism" } for testing/documenting
       2     aspects of that functionality.  */
       3  
       4  #include <stdio.h>
       5  #include <openacc.h>
       6  #include <gomp-constants.h>
       7  
       8  #define NUM_WORKERS 16
       9  #ifdef ACC_DEVICE_TYPE_radeon
      10  /* AMD GCN uses the autovectorizer for the vector dimension: the use
      11     of a function call in vector-partitioned code in this test is not
      12     currently supported.  */
      13  #define NUM_VECTORS 1
      14  #else
      15  #define NUM_VECTORS 32
      16  #endif
      17  #define WIDTH 64
      18  #define HEIGHT 32
      19  
      20  #define WORK_ID(I,N)						\
      21    (acc_on_device (acc_device_not_host)				\
      22     ? __builtin_goacc_parlevel_id (GOMP_DIM_WORKER)				\
      23     : (I % N))
      24  #define VEC_ID(I,N)						\
      25    (acc_on_device (acc_device_not_host)				\
      26     ? __builtin_goacc_parlevel_id (GOMP_DIM_VECTOR)				\
      27     : (I % N))
      28  
      29  #pragma acc routine worker
      30  void __attribute__ ((noinline))
      31    WorkVec (int *ptr, int w, int h, int nw, int nv)
      32  {
      33  #pragma acc loop worker
      34    for (int i = 0; i < h; i++)
      35  #pragma acc loop vector
      36      for (int j = 0; j < w; j++)
      37        ptr[i*w + j] = (WORK_ID (i, nw) << 8) | VEC_ID(j, nv);
      38  }
      39  
      40  int DoWorkVec (int nw)
      41  {
      42    int ary[HEIGHT][WIDTH];
      43    int err = 0;
      44  
      45    for (int ix = 0; ix != HEIGHT; ix++)
      46      for (int jx = 0; jx != WIDTH; jx++)
      47        ary[ix][jx] = 0xdeadbeef;
      48  
      49    printf ("spawning %d ...", nw); fflush (stdout);
      50    
      51  #pragma acc parallel num_workers(nw) vector_length (NUM_VECTORS) copy (ary)
      52    /* { dg-warning "region contains vector partitioned code but is not vector partitioned" "" { target openacc_radeon_accel_selected } .-1 } */
      53    {
      54      WorkVec ((int *)ary, WIDTH, HEIGHT, nw, NUM_VECTORS);
      55    }
      56  
      57    for (int ix = 0; ix != HEIGHT; ix++)
      58      for (int jx = 0; jx != WIDTH; jx++)
      59        {
      60  	int exp = ((ix % nw) << 8) | (jx % NUM_VECTORS);
      61  	
      62  	if (ary[ix][jx] != exp)
      63  	  {
      64  	    printf ("\nary[%d][%d] = %#x expected %#x", ix, jx,
      65  		    ary[ix][jx], exp);
      66  	    err = 1;
      67  	  }
      68        }
      69    printf (err ? " failed\n" : " ok\n");
      70    
      71    return err;
      72  }
      73  
      74  int main ()
      75  {
      76    int err = 0;
      77  
      78    for (int W = 1; W <= NUM_WORKERS; W <<= 1)
      79      err |= DoWorkVec (W);
      80  
      81    return err;
      82  }