1 // RUN: %compile-run-and-check
2
3 #include <omp.h>
4 #include <stdio.h>
5
6 const int MaxThreads = 1024;
7 const int NumThreads = 64;
8
main(int argc,char * argv[])9 int main(int argc, char *argv[]) {
10 int level = -1, activeLevel = -1;
11 // The expected value is -1, initialize to different value.
12 int ancestorTNumNeg = 1, teamSizeNeg = 1;
13 int ancestorTNum0 = -1, teamSize0 = -1;
14 // The expected value is -1, initialize to different value.
15 int ancestorTNum1 = 1, teamSize1 = 1;
16 int check1[MaxThreads];
17 int check2[MaxThreads];
18 int check3[MaxThreads];
19 int check4[MaxThreads];
20 for (int i = 0; i < MaxThreads; i++) {
21 check1[i] = check2[i] = check3[i] = check4[i] = 0;
22 }
23
24 #pragma omp target map(level, activeLevel, ancestorTNumNeg, teamSizeNeg) \
25 map(ancestorTNum0, teamSize0, ancestorTNum1, teamSize1) \
26 map(check1[:], check2[:], check3[:], check4[:])
27 {
28 level = omp_get_level();
29 activeLevel = omp_get_active_level();
30
31 // Expected to return -1.
32 ancestorTNumNeg = omp_get_ancestor_thread_num(-1);
33 teamSizeNeg = omp_get_team_size(-1);
34
35 // Expected to return 0 and 1.
36 ancestorTNum0 = omp_get_ancestor_thread_num(0);
37 teamSize0 = omp_get_team_size(0);
38
39 // Expected to return -1 because the requested level is larger than
40 // the nest level.
41 ancestorTNum1 = omp_get_ancestor_thread_num(1);
42 teamSize1 = omp_get_team_size(1);
43
44 // Expecting active parallel region.
45 #pragma omp parallel num_threads(NumThreads)
46 {
47 int id = omp_get_thread_num();
48 // Multiply return value of omp_get_level by 5 to avoid that this test
49 // passes if both API calls return wrong values.
50 check1[id] += omp_get_level() * 5 + omp_get_active_level();
51
52 // Expected to return 0 and 1.
53 check2[id] += omp_get_ancestor_thread_num(0) + 5 * omp_get_team_size(0);
54 // Expected to return the current thread num.
55 check2[id] += (omp_get_ancestor_thread_num(1) - id);
56 // Expected to return the current number of threads.
57 check2[id] += 3 * omp_get_team_size(1);
58 // Expected to return -1, see above.
59 check2[id] += omp_get_ancestor_thread_num(2) + omp_get_team_size(2);
60
61 // Expecting serialized parallel region.
62 #pragma omp parallel
63 {
64 #pragma omp atomic
65 check3[id] += omp_get_level() * 5 + omp_get_active_level();
66
67 // Expected to return 0 and 1.
68 int check4Inc = omp_get_ancestor_thread_num(0) + 5 * omp_get_team_size(0);
69 // Expected to return the parent thread num.
70 check4Inc += (omp_get_ancestor_thread_num(1) - id);
71 // Expected to return the number of threads in the active parallel region.
72 check4Inc += 3 * omp_get_team_size(1);
73 // Expected to return 0 and 1.
74 check4Inc += omp_get_ancestor_thread_num(2) + 3 * omp_get_team_size(2);
75 // Expected to return -1, see above.
76 check4Inc += omp_get_ancestor_thread_num(3) + omp_get_team_size(3);
77
78 #pragma omp atomic
79 check4[id] += check4Inc;
80 }
81 }
82 }
83
84 // CHECK: target: level = 0, activeLevel = 0
85 printf("target: level = %d, activeLevel = %d\n", level, activeLevel);
86 // CHECK: level = -1: ancestorTNum = -1, teamSize = -1
87 printf("level = -1: ancestorTNum = %d, teamSize = %d\n", ancestorTNumNeg, teamSizeNeg);
88 // CHECK: level = 0: ancestorTNum = 0, teamSize = 1
89 printf("level = 0: ancestorTNum = %d, teamSize = %d\n", ancestorTNum0, teamSize0);
90 // CHECK: level = 1: ancestorTNum = -1, teamSize = -1
91 printf("level = 1: ancestorTNum = %d, teamSize = %d\n", ancestorTNum1, teamSize1);
92
93 // CHECK-NOT: invalid
94 for (int i = 0; i < MaxThreads; i++) {
95 // Check active parallel region:
96 // omp_get_level() = 1, omp_get_active_level() = 1
97 const int Expected1 = 6;
98 if (i < NumThreads) {
99 if (check1[i] != Expected1) {
100 printf("invalid: check1[%d] should be %d, is %d\n", i, Expected1, check1[i]);
101 }
102 } else if (check1[i] != 0) {
103 printf("invalid: check1[%d] should be 0, is %d\n", i, check1[i]);
104 }
105
106 // 5 * 1 + 3 * 64 - 1 - 1 (see above)
107 const int Expected2 = 195;
108 if (i < NumThreads) {
109 if (check2[i] != Expected2) {
110 printf("invalid: check2[%d] should be %d, is %d\n", i, Expected2, check2[i]);
111 }
112 } else if (check2[i] != 0) {
113 printf("invalid: check2[%d] should be 0, is %d\n", i, check2[i]);
114 }
115
116 // Check serialized parallel region:
117 // omp_get_level() = 2, omp_get_active_level() = 1
118 const int Expected3 = 11;
119 if (i < NumThreads) {
120 if (check3[i] != Expected3) {
121 printf("invalid: check3[%d] should be %d, is %d\n", i, Expected3, check3[i]);
122 }
123 } else if (check3[i] != 0) {
124 printf("invalid: check3[%d] should be 0, is %d\n", i, check3[i]);
125 }
126
127 // 5 * 1 + 3 * 64 + 3 * 1 - 1 - 1 (see above)
128 const int Expected4 = 198;
129 if (i < NumThreads) {
130 if (check4[i] != Expected4) {
131 printf("invalid: check4[%d] should be %d, is %d\n", i, Expected4, check4[i]);
132 }
133 } else if (check4[i] != 0) {
134 printf("invalid: check4[%d] should be 0, is %d\n", i, check4[i]);
135 }
136 }
137
138 // Check for paraller level in non-SPMD kernels.
139 level = 0;
140 #pragma omp target teams distribute num_teams(1) thread_limit(32) reduction(+:level)
141 for (int i=0; i<5032; i+=32) {
142 int ub = (i+32 > 5032) ? 5032 : i+32;
143 #pragma omp parallel for schedule(dynamic)
144 for (int j=i ; j < ub; j++) ;
145 level += omp_get_level();
146 }
147 // CHECK: Integral level = 0.
148 printf("Integral level = %d.\n", level);
149
150 return 0;
151 }
152