libgomp/testsuite/libgomp.oacc-c-c++-common/loop-dim-default.c

   1 /* { dg-additional-options "-fopenacc-dim=16:16" } */
   2
   3 #include <openacc.h>
   4 #include <alloca.h>
   5 #include <string.h>
   6 #include <stdio.h>
   7 #include <gomp-constants.h>
   8
   9 #pragma acc routine
  10 static int __attribute__ ((noinline)) coord ()
  11 {
  12   int res = 0;
  13
  14   if (acc_on_device (acc_device_not_host))
  15     {
  16       int g, w, v;
  17
  18       g = __builtin_goacc_parlevel_id (GOMP_DIM_GANG);
  19       w = __builtin_goacc_parlevel_id (GOMP_DIM_WORKER);
  20       v = __builtin_goacc_parlevel_id (GOMP_DIM_VECTOR);
  21       res = (1 << 24) | (g << 16) | (w << 8) | v;
  22     }
  23   return res;
  24 }
  25
  26
  27 int check (const int *ary, int size, int gp, int wp, int vp)
  28 {
  29   int exit = 0;
  30   int ix;
  31   int *gangs = (int *)alloca (gp * sizeof (int));
  32   int *workers = (int *)alloca (wp * sizeof (int));
  33   int *vectors = (int *)alloca (vp * sizeof (int));
  34   int offloaded = 0;
  35
  36   memset (gangs, 0, gp * sizeof (int));
  37   memset (workers, 0, wp * sizeof (int));
  38   memset (vectors, 0, vp * sizeof (int));
  39
  40   for (ix = 0; ix < size; ix++)
  41     {
  42       int g = (ary[ix] >> 16) & 0xff;
  43       int w = (ary[ix] >> 8) & 0xff;
  44       int v = (ary[ix] >> 0) & 0xff;
  45
  46       if (g >= gp || w >= wp || v >= vp)
  47         {
  48           printf ("unexpected cpu %#x used\n", ary[ix]);
  49           exit = 1;
  50         }
  51       else
  52         {
  53           vectors[v]++;
  54           workers[w]++;
  55           gangs[g]++;
  56         }
  57       offloaded += ary[ix] >> 24;
  58     }
  59
  60   if (!offloaded)
  61     return 0;
  62
  63   if (offloaded != size)
  64     {
  65       printf ("offloaded %d times,  expected %d\n", offloaded, size);
  66       return 1;
  67     }
  68
  69   for (ix = 0; ix < gp; ix++)
  70     if (gangs[ix] != gangs[0])
  71       {
  72         printf ("gang %d not used %d times\n", ix, gangs[0]);
  73         exit = 1;
  74       }
  75
  76   for (ix = 0; ix < wp; ix++)
  77     if (workers[ix] != workers[0])
  78       {
  79         printf ("worker %d not used %d times\n", ix, workers[0]);
  80         exit = 1;
  81       }
  82
  83   for (ix = 0; ix < vp; ix++)
  84     if (vectors[ix] != vectors[0])
  85       {
  86         printf ("vector %d not used %d times\n", ix, vectors[0]);
  87         exit = 1;
  88       }
  89
  90   return exit;
  91 }
  92
  93 #define N (32 *32*32)
  94
  95 int test_1 (int gp, int wp, int vp)
  96 {
  97   int ary[N];
  98   int exit = 0;
  99
 100 #pragma acc parallel copyout (ary)
 101   {
 102 #pragma acc loop gang (static:1)
 103     for (int ix = 0; ix < N; ix++)
 104       ary[ix] = coord ();
 105   }
 106
 107   exit |= check (ary, N, gp, 1, 1);
 108
 109 #pragma  acc parallel copyout (ary)
 110   {
 111 #pragma acc loop worker
 112     for (int ix = 0; ix < N; ix++)
 113       ary[ix] = coord ();
 114   }
 115
 116   exit |= check (ary, N, 1, wp, 1);
 117
 118 #pragma  acc parallel copyout (ary)
 119   {
 120 #pragma acc loop vector
 121     for (int ix = 0; ix < N; ix++)
 122       ary[ix] = coord ();
 123   }
 124
 125   exit |= check (ary, N, 1, 1, vp);
 126
 127   return exit;
 128 }
 129
 130 int main ()
 131 {
 132   return test_1 (16, 16, 32);
 133 }