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