1// RUN: %libomptarget-compile-run-and-check-generic
2// RUN: %libomptarget-compileopt-run-and-check-generic
3
4// UNSUPPORTED: aarch64-unknown-linux-gnu
5// UNSUPPORTED: aarch64-unknown-linux-gnu-LTO
6// UNSUPPORTED: x86_64-pc-linux-gnu
7// UNSUPPORTED: x86_64-pc-linux-gnu-LTO
8// UNSUPPORTED: s390x-ibm-linux-gnu
9// UNSUPPORTED: s390x-ibm-linux-gnu-LTO
10
11// clang-format off
12
13#include <omp.h>
14#include <stdio.h>
15
16#define N 100
17#define BLOCK_SHIFT 8
18
19void print(int *A, int size) {
20 for (int i = 0; i < size; ++i) {
21 printf(format: "B%dT%d ", A[i] >> BLOCK_SHIFT, A[i] % (1 << BLOCK_SHIFT));
22 }
23 printf(format: "\n");
24}
25
26int main() {
27 int A[N];
28
29#pragma omp target parallel for map(from:A) num_threads(10) schedule(static, 2)
30 for (int i = 0; i < N; ++i) {
31 A[i] = (omp_get_team_num() << BLOCK_SHIFT) + omp_get_thread_num();
32 }
33 printf(format: "omp target parallel for thread chunk size %d\n", 2);
34 print(A, N);
35
36#pragma omp target teams distribute map(from:A) num_teams(10) \
37 dist_schedule(static, 2)
38 for (int i = 0; i < N; ++i) {
39 A[i] = (omp_get_team_num() << BLOCK_SHIFT) + omp_get_thread_num();
40 }
41 printf(format: "omp target teams distribute block chunk size %d\n", 2);
42 print(A, N);
43
44#pragma omp target teams distribute parallel for map(from:A) \
45 num_teams(10) dist_schedule(static, 2)
46 for (int i = 0; i < N; ++i) {
47 A[i] = (omp_get_team_num() << BLOCK_SHIFT) + omp_get_thread_num();
48 }
49 printf(format: "omp target teams distribute parallel for block chunk size %d ", 2);
50 printf(format: "thread chunk size default\n");
51 print(A, N);
52
53#pragma omp target teams distribute parallel for map(from:A) \
54 num_teams(10) dist_schedule(static, 2) schedule(static, 3)
55 for (int i = 0; i < N; ++i) {
56 A[i] = (omp_get_team_num() << BLOCK_SHIFT) + omp_get_thread_num();
57 }
58 printf(format: "omp target teams distribute parallel for block chunk size %d ", 2);
59 printf(format: "thread chunk size %d\n", 3);
60 print(A, N);
61
62#pragma omp target teams distribute parallel for map(from:A) \
63 num_teams(10) dist_schedule(static, 3) schedule(static, 2)
64 for (int i = 0; i < N; ++i) {
65 A[i] = (omp_get_team_num() << BLOCK_SHIFT) + omp_get_thread_num();
66 }
67 printf(format: "omp target teams distribute parallel for block chunk size %d ", 3);
68 printf(format: "thread chunk size %d\n", 2);
69 print(A, N);
70
71#pragma omp target teams distribute parallel for map(from:A) \
72 num_teams(10) dist_schedule(static, 5) schedule(static, 2)
73 for (int i = 0; i < N; ++i) {
74 A[i] = (omp_get_team_num() << BLOCK_SHIFT) + omp_get_thread_num();
75 }
76 printf(format: "omp target teams distribute parallel for block chunk size %d ", 5);
77 printf(format: "thread chunk size %d\n", 2);
78 print(A, N);
79
80#pragma omp target teams distribute parallel for map(from:A) num_teams(10) \
81 dist_schedule(static, 49) schedule(static, 2)
82 for (int i = 0; i < N; ++i) {
83 A[i] = (omp_get_team_num() << BLOCK_SHIFT) + omp_get_thread_num();
84 }
85 printf(format: "omp target teams distribute parallel for block chunk size %d ", 49);
86 printf(format: "thread chunk size %d\n", 2);
87 print(A, N);
88
89#pragma omp target teams distribute parallel for map(from:A) \
90 num_teams(10) num_threads(10) dist_schedule(static, 29)
91 for (int i = 0; i < N; ++i) {
92 A[i] = (omp_get_team_num() << BLOCK_SHIFT) + omp_get_thread_num();
93 }
94 printf(format: "omp target teams distribute parallel for block chunk size %d ", 29);
95 printf(format: "thread chunk size default\n");
96 print(A, N);
97
98#pragma omp target teams distribute parallel for map(from:A) \
99 num_teams(10) num_threads(10) dist_schedule(static, 101)
100 for (int i = 0; i < N; ++i) {
101 A[i] = (omp_get_team_num() << BLOCK_SHIFT) + omp_get_thread_num();
102 }
103 printf(format: "omp target teams distribute parallel for block chunk size %d ", 101);
104 printf(format: "thread chunk size default\n");
105 print(A, N);
106
107#pragma omp target teams distribute parallel for map(from:A) \
108 num_teams(9) num_threads(10) schedule(static, 101)
109 for (int i = 0; i < N; ++i) {
110 A[i] = (omp_get_team_num() << BLOCK_SHIFT) + omp_get_thread_num();
111 }
112 printf(format: "omp target teams distribute parallel for default block chunk size ");
113 printf(format: "thread chunk size %d\n", 101);
114 print(A, N);
115 return 0;
116}
117//CHECK: omp target parallel for thread chunk size 2
118
119//CHECK-NEXT: B0T0 B0T0 B0T1 B0T1 B0T2 B0T2 B0T3 B0T3 B0T4 B0T4
120//CHECK-SAME: B0T5 B0T5 B0T6 B0T6 B0T7 B0T7 B0T8 B0T8 B0T9 B0T9
121//CHECK-SAME: B0T0 B0T0 B0T1 B0T1 B0T2 B0T2 B0T3 B0T3 B0T4 B0T4
122//CHECK-SAME: B0T5 B0T5 B0T6 B0T6 B0T7 B0T7 B0T8 B0T8 B0T9 B0T9
123//CHECK-SAME: B0T0 B0T0 B0T1 B0T1 B0T2 B0T2 B0T3 B0T3 B0T4 B0T4
124//CHECK-SAME: B0T5 B0T5 B0T6 B0T6 B0T7 B0T7 B0T8 B0T8 B0T9 B0T9
125//CHECK-SAME: B0T0 B0T0 B0T1 B0T1 B0T2 B0T2 B0T3 B0T3 B0T4 B0T4
126//CHECK-SAME: B0T5 B0T5 B0T6 B0T6 B0T7 B0T7 B0T8 B0T8 B0T9 B0T9
127//CHECK-SAME: B0T0 B0T0 B0T1 B0T1 B0T2 B0T2 B0T3 B0T3 B0T4 B0T4
128//CHECK-SAME: B0T5 B0T5 B0T6 B0T6 B0T7 B0T7 B0T8 B0T8 B0T9 B0T9
129
130//CHECK: omp target teams distribute block chunk size 2
131
132//CHECK-NEXT: B0T0 B0T0 B1T0 B1T0 B2T0 B2T0 B3T0 B3T0 B4T0 B4T0
133//CHECK-SAME: B5T0 B5T0 B6T0 B6T0 B7T0 B7T0 B8T0 B8T0 B9T0 B9T0
134//CHECK-SAME: B0T0 B0T0 B1T0 B1T0 B2T0 B2T0 B3T0 B3T0 B4T0 B4T0
135//CHECK-SAME: B5T0 B5T0 B6T0 B6T0 B7T0 B7T0 B8T0 B8T0 B9T0 B9T0
136//CHECK-SAME: B0T0 B0T0 B1T0 B1T0 B2T0 B2T0 B3T0 B3T0 B4T0 B4T0
137//CHECK-SAME: B5T0 B5T0 B6T0 B6T0 B7T0 B7T0 B8T0 B8T0 B9T0 B9T0
138//CHECK-SAME: B0T0 B0T0 B1T0 B1T0 B2T0 B2T0 B3T0 B3T0 B4T0 B4T0
139//CHECK-SAME: B5T0 B5T0 B6T0 B6T0 B7T0 B7T0 B8T0 B8T0 B9T0 B9T0
140//CHECK-SAME: B0T0 B0T0 B1T0 B1T0 B2T0 B2T0 B3T0 B3T0 B4T0 B4T0
141//CHECK-SAME: B5T0 B5T0 B6T0 B6T0 B7T0 B7T0 B8T0 B8T0 B9T0 B9T0
142
143//CHECK: omp target teams distribute parallel for
144//CHECK-SAME: block chunk size 2 thread chunk size default
145
146//CHECK-NEXT: B0T0 B0T1 B1T0 B1T1 B2T0 B2T1 B3T0 B3T1 B4T0 B4T1
147//CHECK-SAME: B5T0 B5T1 B6T0 B6T1 B7T0 B7T1 B8T0 B8T1 B9T0 B9T1
148//CHECK-SAME: B0T0 B0T1 B1T0 B1T1 B2T0 B2T1 B3T0 B3T1 B4T0 B4T1
149//CHECK-SAME: B5T0 B5T1 B6T0 B6T1 B7T0 B7T1 B8T0 B8T1 B9T0 B9T1
150//CHECK-SAME: B0T0 B0T1 B1T0 B1T1 B2T0 B2T1 B3T0 B3T1 B4T0 B4T1
151//CHECK-SAME: B5T0 B5T1 B6T0 B6T1 B7T0 B7T1 B8T0 B8T1 B9T0 B9T1
152
153//CHECK: omp target teams distribute parallel for
154//CHECK-SAME block chunk size 2 thread chunk size 3
155
156//CHECK-NEXT: B0T0 B0T0 B1T0 B1T0 B2T0 B2T0 B3T0 B3T0 B4T0 B4T0
157//CHECK-SAME: B5T0 B5T0 B6T0 B6T0 B7T0 B7T0 B8T0 B8T0 B9T0 B9T0
158//CHECK-SAME: B0T0 B0T0 B1T0 B1T0 B2T0 B2T0 B3T0 B3T0 B4T0 B4T0
159//CHECK-SAME: B5T0 B5T0 B6T0 B6T0 B7T0 B7T0 B8T0 B8T0 B9T0 B9T0
160//CHECK-SAME: B0T0 B0T0 B1T0 B1T0 B2T0 B2T0 B3T0 B3T0 B4T0 B4T0
161//CHECK-SAME: B5T0 B5T0 B6T0 B6T0 B7T0 B7T0 B8T0 B8T0 B9T0 B9T0
162//CHECK-SAME: B0T0 B0T0 B1T0 B1T0 B2T0 B2T0 B3T0 B3T0 B4T0 B4T0
163//CHECK-SAME: B5T0 B5T0 B6T0 B6T0 B7T0 B7T0 B8T0 B8T0 B9T0 B9T0
164//CHECK-SAME: B0T0 B0T0 B1T0 B1T0 B2T0 B2T0 B3T0 B3T0 B4T0 B4T0
165//CHECK-SAME: B5T0 B5T0 B6T0 B6T0 B7T0 B7T0 B8T0 B8T0 B9T0 B9T0
166
167//CHECK: omp target teams distribute parallel for
168//CHECK-SAME: block chunk size 3 thread chunk size 2
169
170//CHECK-NEXT: B0T0 B0T0 B0T1 B1T0 B1T0 B1T1 B2T0 B2T0 B2T1
171//CHECK-SAME: B3T0 B3T0 B3T1 B4T0 B4T0 B4T1
172//CHECK-SAME: B5T0 B5T0 B5T1 B6T0 B6T0 B6T1 B7T0 B7T0 B7T1
173//CHECK-SAME: B8T0 B8T0 B8T1 B9T0 B9T0 B9T1
174//CHECK-SAME: B0T0 B0T0 B0T1 B1T0 B1T0 B1T1 B2T0 B2T0 B2T1
175//CHECK-SAME: B3T0 B3T0 B3T1 B4T0 B4T0 B4T1
176//CHECK-SAME: B5T0 B5T0 B5T1 B6T0 B6T0 B6T1 B7T0 B7T0 B7T1
177//CHECK-SAME: B8T0 B8T0 B8T1 B9T0 B9T0 B9T1
178//CHECK-SAME: B0T0 B0T0 B0T1 B1T0 B1T0 B1T1 B2T0 B2T0 B2T1
179//CHECK-SAME: B3T0 B3T0 B3T1 B4T0 B4T0 B4T1
180//CHECK-SAME: B5T0 B5T0 B5T1 B6T0 B6T0 B6T1 B7T0 B7T0 B7T1
181//CHECK-SAME: B8T0 B8T0 B8T1 B9T0 B9T0 B9T1
182//CHECK-SAME: B0T0 B0T0 B0T1 B1T0 B1T0 B1T1 B2T0 B2T0 B2T1 B3T0
183
184//CHECK: omp target teams distribute parallel for
185//CHECK-SAME: block chunk size 5 thread chunk size 2
186
187//CHECK-NEXT: B0T0 B0T0 B0T1 B0T1 B0T2 B1T0 B1T0 B1T1 B1T1 B1T2
188//CHECK-SAME: B2T0 B2T0 B2T1 B2T1 B2T2 B3T0 B3T0 B3T1 B3T1 B3T2
189//CHECK-SAME: B4T0 B4T0 B4T1 B4T1 B4T2 B5T0 B5T0 B5T1 B5T1 B5T2
190//CHECK-SAME: B6T0 B6T0 B6T1 B6T1 B6T2 B7T0 B7T0 B7T1 B7T1 B7T2
191//CHECK-SAME: B8T0 B8T0 B8T1 B8T1 B8T2 B9T0 B9T0 B9T1 B9T1 B9T2
192//CHECK-SAME: B0T0 B0T0 B0T1 B0T1 B0T2 B1T0 B1T0 B1T1 B1T1 B1T2
193//CHECK-SAME: B2T0 B2T0 B2T1 B2T1 B2T2 B3T0 B3T0 B3T1 B3T1 B3T2
194//CHECK-SAME: B4T0 B4T0 B4T1 B4T1 B4T2 B5T0 B5T0 B5T1 B5T1 B5T2
195//CHECK-SAME: B6T0 B6T0 B6T1 B6T1 B6T2 B7T0 B7T0 B7T1 B7T1 B7T2
196//CHECK-SAME: B8T0 B8T0 B8T1 B8T1 B8T2 B9T0 B9T0 B9T1 B9T1 B9T2
197
198//CHECK: omp target teams distribute parallel for
199//CHECK-SAME: block chunk size 49 thread chunk size 2
200
201//CHECK-NEXT: B0T0 B0T0 B0T1 B0T1 B0T2 B0T2 B0T3 B0T3 B0T4 B0T4 B0T5 B0T5
202//CHECK-SAME: B0T6 B0T6 B0T7 B0T7 B0T8 B0T8 B0T9 B0T9 B0T10 B0T10 B0T11 B0T11
203//CHECK-SAME: B0T12 B0T12 B0T13 B0T13 B0T14 B0T14 B0T15 B0T15 B0T16 B0T16
204//CHECK-SAME: B0T17 B0T17 B0T18 B0T18 B0T19 B0T19 B0T20 B0T20 B0T21 B0T21
205//CHECK-SAME: B0T22 B0T22 B0T23 B0T23 B0T24
206//CHECK-SAME: B1T0 B1T0 B1T1 B1T1 B1T2 B1T2 B1T3 B1T3 B1T4 B1T4 B1T5 B1T5
207//CHECK-SAME: B1T6 B1T6 B1T7 B1T7 B1T8 B1T8 B1T9 B1T9 B1T10 B1T10 B1T11 B1T11
208//CHECK-SAME: B1T12 B1T12 B1T13 B1T13 B1T14 B1T14 B1T15 B1T15 B1T16 B1T16
209//CHECK-SAME: B1T17 B1T17 B1T18 B1T18 B1T19 B1T19 B1T20 B1T20 B1T21 B1T21
210//CHECK-SAME: B1T22 B1T22 B1T23 B1T23 B1T24
211//CHECK-SAME: B2T0 B2T0
212
213//CHECK: omp target teams distribute parallel for
214//CHECK-SAME: block chunk size 29 thread chunk size default
215
216//CHECK-NEXT: B0T0 B0T1 B0T2 B0T3 B0T4 B0T5 B0T6 B0T7 B0T8 B0T9
217//CHECK-SAME: B0T0 B0T1 B0T2 B0T3 B0T4 B0T5 B0T6 B0T7 B0T8 B0T9
218//CHECK-SAME: B0T0 B0T1 B0T2 B0T3 B0T4 B0T5 B0T6 B0T7 B0T8
219//CHECK-SAME: B1T0 B1T1 B1T2 B1T3 B1T4 B1T5 B1T6 B1T7 B1T8 B1T9
220//CHECK-SAME: B1T0 B1T1 B1T2 B1T3 B1T4 B1T5 B1T6 B1T7 B1T8 B1T9
221//CHECK-SAME: B1T0 B1T1 B1T2 B1T3 B1T4 B1T5 B1T6 B1T7 B1T8
222//CHECK-SAME: B2T0 B2T1 B2T2 B2T3 B2T4 B2T5 B2T6 B2T7 B2T8 B2T9
223//CHECK-SAME: B2T0 B2T1 B2T2 B2T3 B2T4 B2T5 B2T6 B2T7 B2T8 B2T9
224//CHECK-SAME: B2T0 B2T1 B2T2 B2T3 B2T4 B2T5 B2T6 B2T7 B2T8
225//CHECK-SAME: B3T0 B3T1 B3T2 B3T3 B3T4 B3T5 B3T6 B3T7 B3T8 B3T9
226//CHECK-SAME: B3T0 B3T1 B3T2
227
228//CHECK: omp target teams distribute parallel for
229//CHECK-SAME: block chunk size 101 thread chunk size default
230
231//CHECK-NEXT: B0T0 B0T1 B0T2 B0T3 B0T4 B0T5 B0T6 B0T7 B0T8 B0T9
232//CHECK-SAME: B0T0 B0T1 B0T2 B0T3 B0T4 B0T5 B0T6 B0T7 B0T8 B0T9
233//CHECK-SAME: B0T0 B0T1 B0T2 B0T3 B0T4 B0T5 B0T6 B0T7 B0T8 B0T9
234//CHECK-SAME: B0T0 B0T1 B0T2 B0T3 B0T4 B0T5 B0T6 B0T7 B0T8 B0T9
235//CHECK-SAME: B0T0 B0T1 B0T2 B0T3 B0T4 B0T5 B0T6 B0T7 B0T8 B0T9
236//CHECK-SAME: B0T0 B0T1 B0T2 B0T3 B0T4 B0T5 B0T6 B0T7 B0T8 B0T9
237//CHECK-SAME: B0T0 B0T1 B0T2 B0T3 B0T4 B0T5 B0T6 B0T7 B0T8 B0T9
238//CHECK-SAME: B0T0 B0T1 B0T2 B0T3 B0T4 B0T5 B0T6 B0T7 B0T8 B0T9
239//CHECK-SAME: B0T0 B0T1 B0T2 B0T3 B0T4 B0T5 B0T6 B0T7 B0T8 B0T9
240//CHECK-SAME: B0T0 B0T1 B0T2 B0T3 B0T4 B0T5 B0T6 B0T7 B0T8 B0T9
241
242//CHECK: omp target teams distribute parallel for
243//CHECK-SAME: default block chunk size thread chunk size 101
244
245//CHECK-NEXT: B0T0 B0T0 B0T0 B0T0 B0T0 B0T0 B0T0 B0T0 B0T0 B0T0
246//CHECK-SAME: B1T0 B1T0 B1T0 B1T0 B1T0 B1T0 B1T0 B1T0 B1T0 B1T0
247//CHECK-SAME: B2T0 B2T0 B2T0 B2T0 B2T0 B2T0 B2T0 B2T0 B2T0 B2T0
248//CHECK-SAME: B3T0 B3T0 B3T0 B3T0 B3T0 B3T0 B3T0 B3T0 B3T0 B3T0
249//CHECK-SAME: B4T0 B4T0 B4T0 B4T0 B4T0 B4T0 B4T0 B4T0 B4T0 B4T0
250//CHECK-SAME: B5T0 B5T0 B5T0 B5T0 B5T0 B5T0 B5T0 B5T0 B5T0 B5T0
251//CHECK-SAME: B6T0 B6T0 B6T0 B6T0 B6T0 B6T0 B6T0 B6T0 B6T0 B6T0
252//CHECK-SAME: B7T0 B7T0 B7T0 B7T0 B7T0 B7T0 B7T0 B7T0 B7T0 B7T0
253//CHECK-SAME: B8T0 B8T0 B8T0 B8T0 B8T0 B8T0 B8T0 B8T0 B8T0 B8T0
254//CHECK-SAME: B0T0 B0T0 B0T0 B0T0 B0T0 B0T0 B0T0 B0T0 B0T0 B0T0
255

source code of offload/test/offloading/workshare_chunk.c