libgomp/testsuite/libgomp.oacc-c-c++-common/routine-wv-2.c

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