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 #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 }