]> git.ipfire.org Git - thirdparty/gcc.git/blame - libgomp/testsuite/libgomp.oacc-c-c++-common/routine-wv-1.c
[openacc] Add __builtin_goacc_parlevel_{id,size}
[thirdparty/gcc.git] / libgomp / testsuite / libgomp.oacc-c-c++-common / routine-wv-1.c
CommitLineData
36fa0f5f 1#include <stdio.h>
1f62d637
TV
2#include <openacc.h>
3#include <gomp-constants.h>
36fa0f5f
NS
4
5#define N (32*32*32+17)
6
7#pragma acc routine worker
8void __attribute__ ((noinline)) worker (int ary[N])
9{
10#pragma acc loop worker vector
11 for (unsigned ix = 0; ix < N; ix++)
12 {
1f62d637 13 if (acc_on_device (acc_device_not_host))
36fa0f5f 14 {
1f62d637 15 int g, w, v;
36fa0f5f 16
1f62d637
TV
17 g = __builtin_goacc_parlevel_id (GOMP_DIM_GANG);
18 w = __builtin_goacc_parlevel_id (GOMP_DIM_WORKER);
19 v = __builtin_goacc_parlevel_id (GOMP_DIM_VECTOR);
36fa0f5f
NS
20 ary[ix] = (g << 16) | (w << 8) | v;
21 }
22 else
23 ary[ix] = ix;
24 }
25}
26
27int main ()
28{
29 int ary[N];
30 int ix;
31 int exit = 0;
32 int ondev = 0;
33
34 for (ix = 0; ix < N;ix++)
35 ary[ix] = -1;
36
37#pragma acc parallel num_workers(32) vector_length(32) copy(ary) copy(ondev)
38 {
1f62d637 39 ondev = acc_on_device (acc_device_not_host);
36fa0f5f
NS
40 worker (ary);
41 }
42
43 for (ix = 0; ix < N; ix++)
44 {
45 int expected = ix;
46 if(ondev)
47 {
48 int g = 0;
49 int w = (ix / 32) % 32;
50 int v = ix % 32;
51
52 expected = (g << 16) | (w << 8) | v;
53 }
54
55 if (ary[ix] != expected)
56 {
57 exit = 1;
58 printf ("ary[%d]=%x expected %x\n", ix, ary[ix], expected);
59 }
60 }
61
62 return exit;
63}