1 /* { dg-do run } */
2 /* { dg-set-target-env-var OMP_NUM_TEAMS_ALL "3" } */
3 /* { dg-set-target-env-var OMP_NUM_TEAMS_DEV "4" } */
4 /* { dg-set-target-env-var OMP_NUM_TEAMS "5" } */
5 /* { dg-set-target-env-var OMP_NUM_TEAMS_DEV_0 "6" } */
6 /* { dg-set-target-env-var OMP_NUM_TEAMS_DEV_1 "7" } */
7 /* { dg-set-target-env-var OMP_NUM_TEAMS_DEV_2 "8" } */
8 /* { dg-set-target-env-var OMP_TEAMS_THREAD_LIMIT_ALL "2" } */
9 /* { dg-set-target-env-var OMP_TEAMS_THREAD_LIMIT_DEV "3" } */
10 /* { dg-set-target-env-var OMP_TEAMS_THREAD_LIMIT "4" } */
11 /* { dg-set-target-env-var OMP_TEAMS_THREAD_LIMIT_DEV_0 "5" } */
12 /* { dg-set-target-env-var OMP_TEAMS_THREAD_LIMIT_DEV_1 "6" } */
13 /* { dg-set-target-env-var OMP_TEAMS_THREAD_LIMIT_DEV_2 "7" } */
14
15 #include <omp.h>
16 #include <stdlib.h>
17 #include <unistd.h>
18
19 int
20 main ()
21 {
22 if (omp_get_max_teams () != 5
23 || omp_get_teams_thread_limit () != 4)
24 abort ();
25
26 #pragma omp teams
27 {
28 if (omp_get_num_teams () > 5
29 || omp_get_team_num () >= 5)
30 abort ();
31 #pragma omp parallel
32 if (omp_get_thread_limit () > 4
33 || omp_get_thread_num () >= 4)
34 abort ();
35 }
36
37 omp_set_num_teams (4);
38 omp_set_teams_thread_limit (3);
39 if (omp_get_max_teams () != 4
40 || omp_get_teams_thread_limit () != 3)
41 abort ();
42
43 #pragma omp teams
44 {
45 if (omp_get_num_teams () > 4
46 || omp_get_team_num () >= 4)
47 abort ();
48 #pragma omp parallel
49 if (omp_get_thread_limit () > 3
50 || omp_get_thread_num () >= 3)
51 abort ();
52 }
53
54 #pragma omp teams num_teams(3) thread_limit(2)
55 {
56 if (omp_get_num_teams () != 3
57 || omp_get_team_num () >= 3)
58 abort ();
59 #pragma omp parallel
60 if (omp_get_thread_limit () > 2
61 || omp_get_thread_num () >= 2)
62 abort ();
63 }
64
65 #pragma omp teams num_teams(5) thread_limit(4)
66 {
67 if (omp_get_num_teams () != 5
68 || omp_get_team_num () >= 5)
69 abort ();
70 #pragma omp parallel
71 if (omp_get_thread_limit () > 4
72 || omp_get_thread_num () >= 4)
73 abort ();
74 }
75
76 int num_devices = omp_get_num_devices () > 3 ? 3 : omp_get_num_devices ();
77
78 for (int i = 0; i < num_devices; i++)
79 {
80 #pragma omp target device (i)
81 if (omp_get_max_teams () != 6 + i
82 || omp_get_teams_thread_limit () != 5 + i)
83 abort ();
84
85 #pragma omp target device (i)
86 #pragma omp teams
87 #pragma omp parallel
88 if (omp_get_thread_limit () > 5 + i
89 || omp_get_thread_num () >= 5 + i)
90 abort ();
91
92 #pragma omp target device (i)
93 {
94 omp_set_num_teams (5 + i);
95 omp_set_teams_thread_limit (4 + i);
96 if (omp_get_max_teams () != 5 + i
97 || omp_get_teams_thread_limit () != 4 + i)
98 abort ();
99 }
100
101 /* omp_set_num_teams and omp_set_teams_thread_limit above set the value
102 of nteams-var and teams-thread-limit-var ICVs on device 'i', which has
103 scope 'device' and should be avaible in subsequent target regions. */
104 #pragma omp target device (i)
105 if (omp_get_max_teams () != 5 + i
106 || omp_get_teams_thread_limit () != 4 + i)
107 abort ();
108
109 #pragma omp target device (i)
110 #pragma omp teams
111 {
112 if (omp_get_num_teams () > 5 + i
113 || omp_get_team_num () >= 5 + i)
114 abort ();
115 #pragma omp parallel
116 if (omp_get_thread_limit () > 4 + i
117 || omp_get_thread_num () >= 4 + i)
118 abort ();
119 }
120
121 #pragma omp target device (i)
122 #pragma omp teams num_teams(6 + i) thread_limit(5 + i)
123 {
124 if (omp_get_num_teams () > 6 + i
125 || omp_get_team_num () >= 6 + i)
126 abort ();
127 #pragma omp parallel
128 if (omp_get_thread_limit () > 5 + i
129 || omp_get_thread_num () >= 5 + i
130 || omp_get_num_teams () > 6 + i
131 || omp_get_team_num () >= 6 + i)
132 abort ();
133 }
134
135 #pragma omp target device (i)
136 #pragma omp teams num_teams(4 + i) thread_limit(3 + i)
137 {
138 if (omp_get_num_teams () > 4 + i
139 || omp_get_team_num () >= 4 + i)
140 abort ();
141 #pragma omp parallel
142 if (omp_get_thread_limit () > 3 + i
143 || omp_get_thread_num () >= 3 + i
144 || omp_get_num_teams () > 4 + i
145 || omp_get_team_num () >= 4 + i)
146 abort ();
147 }
148
149 #pragma omp target device (i)
150 #pragma omp teams thread_limit(3 + i) num_teams(4 + i)
151 {
152 if (omp_get_num_teams () > 4 + i
153 || omp_get_team_num () >= 4 + i)
154 abort ();
155 #pragma omp parallel
156 if (omp_get_thread_limit () > 3 + i
157 || omp_get_thread_num () >= 3 + i
158 || omp_get_num_teams () > 4 + i
159 || omp_get_team_num () >= 4 + i)
160 abort ();
161 }
162
163 /* The NUM_TEAMS and THREAD_LIMIT clauses should not change the values
164 of the corresponding ICVs. */
165 #pragma omp target device (i)
166 if (omp_get_max_teams () != 5 + i
167 || omp_get_teams_thread_limit () != 4 + i)
168 abort ();
169
170 /* This tests a large number of teams and threads. If it is larger than
171 2^15+1 then the according argument in the kernels arguments list
172 is encoded with two items instead of one. */
173 intptr_t large_num_teams = 66000;
174 intptr_t large_threads_limit = 67000;
175 #pragma omp target device (i)
176 {
177 omp_set_num_teams (large_num_teams + i);
178 omp_set_teams_thread_limit (large_threads_limit + i);
179 if (omp_get_max_teams () != large_num_teams + i
180 || omp_get_teams_thread_limit () != large_threads_limit + i)
181 abort ();
182 }
183
184 #pragma omp target device (i)
185 if (omp_get_max_teams () != large_num_teams + i
186 || omp_get_teams_thread_limit () != large_threads_limit + i)
187 abort ();
188
189 #pragma omp target device (i)
190 #pragma omp teams
191 {
192 if (omp_get_num_teams () > large_num_teams + i
193 || omp_get_team_num () >= large_num_teams + i)
194 abort ();
195 #pragma omp parallel
196 if (omp_get_thread_limit () > large_threads_limit + i
197 || omp_get_thread_num () >= large_threads_limit + i)
198 abort ();
199 }
200 }
201
202 return 0;
203 }