xref: /aosp_15_r20/external/ComputeLibrary/cl_kernels/common/fft.clembed (revision c217d954acce2dbc11938adb493fc0abd69584f3)
1R"(
2
3#ifndef ARM_COMPUTE_HELPER_H
4#define ARM_COMPUTE_HELPER_H
5
6
7
8
9#define STORE_ROW_1(N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z) \
10    VSTORE(N0)                                                 \
11    (BASENAME##0, 0, (__global DATA_TYPE *)(PTR + 0 * STRIDE_Y + Z##0));
12
13#define STORE_ROW_2(N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z) \
14    STORE_ROW_1(N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z)     \
15    VSTORE(N0)                                                 \
16    (BASENAME##1, 0, (__global DATA_TYPE *)(PTR + 1 * STRIDE_Y + Z##1));
17
18#define STORE_ROW_3(N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z) \
19    STORE_ROW_2(N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z)     \
20    VSTORE(N0)                                                 \
21    (BASENAME##2, 0, (__global DATA_TYPE *)(PTR + 2 * STRIDE_Y + Z##2));
22
23#define STORE_ROW_4(N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z) \
24    STORE_ROW_3(N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z)     \
25    VSTORE(N0)                                                 \
26    (BASENAME##3, 0, (__global DATA_TYPE *)(PTR + 3 * STRIDE_Y + Z##3));
27
28#define STORE_ROW_5(N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z) \
29    STORE_ROW_4(N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z)     \
30    VSTORE(N0)                                                 \
31    (BASENAME##4, 0, (__global DATA_TYPE *)(PTR + 4 * STRIDE_Y + Z##4));
32
33#define STORE_ROW_6(N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z) \
34    STORE_ROW_5(N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z)     \
35    VSTORE(N0)                                                 \
36    (BASENAME##5, 0, (__global DATA_TYPE *)(PTR + 5 * STRIDE_Y + Z##5));
37
38#define STORE_ROW_7(N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z) \
39    STORE_ROW_6(N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z)     \
40    VSTORE(N0)                                                 \
41    (BASENAME##6, 0, (__global DATA_TYPE *)(PTR + 6 * STRIDE_Y + Z##6));
42
43#define STORE_ROW_8(N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z) \
44    STORE_ROW_7(N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z)     \
45    VSTORE(N0)                                                 \
46    (BASENAME##7, 0, (__global DATA_TYPE *)(PTR + 7 * STRIDE_Y + Z##7));
47
48#define STORE_ROW_9(N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z) \
49    STORE_ROW_8(N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z)     \
50    VSTORE(N0)                                                 \
51    (BASENAME##8, 0, (__global DATA_TYPE *)(PTR + 8 * STRIDE_Y + Z##8));
52
53#define STORE_ROW_10(N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z) \
54    STORE_ROW_9(N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z)      \
55    VSTORE(N0)                                                  \
56    (BASENAME##9, 0, (__global DATA_TYPE *)(PTR + 9 * STRIDE_Y + Z##9));
57
58#define STORE_ROW_11(N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z) \
59    STORE_ROW_10(N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z)     \
60    VSTORE(N0)                                                  \
61    (BASENAME##A, 0, (__global DATA_TYPE *)(PTR + 10 * STRIDE_Y + Z##A));
62
63#define STORE_ROW_12(N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z) \
64    STORE_ROW_11(N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z)     \
65    VSTORE(N0)                                                  \
66    (BASENAME##B, 0, (__global DATA_TYPE *)(PTR + 11 * STRIDE_Y + Z##B));
67
68#define STORE_ROW_13(N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z) \
69    STORE_ROW_12(N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z)     \
70    VSTORE(N0)                                                  \
71    (BASENAME##C, 0, (__global DATA_TYPE *)(PTR + 12 * STRIDE_Y + Z##C));
72
73#define STORE_ROW_14(N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z) \
74    STORE_ROW_13(N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z)     \
75    VSTORE(N0)                                                  \
76    (BASENAME##D, 0, (__global DATA_TYPE *)(PTR + 13 * STRIDE_Y + Z##D));
77
78#define STORE_ROW_15(N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z) \
79    STORE_ROW_14(N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z)     \
80    VSTORE(N0)                                                  \
81    (BASENAME##E, 0, (__global DATA_TYPE *)(PTR + 14 * STRIDE_Y + Z##E));
82
83#define STORE_ROW_16(N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z) \
84    STORE_ROW_15(N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z)     \
85    VSTORE(N0)                                                  \
86    (BASENAME##F, 0, (__global DATA_TYPE *)(PTR + 15 * STRIDE_Y + Z##F));
87
88
89
90#define CONVERT_STORE_ROW_1(N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z) \
91    VSTORE(N0)                                                         \
92    (CONVERT_SAT((BASENAME##0), VEC_DATA_TYPE(DATA_TYPE, N0)), 0, (__global DATA_TYPE *)(PTR + 0 * STRIDE_Y + Z##0));
93
94#define CONVERT_STORE_ROW_2(N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z) \
95    CONVERT_STORE_ROW_1(N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z)     \
96    VSTORE(N0)                                                         \
97    (CONVERT_SAT((BASENAME##1), VEC_DATA_TYPE(DATA_TYPE, N0)), 0, (__global DATA_TYPE *)(PTR + 1 * STRIDE_Y + Z##1));
98
99#define CONVERT_STORE_ROW_3(N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z) \
100    CONVERT_STORE_ROW_2(N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z)     \
101    VSTORE(N0)                                                         \
102    (CONVERT_SAT((BASENAME##2), VEC_DATA_TYPE(DATA_TYPE, N0)), 0, (__global DATA_TYPE *)(PTR + 2 * STRIDE_Y + Z##2));
103
104#define CONVERT_STORE_ROW_4(N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z) \
105    CONVERT_STORE_ROW_3(N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z)     \
106    VSTORE(N0)                                                         \
107    (CONVERT_SAT((BASENAME##3), VEC_DATA_TYPE(DATA_TYPE, N0)), 0, (__global DATA_TYPE *)(PTR + 3 * STRIDE_Y + Z##3));
108
109#define CONVERT_STORE_ROW_5(N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z) \
110    CONVERT_STORE_ROW_4(N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z)     \
111    VSTORE(N0)                                                         \
112    (CONVERT_SAT((BASENAME##4), VEC_DATA_TYPE(DATA_TYPE, N0)), 0, (__global DATA_TYPE *)(PTR + 4 * STRIDE_Y + Z##4));
113
114#define CONVERT_STORE_ROW_6(N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z) \
115    CONVERT_STORE_ROW_5(N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z)     \
116    VSTORE(N0)                                                         \
117    (CONVERT_SAT((BASENAME##5), VEC_DATA_TYPE(DATA_TYPE, N0)), 0, (__global DATA_TYPE *)(PTR + 5 * STRIDE_Y + Z##5));
118
119#define CONVERT_STORE_ROW_7(N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z) \
120    CONVERT_STORE_ROW_6(N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z)     \
121    VSTORE(N0)                                                         \
122    (CONVERT_SAT((BASENAME##6), VEC_DATA_TYPE(DATA_TYPE, N0)), 0, (__global DATA_TYPE *)(PTR + 6 * STRIDE_Y + Z##6));
123
124#define CONVERT_STORE_ROW_8(N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z) \
125    CONVERT_STORE_ROW_7(N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z)     \
126    VSTORE(N0)                                                         \
127    (CONVERT_SAT((BASENAME##7), VEC_DATA_TYPE(DATA_TYPE, N0)), 0, (__global DATA_TYPE *)(PTR + 7 * STRIDE_Y + Z##7));
128
129#define CONVERT_STORE_ROW_9(N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z) \
130    CONVERT_STORE_ROW_8(N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z)     \
131    VSTORE(N0)                                                         \
132    (CONVERT_SAT((BASENAME##8), VEC_DATA_TYPE(DATA_TYPE, N0)), 0, (__global DATA_TYPE *)(PTR + 8 * STRIDE_Y + Z##8));
133
134#define CONVERT_STORE_ROW_10(N0, DATA, BASENAME, PTR, STRIDE_Y, Z) \
135    CONVERT_STORE_ROW_9(N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z) \
136    VSTORE(N0)                                                     \
137    (CONVERT_SAT((BASENAME##9), VEC_DATA_TYPE(DATA_TYPE, N0)), 0, (__global DATA_TYPE *)(PTR + 9 * STRIDE_Y + Z##9));
138
139#define CONVERT_STORE_ROW_11(N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z) \
140    CONVERT_STORE_ROW_10(N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z)     \
141    VSTORE(N0)                                                          \
142    (CONVERT_SAT((BASENAME##A), VEC_DATA_TYPE(DATA_TYPE, N0)), 0, (__global DATA_TYPE *)(PTR + 10 * STRIDE_Y + Z##A));
143
144#define CONVERT_STORE_ROW_12(N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z) \
145    CONVERT_STORE_ROW_11(N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z)     \
146    VSTORE(N0)                                                          \
147    (CONVERT_SAT((BASENAME##B), VEC_DATA_TYPE(DATA_TYPE, N0)), 0, (__global DATA_TYPE *)(PTR + 11 * STRIDE_Y + Z##B));
148
149#define CONVERT_STORE_ROW_13(N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z) \
150    CONVERT_STORE_ROW_12(N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z)     \
151    VSTORE(N0)                                                          \
152    (CONVERT_SAT((BASENAME##C), VEC_DATA_TYPE(DATA_TYPE, N0)), 0, (__global DATA_TYPE *)(PTR + 12 * STRIDE_Y + Z##C));
153
154#define CONVERT_STORE_ROW_14(N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z) \
155    CONVERT_STORE_ROW_13(N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z)     \
156    VSTORE(N0)                                                          \
157    (CONVERT_SAT((BASENAME##D), VEC_DATA_TYPE(DATA_TYPE, N0)), 0, (__global DATA_TYPE *)(PTR + 13 * STRIDE_Y + Z##D));
158
159#define CONVERT_STORE_ROW_15(N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z) \
160    CONVERT_STORE_ROW_14(N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z)     \
161    VSTORE(N0)                                                          \
162    (CONVERT_SAT((BASENAME##E), VEC_DATA_TYPE(DATA_TYPE, N0)), 0, (__global DATA_TYPE *)(PTR + 14 * STRIDE_Y + Z##E));
163
164#define CONVERT_STORE_ROW_16(N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z) \
165    CONVERT_STORE_ROW_15(N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z)     \
166    VSTORE(N0)                                                          \
167    (CONVERT_SAT((BASENAME##F), VEC_DATA_TYPE(DATA_TYPE, N0)), 0, (__global DATA_TYPE *)(PTR + 15 * STRIDE_Y + Z##F));
168
169
170
171
172#define STORE_BLOCK_STR(M0, N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z) STORE_ROW_##M0(N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z)
173#define STORE_BLOCK(M0, N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z) STORE_BLOCK_STR(M0, N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z)
174
175
176
177#define CONVERT_STORE_BLOCK_STR(M0, N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z) CONVERT_STORE_ROW_##M0(N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z)
178#define CONVERT_STORE_BLOCK(M0, N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z) CONVERT_STORE_BLOCK_STR(M0, N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z)
179
180
181
182#define STORE_ROW_PARTIAL_1(N0, STORE_N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z) \
183    VSTORE_PARTIAL(N0, STORE_N0)                                                 \
184    (BASENAME##0, 0, (__global DATA_TYPE *)(PTR + 0 * STRIDE_Y + Z##0));
185
186#define STORE_ROW_PARTIAL_2(N0, STORE_N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z) \
187    STORE_ROW_PARTIAL_1(N0, STORE_N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z)     \
188    VSTORE_PARTIAL(N0, STORE_N0)                                                 \
189    (BASENAME##1, 0, (__global DATA_TYPE *)(PTR + 1 * STRIDE_Y + Z##1));
190
191#define STORE_ROW_PARTIAL_3(N0, STORE_N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z) \
192    STORE_ROW_PARTIAL_2(N0, STORE_N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z)     \
193    VSTORE_PARTIAL(N0, STORE_N0)                                                 \
194    (BASENAME##2, 0, (__global DATA_TYPE *)(PTR + 2 * STRIDE_Y + Z##2));
195
196#define STORE_ROW_PARTIAL_4(N0, STORE_N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z) \
197    STORE_ROW_PARTIAL_3(N0, STORE_N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z)     \
198    VSTORE_PARTIAL(N0, STORE_N0)                                                 \
199    (BASENAME##3, 0, (__global DATA_TYPE *)(PTR + 3 * STRIDE_Y + Z##3));
200
201#define STORE_ROW_PARTIAL_5(N0, STORE_N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z) \
202    STORE_ROW_PARTIAL_4(N0, STORE_N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z)     \
203    VSTORE_PARTIAL(N0, STORE_N0)                                                 \
204    (BASENAME##4, 0, (__global DATA_TYPE *)(PTR + 4 * STRIDE_Y + Z##4));
205
206#define STORE_ROW_PARTIAL_6(N0, STORE_N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z) \
207    STORE_ROW_PARTIAL_5(N0, STORE_N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z)     \
208    VSTORE_PARTIAL(N0, STORE_N0)                                                 \
209    (BASENAME##5, 0, (__global DATA_TYPE *)(PTR + 5 * STRIDE_Y + Z##5));
210
211#define STORE_ROW_PARTIAL_7(N0, STORE_N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z) \
212    STORE_ROW_PARTIAL_6(N0, STORE_N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z)     \
213    VSTORE_PARTIAL(N0, STORE_N0)                                                 \
214    (BASENAME##6, 0, (__global DATA_TYPE *)(PTR + 6 * STRIDE_Y + Z##6));
215
216#define STORE_ROW_PARTIAL_8(N0, STORE_N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z) \
217    STORE_ROW_PARTIAL_7(N0, STORE_N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z)     \
218    VSTORE_PARTIAL(N0, STORE_N0)                                                 \
219    (BASENAME##7, 0, (__global DATA_TYPE *)(PTR + 7 * STRIDE_Y + Z##7));
220
221#define STORE_ROW_PARTIAL_9(N0, STORE_N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z) \
222    STORE_ROW_PARTIAL_8(N0, STORE_N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z)     \
223    VSTORE_PARTIAL(N0, STORE_N0)                                                 \
224    (BASENAME##8, 0, (__global DATA_TYPE *)(PTR + 8 * STRIDE_Y + Z##8));
225
226#define STORE_ROW_PARTIAL_10(N0, STORE_N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z) \
227    STORE_ROW_PARTIAL_9(N0, STORE_N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z)      \
228    VSTORE_PARTIAL(N0, STORE_N0)                                                  \
229    (BASENAME##9, 0, (__global DATA_TYPE *)(PTR + 9 * STRIDE_Y + Z##9));
230
231#define STORE_ROW_PARTIAL_11(N0, STORE_N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z) \
232    STORE_ROW_PARTIAL_10(N0, STORE_N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z)     \
233    VSTORE_PARTIAL(N0, STORE_N0)                                                  \
234    (BASENAME##A, 0, (__global DATA_TYPE *)(PTR + 10 * STRIDE_Y + Z##A));
235
236#define STORE_ROW_PARTIAL_12(N0, STORE_N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z) \
237    STORE_ROW_PARTIAL_11(N0, STORE_N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z)     \
238    VSTORE_PARTIAL(N0, STORE_N0)                                                  \
239    (BASENAME##B, 0, (__global DATA_TYPE *)(PTR + 11 * STRIDE_Y + Z##B));
240
241#define STORE_ROW_PARTIAL_13(N0, STORE_N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z) \
242    STORE_ROW_PARTIAL_12(N0, STORE_N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z)     \
243    VSTORE_PARTIAL(N0, STORE_N0)                                                  \
244    (BASENAME##C, 0, (__global DATA_TYPE *)(PTR + 12 * STRIDE_Y + Z##C));
245
246#define STORE_ROW_PARTIAL_14(N0, STORE_N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z) \
247    STORE_ROW_PARTIAL_13(N0, STORE_N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z)     \
248    VSTORE_PARTIAL(N0, STORE_N0)                                                  \
249    (BASENAME##D, 0, (__global DATA_TYPE *)(PTR + 13 * STRIDE_Y + Z##D));
250
251#define STORE_ROW_PARTIAL_15(N0, STORE_N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z) \
252    STORE_ROW_PARTIAL_14(N0, STORE_N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z)     \
253    VSTORE_PARTIAL(N0, STORE_N0)                                                  \
254    (BASENAME##E, 0, (__global DATA_TYPE *)(PTR + 14 * STRIDE_Y + Z##E));
255
256#define STORE_ROW_PARTIAL_16(N0, STORE_N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z) \
257    STORE_ROW_PARTIAL_15(N0, STORE_N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z)     \
258    VSTORE_PARTIAL(N0, STORE_N0)                                                  \
259    (BASENAME##F, 0, (__global DATA_TYPE *)(PTR + 15 * STRIDE_Y + Z##F));
260
261
262
263#define STORE_BLOCK_PARTIAL_STR(STORE_M0, STORE_N0, N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z) STORE_ROW_PARTIAL_##STORE_M0(N0, STORE_N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z)
264#define STORE_BLOCK_PARTIAL(STORE_M0, STORE_N0, N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z) STORE_BLOCK_PARTIAL_STR(STORE_M0, STORE_N0, N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z)
265
266#define STORE_BLOCK_PARTIAL_IN_X_AND_Y(M0, N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z, PARTIAL_STORE_M0, PARTIAL_STORE_N0, PARTIAL_COND_Y, PARTIAL_COND_X) \
267    if(!(PARTIAL_COND_X) && !(PARTIAL_COND_Y))                                                                                                            \
268    {                                                                                                                                                     \
269        STORE_BLOCK_PARTIAL(M0, N0, N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z);                                                                           \
270    }                                                                                                                                                     \
271    else if((PARTIAL_COND_Y) && !(PARTIAL_COND_X))                                                                                                        \
272    {                                                                                                                                                     \
273        STORE_BLOCK_PARTIAL(PARTIAL_STORE_M0, N0, N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z);                                                             \
274    }                                                                                                                                                     \
275    else if(!(PARTIAL_COND_Y) && (PARTIAL_COND_X))                                                                                                        \
276    {                                                                                                                                                     \
277        STORE_BLOCK_PARTIAL(M0, PARTIAL_STORE_N0, N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z);                                                             \
278    }                                                                                                                                                     \
279    else                                                                                                                                                  \
280    {                                                                                                                                                     \
281        STORE_BLOCK_PARTIAL(PARTIAL_STORE_M0, PARTIAL_STORE_N0, N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z);                                               \
282    }
283
284#define STORE_BLOCK_PARTIAL_IN_X(M0, N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z, PARTIAL_STORE_N0, PARTIAL_COND_X) \
285    if(!(PARTIAL_COND_X))                                                                                         \
286    {                                                                                                             \
287        STORE_BLOCK_PARTIAL(M0, N0, N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z);                                   \
288    }                                                                                                             \
289    else                                                                                                          \
290    {                                                                                                             \
291        STORE_BLOCK_PARTIAL(M0, PARTIAL_STORE_N0, N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z);                     \
292    }
293
294#define STORE_BLOCK_PARTIAL_IN_Y(M0, N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z, PARTIAL_STORE_M0, PARTIAL_COND_Y) \
295    if(!(PARTIAL_COND_Y))                                                                                         \
296    {                                                                                                             \
297        STORE_BLOCK_PARTIAL(M0, N0, N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z);                                   \
298    }                                                                                                             \
299    else                                                                                                          \
300    {                                                                                                             \
301        STORE_BLOCK_PARTIAL(PARTIAL_STORE_M0, N0, N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z);                     \
302    }
303
304
305#if defined(PARTIAL_STORE_M0) && defined(PARTIAL_STORE_N0)
306
307
308#if PARTIAL_STORE_M0 == 0 && PARTIAL_STORE_N0 == 0
309
310#define STORE_BLOCK_BOUNDARY_AWARE(M0, N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z, PARTIAL_STORE_M0, PARTIAL_STORE_N0, PARTIAL_COND_Y, PARTIAL_COND_X) \
311    STORE_BLOCK(M0, N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z)
312
313#elif PARTIAL_STORE_M0 > 0 && PARTIAL_STORE_N0 == 0
314
315#define STORE_BLOCK_BOUNDARY_AWARE(M0, N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z, PARTIAL_STORE_M0, PARTIAL_STORE_N0, PARTIAL_COND_Y, PARTIAL_COND_X) \
316    STORE_BLOCK_PARTIAL_IN_Y(M0, N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z, PARTIAL_STORE_M0, PARTIAL_COND_Y)
317
318#elif PARTIAL_STORE_M0 == 0 && PARTIAL_STORE_N0 > 0
319
320#define STORE_BLOCK_BOUNDARY_AWARE(M0, N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z, PARTIAL_STORE_M0, PARTIAL_STORE_N0, PARTIAL_COND_Y, PARTIAL_COND_X) \
321    STORE_BLOCK_PARTIAL_IN_X(M0, N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z, PARTIAL_STORE_N0, PARTIAL_COND_X)
322
323#else
324
325#define STORE_BLOCK_BOUNDARY_AWARE(M0, N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z, PARTIAL_STORE_M0, PARTIAL_STORE_N0, PARTIAL_COND_Y, PARTIAL_COND_X) \
326    STORE_BLOCK_PARTIAL_IN_X_AND_Y(M0, N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z, PARTIAL_STORE_M0, PARTIAL_STORE_N0, PARTIAL_COND_Y, PARTIAL_COND_X)
327
328#endif
329
330#endif
331
332
333#if defined(PARTIAL_STORE_M0)
334
335#define COMPUTE_M0_START_ROW(y, M0, PARTIAL_STORE_M0) \
336    ((uint)(max(0, (int)(y * M0) - (int)((M0 - PARTIAL_STORE_M0) % M0))))
337#else
338#define COMPUTE_M0_START_ROW(y, M0, PARTIAL_STORE_M0) \
339    ((uint)(y * M0))
340#endif
341
342
343
344#define STORE_VECTOR_SELECT(basename, data_type, ptr, vec_size, leftover, cond) \
345    STORE_BLOCK_PARTIAL_IN_X(1, vec_size, data_type, basename, ptr, 0, 0, leftover, cond)
346
347
348#if defined(ARM_COMPUTE_OPENCL_FP16_ENABLED) && defined(cl_khr_fp16)
349#pragma OPENCL EXTENSION cl_khr_fp16 : enable
350#endif
351
352#if defined(ARM_COMPUTE_OPENCL_DOT8_ENABLED) && defined(cl_arm_integer_dot_product_int8)
353#pragma OPENCL EXTENSION cl_arm_integer_dot_product_int8 : enable
354#endif
355
356#if defined(ARM_COMPUTE_OPENCL_DOT8_ACC_ENABLED) && defined(cl_arm_integer_dot_product_accumulate_int8)
357#pragma OPENCL EXTENSION cl_arm_integer_dot_product_accumulate_int8 : enable
358#endif
359
360#if defined(ARM_COMPUTE_DEBUG_ENABLED) && defined(cl_arm_printf)
361#pragma OPENCL EXTENSION cl_arm_printf : enable
362#endif
363
364#define GPU_ARCH_MIDGARD 0x100
365#define GPU_ARCH_BIFROST 0x200
366#define GPU_ARCH_VALHALL 0x300
367
368
369#define CONCAT(a, b) a##b
370
371
372#define EXPAND(x) x
373
374
375#define CLAMP(x, min_val, max_val) min(max(x, min_val), max_val)
376
377
378#define REV1(x) ((x))
379#define REV2(x) ((x).s10)
380#define REV3(x) ((x).s210)
381#define REV4(x) ((x).s3210)
382#define REV8(x) ((x).s76543210)
383#define REV16(x) ((x).sFEDCBA9876543210)
384
385
386
387#define REVERSE_STR(x, s) REV##s((x))
388#define REVERSE(x, s) REVERSE_STR(x, s)
389
390
391
392#define ROT1_0(x) ((x))
393#define ROT1_1(x) ((x))
394
395#define ROT2_0(x) ((x))
396#define ROT2_1(x) ((x).s10)
397#define ROT2_2(x) ((x))
398
399#define ROT3_0(x) ((x))
400#define ROT3_1(x) ((x).s201)
401#define ROT3_2(x) ((x).s120)
402#define ROT3_3(x) ((x))
403
404#define ROT4_0(x) ((x))
405#define ROT4_1(x) ((x).s3012)
406#define ROT4_2(x) ((x).s2301)
407#define ROT4_3(x) ((x).s1230)
408#define ROT4_4(x) ((x))
409
410#define ROT8_0(x) ((x))
411#define ROT8_1(x) ((x).s70123456)
412#define ROT8_2(x) ((x).s67012345)
413#define ROT8_3(x) ((x).s56701234)
414#define ROT8_4(x) ((x).s45670123)
415#define ROT8_5(x) ((x).s34567012)
416#define ROT8_6(x) ((x).s23456701)
417#define ROT8_7(x) ((x).s12345670)
418#define ROT8_8(x) ((x))
419
420#define ROT16_0(x) ((x))
421#define ROT16_1(x) ((x).sF0123456789ABCDE)
422#define ROT16_2(x) ((x).sEF0123456789ABCD)
423#define ROT16_3(x) ((x).sDEF0123456789ABC)
424#define ROT16_4(x) ((x).sCDEF0123456789AB)
425#define ROT16_5(x) ((x).sBCDEF0123456789A)
426#define ROT16_6(x) ((x).sABCDEF0123456789)
427#define ROT16_7(x) ((x).s9ABCDEF012345678)
428#define ROT16_8(x) ((x).s89ABCDEF01234567)
429#define ROT16_9(x) ((x).s789ABCDEF0123456)
430#define ROT16_10(x) ((x).s6789ABCDEF012345)
431#define ROT16_11(x) ((x).s56789ABCDEF01234)
432#define ROT16_12(x) ((x).s456789ABCDEF0123)
433#define ROT16_13(x) ((x).s3456789ABCDEF012)
434#define ROT16_14(x) ((x).s23456789ABCDEF01)
435#define ROT16_15(x) ((x).s123456789ABCDEF0)
436#define ROT16_16(x) ((x))
437
438
439
440#define ROTATE_STR(x, s, n) ROT##s##_##n(x)
441#define ROTATE(x, s, n) ROTATE_STR(x, s, n)
442
443
444
445#define V_OFFS1(dt) (dt##1)(0)
446#define V_OFFS2(dt) (dt##2)(0, 1)
447#define V_OFFS3(dt) (dt##3)(0, 1, 2)
448#define V_OFFS4(dt) (dt##4)(0, 1, 2, 3)
449#define V_OFFS8(dt) (dt##8)(0, 1, 2, 3, 4, 5, 6, 7)
450#define V_OFFS16(dt) (dt##16)(0, 1, 2, 3, 4, 5, 6, 7, 8, 9, 10, 11, 12, 13, 14, 15)
451
452
453
454#define VEC_OFFS_STR(dt, s) V_OFFS##s(dt)
455#define VEC_OFFS(dt, s) VEC_OFFS_STR(dt, s)
456
457
458#define VLOAD_STR(size) vload##size
459#define VLOAD(size) VLOAD_STR(size)
460
461
462#define VLOAD_PARTIAL_STR(size, load_size) vload_partial_##size##_##load_size
463#define VLOAD_PARTIAL(size, load_size) VLOAD_PARTIAL_STR(size, load_size)
464
465#define NO_LOAD(data, offs, ptr) \
466    {                            \
467    }
468
469
470#define vload_partial_1_0 NO_LOAD
471#define vload_partial_1_1 vload1
472#define vload_partial_1_2 NO_LOAD
473#define vload_partial_1_3 NO_LOAD
474#define vload_partial_1_4 NO_LOAD
475#define vload_partial_1_5 NO_LOAD
476#define vload_partial_1_6 NO_LOAD
477#define vload_partial_1_7 NO_LOAD
478#define vload_partial_1_8 NO_LOAD
479#define vload_partial_1_9 NO_LOAD
480#define vload_partial_1_10 NO_LOAD
481#define vload_partial_1_11 NO_LOAD
482#define vload_partial_1_12 NO_LOAD
483#define vload_partial_1_13 NO_LOAD
484#define vload_partial_1_14 NO_LOAD
485#define vload_partial_1_15 NO_LOAD
486#define vload_partial_1_16 NO_LOAD
487
488#define vload_partial_2_0 NO_LOAD
489#define vload_partial_2_1 vload_partial_1
490#define vload_partial_2_2 vload_partial_2
491#define vload_partial_2_3 NO_LOAD
492#define vload_partial_2_4 NO_LOAD
493#define vload_partial_2_5 NO_LOAD
494#define vload_partial_2_6 NO_LOAD
495#define vload_partial_2_7 NO_LOAD
496#define vload_partial_2_8 NO_LOAD
497#define vload_partial_2_9 NO_LOAD
498#define vload_partial_2_10 NO_LOAD
499#define vload_partial_2_11 NO_LOAD
500#define vload_partial_2_12 NO_LOAD
501#define vload_partial_2_13 NO_LOAD
502#define vload_partial_2_14 NO_LOAD
503#define vload_partial_2_15 NO_LOAD
504#define vload_partial_2_16 NO_LOAD
505
506#define vload_partial_3_0 NO_LOAD
507#define vload_partial_3_1 vload_partial_1
508#define vload_partial_3_2 vload_partial_2
509#define vload_partial_3_3 vload_partial_3
510#define vload_partial_3_4 NO_LOAD
511#define vload_partial_3_5 NO_LOAD
512#define vload_partial_3_6 NO_LOAD
513#define vload_partial_3_7 NO_LOAD
514#define vload_partial_3_8 NO_LOAD
515#define vload_partial_3_9 NO_LOAD
516#define vload_partial_3_10 NO_LOAD
517#define vload_partial_3_11 NO_LOAD
518#define vload_partial_3_12 NO_LOAD
519#define vload_partial_3_13 NO_LOAD
520#define vload_partial_3_14 NO_LOAD
521#define vload_partial_3_15 NO_LOAD
522#define vload_partial_3_16 NO_LOAD
523
524#define vload_partial_4_0 NO_LOAD
525#define vload_partial_4_1 vload_partial_1
526#define vload_partial_4_2 vload_partial_2
527#define vload_partial_4_3 vload_partial_3
528#define vload_partial_4_4 vload_partial_4
529#define vload_partial_4_5 NO_LOAD
530#define vload_partial_4_6 NO_LOAD
531#define vload_partial_4_7 NO_LOAD
532#define vload_partial_4_8 NO_LOAD
533#define vload_partial_4_9 NO_LOAD
534#define vload_partial_4_10 NO_LOAD
535#define vload_partial_4_11 NO_LOAD
536#define vload_partial_4_12 NO_LOAD
537#define vload_partial_4_13 NO_LOAD
538#define vload_partial_4_14 NO_LOAD
539#define vload_partial_4_15 NO_LOAD
540#define vload_partial_4_16 NO_LOAD
541
542#define vload_partial_8_0 NO_LOAD
543#define vload_partial_8_1 vload_partial_1
544#define vload_partial_8_2 vload_partial_2
545#define vload_partial_8_3 vload_partial_3
546#define vload_partial_8_4 vload_partial_4
547#define vload_partial_8_5 vload_partial_5
548#define vload_partial_8_6 vload_partial_6
549#define vload_partial_8_7 vload_partial_7
550#define vload_partial_8_8 vload_partial_8
551#define vload_partial_8_9 NO_LOAD
552#define vload_partial_8_10 NO_LOAD
553#define vload_partial_8_11 NO_LOAD
554#define vload_partial_8_12 NO_LOAD
555#define vload_partial_8_13 NO_LOAD
556#define vload_partial_8_14 NO_LOAD
557#define vload_partial_8_15 NO_LOAD
558#define vload_partial_8_16 NO_LOAD
559
560#define vload_partial_16_0 NO_LOAD
561#define vload_partial_16_1 vload_partial_1
562#define vload_partial_16_2 vload_partial_2
563#define vload_partial_16_3 vload_partial_3
564#define vload_partial_16_4 vload_partial_4
565#define vload_partial_16_5 vload_partial_5
566#define vload_partial_16_6 vload_partial_6
567#define vload_partial_16_7 vload_partial_7
568#define vload_partial_16_8 vload_partial_8
569#define vload_partial_16_9 vload_partial_9
570#define vload_partial_16_10 vload_partial_10
571#define vload_partial_16_11 vload_partial_11
572#define vload_partial_16_12 vload_partial_12
573#define vload_partial_16_13 vload_partial_13
574#define vload_partial_16_14 vload_partial_14
575#define vload_partial_16_15 vload_partial_15
576#define vload_partial_16_16 vload_partial_16
577
578
579#define vload_partial_1(DATA, OFFSET, PTR) \
580    DATA.s0 = vload1(OFFSET, PTR);
581
582#define vload_partial_2(DATA, OFFSET, PTR) \
583    DATA.s01 = vload2(OFFSET, PTR);
584
585#define vload_partial_3(DATA, OFFSET, PTR) \
586    DATA.s012 = vload3(OFFSET, PTR);
587
588#define vload_partial_4(DATA, OFFSET, PTR) \
589    DATA.s0123 = vload4(OFFSET, PTR);
590
591#define vload_partial_5(DATA, OFFSET, PTR)    \
592    vload_partial_4(DATA.s0123, OFFSET, PTR); \
593    DATA.s4 = vload1(OFFSET, PTR + 4);
594
595#define vload_partial_6(DATA, OFFSET, PTR)    \
596    vload_partial_4(DATA.s0123, OFFSET, PTR); \
597    vload_partial_2(DATA.s45, OFFSET, PTR + 4);
598
599#define vload_partial_7(DATA, OFFSET, PTR)    \
600    vload_partial_4(DATA.s0123, OFFSET, PTR); \
601    vload_partial_3(DATA.s456, OFFSET, PTR + 4);
602
603#define vload_partial_8(DATA, OFFSET, PTR) \
604    DATA.s01234567 = vload8(OFFSET, PTR);
605
606#define vload_partial_9(DATA, OFFSET, PTR)        \
607    vload_partial_8(DATA.s01234567, OFFSET, PTR); \
608    DATA.s8 = vload1(OFFSET, PTR + 8);
609
610#define vload_partial_10(DATA, OFFSET, PTR)       \
611    vload_partial_8(DATA.s01234567, OFFSET, PTR); \
612    vload_partial_2(DATA.s89, OFFSET, PTR + 8);
613
614#define vload_partial_11(DATA, OFFSET, PTR)       \
615    vload_partial_8(DATA.s01234567, OFFSET, PTR); \
616    vload_partial_3(DATA.s89A, OFFSET, PTR + 8);
617
618#define vload_partial_12(DATA, OFFSET, PTR)       \
619    vload_partial_8(DATA.s01234567, OFFSET, PTR); \
620    vload_partial_4(DATA.s89AB, OFFSET, PTR + 8);
621
622#define vload_partial_13(DATA, OFFSET, PTR)       \
623    vload_partial_8(DATA.s01234567, OFFSET, PTR); \
624    vload_partial_5(DATA.s89ABCDEF, OFFSET, PTR + 8);
625
626#define vload_partial_14(DATA, OFFSET, PTR)       \
627    vload_partial_8(DATA.s01234567, OFFSET, PTR); \
628    vload_partial_6(DATA.s89ABCDEF, OFFSET, PTR + 8);
629
630#define vload_partial_15(DATA, OFFSET, PTR)       \
631    vload_partial_8(DATA.s01234567, OFFSET, PTR); \
632    vload_partial_7(DATA.s89ABCDEF, OFFSET, PTR + 8);
633
634#define vload_partial_16(DATA, OFFSET, PTR) \
635    DATA = vload16(OFFSET, PTR);
636
637
638
639#define PIXEL_UNIT4 1
640#define PIXEL_UNIT8 2
641#define PIXEL_UNIT16 4
642
643
644#define CONVERT_VECTOR_SIZE_TO_PIXEL_UNIT_STR(vec_size) PIXEL_UNIT##vec_size
645#define CONVERT_VECTOR_SIZE_TO_PIXEL_UNIT(vec_size) CONVERT_VECTOR_SIZE_TO_PIXEL_UNIT_STR(vec_size)
646
647
648#define read_image2d_floatx1(img, x_coord, y_coord) (float4)(read_imagef(img, (int2)(x_coord, y_coord)));
649#define read_image2d_floatx2(img, x_coord, y_coord) (float8)(read_imagef(img, (int2)(x_coord, y_coord)), read_imagef(img, (int2)(x_coord + 1, y_coord)));
650#define read_image2d_floatx4(img, x_coord, y_coord) (float16)(read_imagef(img, (int2)(x_coord, y_coord)), read_imagef(img, (int2)(x_coord + 1, y_coord)), read_imagef(img, (int2)(x_coord + 2, y_coord)), read_imagef(img, (int2)(x_coord + 3, y_coord)));
651
652#if defined(ARM_COMPUTE_OPENCL_FP16_ENABLED) && defined(cl_khr_fp16)
653#define read_image2d_halfx1(img, x_coord, y_coord) (half4)(read_imageh(img, (int2)(x_coord, y_coord)));
654#define read_image2d_halfx2(img, x_coord, y_coord) (half8)(read_imageh(img, (int2)(x_coord, y_coord)), read_imageh(img, (int2)(x_coord + 1, y_coord)));
655#define read_image2d_halfx4(img, x_coord, y_coord) (half16)(read_imageh(img, (int2)(x_coord, y_coord)), read_imageh(img, (int2)(x_coord + 1, y_coord)), read_imageh(img, (int2)(x_coord + 2, y_coord)), read_imageh(img, (int2)(x_coord + 3, y_coord)));
656#endif
657
658#define write_image2d_floatx1(img, x_coord, y_coord, values) (write_imagef(img, (int2)(x_coord, y_coord), values));
659#define write_image2d_floatx2(img, x_coord, y_coord, values) (write_imagef(img, (int2)(x_coord, y_coord), values.s0123), write_imagef(img, (int2)(x_coord + 1, y_coord), values.s4567));
660#define write_image2d_floatx4(img, x_coord, y_coord, values) (write_imagef(img, (int2)(x_coord, y_coord), values.s0123), write_imagef(img, (int2)(x_coord + 1, y_coord), values.s4567), write_imagef(img, (int2)(x_coord + 2, y_coord), values.s89AB), write_imagef(img, (int2)(x_coord + 3, y_coord), values.sCDEF));
661
662#if defined(ARM_COMPUTE_OPENCL_FP16_ENABLED) && defined(cl_khr_fp16)
663#define write_image2d_halfx1(img, x_coord, y_coord, values) (write_imageh(img, (int2)(x_coord, y_coord), values));
664#define write_image2d_halfx2(img, x_coord, y_coord, values) (write_imageh(img, (int2)(x_coord, y_coord), values.s0123), write_imageh(img, (int2)(x_coord + 1, y_coord), values.s4567));
665#define write_image2d_halfx4(img, x_coord, y_coord, values) (write_imageh(img, (int2)(x_coord, y_coord), values.s0123), write_imageh(img, (int2)(x_coord + 1, y_coord), values.s4567), write_imageh(img, (int2)(x_coord + 2, y_coord), values.s89AB), write_imageh(img, (int2)(x_coord + 3, y_coord), values.sCDEF));
666#endif
667
668
669#define READ_IMAGE2D_STR(data_type, n0, img, x_coord, y_coord) read_image2d_##data_type##x##n0(img, x_coord, y_coord)
670#define READ_IMAGE2D(data_type, n0, img, x_coord, y_coord) READ_IMAGE2D_STR(data_type, n0, img, x_coord, y_coord)
671
672
673#define WRITE_IMAGE2D_STR(data_type, n0, img, x_coord, y_coord, values) write_image2d_##data_type##x##n0(img, x_coord, y_coord, values)
674#define WRITE_IMAGE2D(data_type, n0, img, x_coord, y_coord, values) WRITE_IMAGE2D_STR(data_type, n0, img, x_coord, y_coord, values)
675
676#define VSTORE_STR(size) vstore##size
677#define VSTORE(size) VSTORE_STR(size)
678
679#define float1 float
680#define half1 half
681#define char1 char
682#define uchar1 uchar
683#define short1 short
684#define ushort1 ushort
685#define int1 int
686#define uint1 uint
687#define long1 long
688#define ulong1 ulong
689#define double1 double
690
691#define vload1(OFFSET, PTR) *(OFFSET + PTR)
692#define vstore1(DATA, OFFSET, PTR) *(OFFSET + PTR) = DATA
693
694
695#define VSTORE_PARTIAL_STR(size, store_size) vstore_partial_##size##_##store_size
696#define VSTORE_PARTIAL(size, store_size) VSTORE_PARTIAL_STR(size, store_size)
697
698#define NO_STORE(data, offs, ptr) \
699    {                             \
700    }
701
702
703#define vstore_partial_1_0 NO_STORE
704#define vstore_partial_1_1 vstore1
705#define vstore_partial_1_2 NO_STORE
706#define vstore_partial_1_3 NO_STORE
707#define vstore_partial_1_4 NO_STORE
708#define vstore_partial_1_5 NO_STORE
709#define vstore_partial_1_6 NO_STORE
710#define vstore_partial_1_7 NO_STORE
711#define vstore_partial_1_8 NO_STORE
712#define vstore_partial_1_9 NO_STORE
713#define vstore_partial_1_10 NO_STORE
714#define vstore_partial_1_11 NO_STORE
715#define vstore_partial_1_12 NO_STORE
716#define vstore_partial_1_13 NO_STORE
717#define vstore_partial_1_14 NO_STORE
718#define vstore_partial_1_15 NO_STORE
719#define vstore_partial_1_16 NO_STORE
720
721#define vstore_partial_2_0 NO_STORE
722#define vstore_partial_2_1 vstore_partial_1
723#define vstore_partial_2_2 vstore_partial_2
724#define vstore_partial_2_3 NO_STORE
725#define vstore_partial_2_4 NO_STORE
726#define vstore_partial_2_5 NO_STORE
727#define vstore_partial_2_6 NO_STORE
728#define vstore_partial_2_7 NO_STORE
729#define vstore_partial_2_8 NO_STORE
730#define vstore_partial_2_9 NO_STORE
731#define vstore_partial_2_10 NO_STORE
732#define vstore_partial_2_11 NO_STORE
733#define vstore_partial_2_12 NO_STORE
734#define vstore_partial_2_13 NO_STORE
735#define vstore_partial_2_14 NO_STORE
736#define vstore_partial_2_15 NO_STORE
737#define vstore_partial_2_16 NO_STORE
738
739#define vstore_partial_3_0 NO_STORE
740#define vstore_partial_3_1 vstore_partial_1
741#define vstore_partial_3_2 vstore_partial_2
742#define vstore_partial_3_3 vstore_partial_3
743#define vstore_partial_3_4 NO_STORE
744#define vstore_partial_3_5 NO_STORE
745#define vstore_partial_3_6 NO_STORE
746#define vstore_partial_3_7 NO_STORE
747#define vstore_partial_3_8 NO_STORE
748#define vstore_partial_3_9 NO_STORE
749#define vstore_partial_3_10 NO_STORE
750#define vstore_partial_3_11 NO_STORE
751#define vstore_partial_3_12 NO_STORE
752#define vstore_partial_3_13 NO_STORE
753#define vstore_partial_3_14 NO_STORE
754#define vstore_partial_3_15 NO_STORE
755#define vstore_partial_3_16 NO_STORE
756
757#define vstore_partial_4_0 NO_STORE
758#define vstore_partial_4_1 vstore_partial_1
759#define vstore_partial_4_2 vstore_partial_2
760#define vstore_partial_4_3 vstore_partial_3
761#define vstore_partial_4_4 vstore_partial_4
762#define vstore_partial_4_5 NO_STORE
763#define vstore_partial_4_6 NO_STORE
764#define vstore_partial_4_7 NO_STORE
765#define vstore_partial_4_8 NO_STORE
766#define vstore_partial_4_9 NO_STORE
767#define vstore_partial_4_10 NO_STORE
768#define vstore_partial_4_11 NO_STORE
769#define vstore_partial_4_12 NO_STORE
770#define vstore_partial_4_13 NO_STORE
771#define vstore_partial_4_14 NO_STORE
772#define vstore_partial_4_15 NO_STORE
773#define vstore_partial_4_16 NO_STORE
774
775#define vstore_partial_8_0 NO_STORE
776#define vstore_partial_8_1 vstore_partial_1
777#define vstore_partial_8_2 vstore_partial_2
778#define vstore_partial_8_3 vstore_partial_3
779#define vstore_partial_8_4 vstore_partial_4
780#define vstore_partial_8_5 vstore_partial_5
781#define vstore_partial_8_6 vstore_partial_6
782#define vstore_partial_8_7 vstore_partial_7
783#define vstore_partial_8_8 vstore_partial_8
784#define vstore_partial_8_9 NO_STORE
785#define vstore_partial_8_10 NO_STORE
786#define vstore_partial_8_11 NO_STORE
787#define vstore_partial_8_12 NO_STORE
788#define vstore_partial_8_13 NO_STORE
789#define vstore_partial_8_14 NO_STORE
790#define vstore_partial_8_15 NO_STORE
791#define vstore_partial_8_16 NO_STORE
792
793#define vstore_partial_16_0 NO_STORE
794#define vstore_partial_16_1 vstore_partial_1
795#define vstore_partial_16_2 vstore_partial_2
796#define vstore_partial_16_3 vstore_partial_3
797#define vstore_partial_16_4 vstore_partial_4
798#define vstore_partial_16_5 vstore_partial_5
799#define vstore_partial_16_6 vstore_partial_6
800#define vstore_partial_16_7 vstore_partial_7
801#define vstore_partial_16_8 vstore_partial_8
802#define vstore_partial_16_9 vstore_partial_9
803#define vstore_partial_16_10 vstore_partial_10
804#define vstore_partial_16_11 vstore_partial_11
805#define vstore_partial_16_12 vstore_partial_12
806#define vstore_partial_16_13 vstore_partial_13
807#define vstore_partial_16_14 vstore_partial_14
808#define vstore_partial_16_15 vstore_partial_15
809#define vstore_partial_16_16 vstore_partial_16
810
811
812#define vstore_partial_1(DATA, OFFSET, PTR) \
813    vstore1(DATA.s0, OFFSET, PTR);
814
815#define vstore_partial_2(DATA, OFFSET, PTR) \
816    vstore2(DATA.s01, OFFSET, PTR);
817
818#define vstore_partial_3(DATA, OFFSET, PTR) \
819    vstore3(DATA.s012, OFFSET, PTR);
820
821#define vstore_partial_4(DATA, OFFSET, PTR) \
822    vstore4(DATA.s0123, OFFSET, PTR);
823
824#define vstore_partial_5(DATA, OFFSET, PTR)    \
825    vstore_partial_4(DATA.s0123, OFFSET, PTR); \
826    vstore1(DATA.s4, OFFSET, PTR + 4);
827
828#define vstore_partial_6(DATA, OFFSET, PTR)    \
829    vstore_partial_4(DATA.s0123, OFFSET, PTR); \
830    vstore_partial_2(DATA.s45, OFFSET, PTR + 4);
831
832#define vstore_partial_7(DATA, OFFSET, PTR)    \
833    vstore_partial_4(DATA.s0123, OFFSET, PTR); \
834    vstore_partial_3(DATA.s456, OFFSET, PTR + 4);
835
836#define vstore_partial_8(DATA, OFFSET, PTR) \
837    vstore8(DATA.s01234567, OFFSET, PTR);
838
839#define vstore_partial_9(DATA, OFFSET, PTR)        \
840    vstore_partial_8(DATA.s01234567, OFFSET, PTR); \
841    vstore1(DATA.s8, OFFSET, PTR + 8);
842
843#define vstore_partial_10(DATA, OFFSET, PTR)       \
844    vstore_partial_8(DATA.s01234567, OFFSET, PTR); \
845    vstore_partial_2(DATA.s89, OFFSET, PTR + 8);
846
847#define vstore_partial_11(DATA, OFFSET, PTR)       \
848    vstore_partial_8(DATA.s01234567, OFFSET, PTR); \
849    vstore_partial_3(DATA.s89a, OFFSET, PTR + 8);
850
851#define vstore_partial_12(DATA, OFFSET, PTR)       \
852    vstore_partial_8(DATA.s01234567, OFFSET, PTR); \
853    vstore_partial_4(DATA.s89ab, OFFSET, PTR + 8);
854
855#define vstore_partial_13(DATA, OFFSET, PTR)       \
856    vstore_partial_8(DATA.s01234567, OFFSET, PTR); \
857    vstore_partial_5(DATA.s89abcdef, OFFSET, PTR + 8);
858
859#define vstore_partial_14(DATA, OFFSET, PTR)       \
860    vstore_partial_8(DATA.s01234567, OFFSET, PTR); \
861    vstore_partial_6(DATA.s89abcdef, OFFSET, PTR + 8);
862
863#define vstore_partial_15(DATA, OFFSET, PTR)       \
864    vstore_partial_8(DATA.s01234567, OFFSET, PTR); \
865    vstore_partial_7(DATA.s89abcdef, OFFSET, PTR + 8);
866
867#define vstore_partial_16(DATA, OFFSET, PTR) \
868    vstore16(DATA, OFFSET, PTR);
869
870
871
872
873
874#define convert_float_sat convert_float
875#define convert_float1_sat convert_float
876#define convert_float2_sat convert_float2
877#define convert_float3_sat convert_float3
878#define convert_float4_sat convert_float4
879#define convert_float8_sat convert_float8
880#define convert_float16_sat convert_float16
881#define convert_half_sat convert_float
882#define convert_half1_sat convert_half
883#define convert_half2_sat convert_half2
884#define convert_half3_sat convert_half3
885#define convert_half4_sat convert_half4
886#define convert_half8_sat convert_half8
887#define convert_half16_sat convert_half16
888
889#define convert_float1 convert_float
890#define convert_half1 convert_half
891#define convert_char1 convert_char
892#define convert_uchar1 convert_uchar
893#define convert_short1 convert_short
894#define convert_ushort1 convert_ushort
895#define convert_int1 convert_int
896#define convert_uint1 convert_uint
897#define convert_long1 convert_long
898#define convert_ulong1 convert_ulong
899#define convert_double1 convert_double
900
901#define convert_char1_sat convert_char_sat
902#define convert_uchar1_sat convert_uchar_sat
903#define convert_uchar2_sat convert_uchar2_sat
904#define convert_uchar3_sat convert_uchar3_sat
905#define convert_uchar4_sat convert_uchar4_sat
906#define convert_uchar8_sat convert_uchar8_sat
907#define convert_uchar16_sat convert_uchar16_sat
908#define convert_short1_sat convert_short_sat
909#define convert_ushort1_sat convert_ushort_sat
910#define convert_int1_sat convert_int_sat
911#define convert_uint1_sat convert_uint_sat
912#define convert_long1_sat convert_long_sat
913#define convert_ulong1_sat convert_ulong_sat
914#define convert_double1_sat convert_double_sat
915
916#define VEC_DATA_TYPE_STR(type, size) type##size
917#define VEC_DATA_TYPE(type, size) VEC_DATA_TYPE_STR(type, size)
918
919#define CONVERT_STR(x, type) (convert_##type((x)))
920#define CONVERT(x, type) CONVERT_STR(x, type)
921
922#define CONVERT_SAT_STR(x, type) (convert_##type##_sat((x)))
923#define CONVERT_SAT(x, type) CONVERT_SAT_STR(x, type)
924
925#define CONVERT_SAT_ROUND_STR(x, type, round) (convert_##type##_sat_##round((x)))
926#define CONVERT_SAT_ROUND(x, type, round) CONVERT_SAT_ROUND_STR(x, type, round)
927
928#define select_vec_dt_uchar(size) uchar##size
929#define select_vec_dt_char(size) char##size
930#define select_vec_dt_ushort(size) ushort##size
931#define select_vec_dt_short(size) short##size
932#define select_vec_dt_half(size) short##size
933#define select_vec_dt_uint(size) uint##size
934#define select_vec_dt_int(size) int##size
935#define select_vec_dt_float(size) int##size
936#define select_vec_dt_ulong(size) ulong##size
937#define select_vec_dt_long(size) long##size
938
939#define SELECT_VEC_DATA_TYPE_STR(type, size) select_vec_dt_##type(size)
940#define SELECT_VEC_DATA_TYPE(type, size) SELECT_VEC_DATA_TYPE_STR(type, size)
941#define SELECT_DATA_TYPE(type) SELECT_VEC_DATA_TYPE_STR(type, 1)
942
943#define signed_int_vec_dt_uchar(size) char##size
944#define signed_int_vec_dt_char(size) char##size
945#define signed_int_vec_dt_ushort(size) short##size
946#define signed_int_vec_dt_short(size) short##size
947#define signed_int_vec_dt_half(size) short##size
948#define signed_int_vec_dt_uint(size) int##size
949#define signed_int_vec_dt_int(size) int##size
950#define signed_int_vec_dt_float(size) int##size
951#define signed_int_vec_dt_ulong(size) long##size
952#define signed_int_vec_dt_long(size) long##size
953
954#define SIGNED_INT_VEC_DATA_TYPE_STR(type, size) signed_int_vec_dt_##type(size)
955#define SIGNED_INT_VEC_DATA_TYPE(type, size) SIGNED_INT_VEC_DATA_TYPE_STR(type, size)
956#define SIGNED_INT_DATA_TYPE(type) SIGNED_INT_VEC_DATA_TYPE_STR(type, 1)
957
958#define sum_reduce_1(x) (x)
959#define sum_reduce_2(x) ((x).s0) + ((x).s1)
960#define sum_reduce_3(x) sum_reduce_2((x).s01) + ((x).s2)
961#define sum_reduce_4(x) sum_reduce_2((x).s01) + sum_reduce_2((x).s23)
962#define sum_reduce_8(x) sum_reduce_4((x).s0123) + sum_reduce_4((x).s4567)
963#define sum_reduce_16(x) sum_reduce_8((x).s01234567) + sum_reduce_8((x).s89ABCDEF)
964
965#define SUM_REDUCE_STR(x, size) sum_reduce_##size(x)
966#define SUM_REDUCE(x, size) SUM_REDUCE_STR(x, size)
967
968#define prod_reduce_1(x) (x)
969#define prod_reduce_2(x) ((x).s0) * ((x).s1)
970#define prod_reduce_3(x) prod_reduce_2((x).s01) * ((x).s2)
971#define prod_reduce_4(x) prod_reduce_2((x).s01) * prod_reduce_2((x).s23)
972#define prod_reduce_8(x) prod_reduce_4((x).s0123) * prod_reduce_4((x).s4567)
973#define prod_reduce_16(x) prod_reduce_8((x).s01234567) * prod_reduce_8((x).s89ABCDEF)
974
975#define PROD_REDUCE_STR(x, size) prod_reduce_##size(x)
976#define PROD_REDUCE(x, size) PROD_REDUCE_STR(x, size)
977
978#define max_reduce_1(x) (x)
979#define max_reduce_2(x) max(((x).s0), ((x).s1))
980#define max_reduce_3(x) max(max_reduce_2((x).s01), ((x).s2))
981#define max_reduce_4(x) max(max_reduce_2((x).s01), max_reduce_2((x).s23))
982#define max_reduce_8(x) max(max_reduce_4((x).s0123), max_reduce_4((x).s4567))
983#define max_reduce_16(x) max(max_reduce_8((x).s01234567), max_reduce_8((x).s89ABCDEF))
984
985#define MAX_REDUCE_STR(x, size) max_reduce_##size(x)
986#define MAX_REDUCE(x, size) MAX_REDUCE_STR(x, size)
987
988#define VECTOR_DECLARATION(name)     \
989    __global uchar *name##_ptr,      \
990    uint        name##_stride_x, \
991    uint        name##_step_x,   \
992    uint        name##_offset_first_element_in_bytes
993
994#define IMAGE_DECLARATION(name)      \
995    __global uchar *name##_ptr,      \
996    uint        name##_stride_x, \
997    uint        name##_step_x,   \
998    uint        name##_stride_y, \
999    uint        name##_step_y,   \
1000    uint        name##_offset_first_element_in_bytes
1001
1002#define TENSOR3D_DECLARATION(name)   \
1003    __global uchar *name##_ptr,      \
1004    uint        name##_stride_x, \
1005    uint        name##_step_x,   \
1006    uint        name##_stride_y, \
1007    uint        name##_step_y,   \
1008    uint        name##_stride_z, \
1009    uint        name##_step_z,   \
1010    uint        name##_offset_first_element_in_bytes
1011
1012#define TENSOR4D_DECLARATION(name)   \
1013    __global uchar *name##_ptr,      \
1014    uint        name##_stride_x, \
1015    uint        name##_step_x,   \
1016    uint        name##_stride_y, \
1017    uint        name##_step_y,   \
1018    uint        name##_stride_z, \
1019    uint        name##_step_z,   \
1020    uint        name##_stride_w, \
1021    uint        name##_step_w,   \
1022    uint        name##_offset_first_element_in_bytes
1023
1024#define TENSOR5D_DECLARATION(name)   \
1025    __global uchar *name##_ptr,      \
1026    uint        name##_stride_x, \
1027    uint        name##_step_x,   \
1028    uint        name##_stride_y, \
1029    uint        name##_step_y,   \
1030    uint        name##_stride_z, \
1031    uint        name##_step_z,   \
1032    uint        name##_stride_w, \
1033    uint        name##_step_w,   \
1034    uint        name##_stride_v, \
1035    uint        name##_step_v,   \
1036    uint        name##_offset_first_element_in_bytes
1037
1038#define CONVERT_TO_VECTOR_STRUCT(name) \
1039    update_vector_workitem_ptr(name##_ptr, name##_offset_first_element_in_bytes, name##_stride_x, name##_step_x)
1040
1041#define CONVERT_TO_VECTOR_STRUCT_NO_STEP(name) \
1042    update_vector_workitem_ptr(name##_ptr, name##_offset_first_element_in_bytes, name##_stride_x, 0)
1043
1044#define CONVERT_TO_IMAGE_STRUCT(name) \
1045    update_image_workitem_ptr(name##_ptr, name##_offset_first_element_in_bytes, name##_stride_x, name##_step_x, name##_stride_y, name##_step_y)
1046
1047#define CONVERT_TO_IMAGE_STRUCT_NO_STEP(name) \
1048    update_image_workitem_ptr(name##_ptr, name##_offset_first_element_in_bytes, name##_stride_x, 0, name##_stride_y, 0)
1049
1050#define CONVERT_TENSOR3D_TO_IMAGE_STRUCT(name) \
1051    update_image_from_tensor3D_workitem_ptr(name##_ptr, name##_offset_first_element_in_bytes, name##_stride_x, name##_step_x, name##_stride_y, name##_step_y, name##_stride_z, name##_step_z)
1052
1053#define CONVERT_TENSOR3D_TO_IMAGE_STRUCT_NO_STEP(name) \
1054    update_image_from_tensor3D_workitem_ptr(name##_ptr, name##_offset_first_element_in_bytes, name##_stride_x, 0, name##_stride_y, 0, name##_stride_z, name##_step_z)
1055
1056#define CONVERT_TENSOR3D_TO_IMAGE_STRUCT(name) \
1057    update_image_from_tensor3D_workitem_ptr(name##_ptr, name##_offset_first_element_in_bytes, name##_stride_x, name##_step_x, name##_stride_y, name##_step_y, name##_stride_z, name##_step_z)
1058
1059#define CONVERT_TO_TENSOR3D_STRUCT(name)                                                                                                           \
1060    update_tensor3D_workitem_ptr(name##_ptr, name##_offset_first_element_in_bytes, name##_stride_x, name##_step_x, name##_stride_y, name##_step_y, \
1061                                 name##_stride_z, name##_step_z)
1062
1063#define CONVERT_TO_TENSOR3D_STRUCT_NO_STEP(name) \
1064    update_tensor3D_workitem_ptr(name##_ptr, name##_offset_first_element_in_bytes, name##_stride_x, 0, name##_stride_y, 0, name##_stride_z, 0)
1065
1066#define CONVERT_TO_TENSOR4D_STRUCT(name, mod_size)                                                                                                 \
1067    update_tensor4D_workitem_ptr(name##_ptr, name##_offset_first_element_in_bytes, name##_stride_x, name##_step_x, name##_stride_y, name##_step_y, \
1068                                 name##_stride_z, name##_step_z, name##_stride_w, name##_step_w, mod_size)
1069
1070#define CONVERT_TO_TENSOR4D_STRUCT_NO_STEP(name, mod_size) \
1071    update_tensor4D_workitem_ptr(name##_ptr, name##_offset_first_element_in_bytes, name##_stride_x, 0, name##_stride_y, 0, name##_stride_z, 0, name##_stride_w, 0, mod_size)
1072
1073#define CONVERT_TO_TENSOR3D_STRUCT_NO_UPDATE_PTR(name)                                                                                       \
1074    tensor3D_ptr_no_update(name##_ptr, name##_offset_first_element_in_bytes, name##_stride_x, name##_step_x, name##_stride_y, name##_step_y, \
1075                           name##_stride_z, name##_step_z)
1076
1077
1078typedef struct Vector
1079{
1080    __global uchar *ptr;
1081    int             offset_first_element_in_bytes;
1082    int             stride_x;
1083} Vector;
1084
1085
1086typedef struct Image
1087{
1088    __global uchar *ptr;
1089    int             offset_first_element_in_bytes;
1090    int             stride_x;
1091    int             stride_y;
1092} Image;
1093
1094
1095typedef struct Tensor3D
1096{
1097    __global uchar *ptr;
1098    int             offset_first_element_in_bytes;
1099    int             stride_x;
1100    int             stride_y;
1101    int             stride_z;
1102} Tensor3D;
1103
1104
1105typedef struct Tensor4D
1106{
1107    __global uchar *ptr;
1108    int             offset_first_element_in_bytes;
1109    int             stride_x;
1110    int             stride_y;
1111    int             stride_z;
1112    int             stride_w;
1113} Tensor4D;
1114
1115
1116inline Vector update_vector_workitem_ptr(__global uchar *ptr, uint offset_first_element_in_bytes, uint stride_x, uint step_x)
1117{
1118    Vector vector =
1119    {
1120        .ptr                           = ptr,
1121        .offset_first_element_in_bytes = offset_first_element_in_bytes,
1122        .stride_x                      = stride_x,
1123    };
1124    vector.ptr += vector.offset_first_element_in_bytes + get_global_id(0) * step_x;
1125    return vector;
1126}
1127
1128
1129inline Image update_image_workitem_ptr(__global uchar *ptr, uint offset_first_element_in_bytes, uint stride_x, uint step_x, uint stride_y, uint step_y)
1130{
1131    Image img =
1132    {
1133        .ptr                           = ptr,
1134        .offset_first_element_in_bytes = offset_first_element_in_bytes,
1135        .stride_x                      = stride_x,
1136        .stride_y                      = stride_y
1137    };
1138    img.ptr += img.offset_first_element_in_bytes + get_global_id(0) * step_x + get_global_id(1) * step_y;
1139    return img;
1140}
1141
1142
1143inline Image update_image_from_tensor3D_workitem_ptr(__global uchar *ptr, uint offset_first_element_in_bytes, uint stride_x, uint step_x, uint stride_y, uint step_y, uint stride_z, uint step_z)
1144{
1145    Image img =
1146    {
1147        .ptr                           = ptr,
1148        .offset_first_element_in_bytes = offset_first_element_in_bytes,
1149        .stride_x                      = stride_x,
1150        .stride_y                      = stride_y
1151    };
1152    img.ptr += img.offset_first_element_in_bytes + get_global_id(0) * step_x + get_global_id(1) * step_y + get_global_id(2) * step_z;
1153    return img;
1154}
1155
1156
1157inline Tensor3D update_tensor3D_workitem_ptr(__global uchar *ptr, uint offset_first_element_in_bytes, uint stride_x, uint step_x, uint stride_y, uint step_y, uint stride_z, uint step_z)
1158{
1159    Tensor3D tensor =
1160    {
1161        .ptr                           = ptr,
1162        .offset_first_element_in_bytes = offset_first_element_in_bytes,
1163        .stride_x                      = stride_x,
1164        .stride_y                      = stride_y,
1165        .stride_z                      = stride_z
1166    };
1167    tensor.ptr += tensor.offset_first_element_in_bytes + get_global_id(0) * step_x + get_global_id(1) * step_y + get_global_id(2) * step_z;
1168    return tensor;
1169}
1170
1171
1172inline Tensor3D tensor3D_ptr_no_update(__global uchar *ptr, uint offset_first_element_in_bytes, uint stride_x, uint step_x, uint stride_y, uint step_y, uint stride_z, uint step_z)
1173{
1174    Tensor3D tensor =
1175    {
1176        .ptr                           = ptr,
1177        .offset_first_element_in_bytes = offset_first_element_in_bytes,
1178        .stride_x                      = stride_x,
1179        .stride_y                      = stride_y,
1180        .stride_z                      = stride_z
1181    };
1182    return tensor;
1183}
1184
1185inline Tensor4D update_tensor4D_workitem_ptr(__global uchar *ptr, uint offset_first_element_in_bytes, uint stride_x, uint step_x, uint stride_y, uint step_y, uint stride_z, uint step_z, uint stride_w,
1186                                             uint step_w,
1187                                             uint mod_size)
1188{
1189    Tensor4D tensor =
1190    {
1191        .ptr                           = ptr,
1192        .offset_first_element_in_bytes = offset_first_element_in_bytes,
1193        .stride_x                      = stride_x,
1194        .stride_y                      = stride_y,
1195        .stride_z                      = stride_z,
1196        .stride_w                      = stride_w
1197    };
1198
1199    tensor.ptr += tensor.offset_first_element_in_bytes + get_global_id(0) * step_x + get_global_id(1) * step_y + (get_global_id(2) % mod_size) * step_z + (get_global_id(2) / mod_size) * step_w;
1200    return tensor;
1201}
1202
1203
1204inline __global const uchar *vector_offset(const Vector *vec, int x)
1205{
1206    return vec->ptr + x * vec->stride_x;
1207}
1208
1209
1210inline __global uchar *offset(const Image *img, int x, int y)
1211{
1212    return img->ptr + x * img->stride_x + y * img->stride_y;
1213}
1214
1215
1216inline __global const uchar *tensor3D_offset(const Tensor3D *tensor, int x, int y, int z)
1217{
1218    return tensor->ptr + x * tensor->stride_x + y * tensor->stride_y + z * tensor->stride_z;
1219}
1220
1221
1222inline __global const uchar *tensor4D_offset(const Tensor4D *tensor, int x, int y, int z, int w)
1223{
1224    return tensor->ptr + x * tensor->stride_x + y * tensor->stride_y + z * tensor->stride_z + w * tensor->stride_w;
1225}
1226
1227
1228inline __global const uchar *tensor3D_index2ptr(const Tensor3D *tensor, uint width, uint height, uint depth, uint index)
1229{
1230    uint num_elements = width * height;
1231
1232    const uint z = index / num_elements;
1233
1234    index %= num_elements;
1235
1236    const uint y = index / width;
1237
1238    index %= width;
1239
1240    const uint x = index;
1241
1242    return tensor->ptr + x * tensor->stride_x + y * tensor->stride_y + z * tensor->stride_z + tensor->offset_first_element_in_bytes;
1243}
1244
1245#endif
1246
1247#if defined(DATA_TYPE)
1248
1249#define TWIDDLE_FACTOR_MULTIPLICATION(phi, input)  \
1250    {                                              \
1251        VEC_DATA_TYPE(DATA_TYPE, 2)                \
1252        w, tmp;                                    \
1253        w.x   = cos(phi);                          \
1254        w.y   = sin(phi);                          \
1255        tmp.x = (w.x * input.x) - (w.y * input.y); \
1256        tmp.y = (w.x * input.y) + (w.y * input.x); \
1257        input = tmp;                               \
1258    }
1259
1260
1261#define DFT_2(c0, c1)               \
1262    {                               \
1263        VEC_DATA_TYPE(DATA_TYPE, 2) \
1264        v0;                         \
1265        v0 = c0;                    \
1266        c0 = v0 + c1;               \
1267        c1 = v0 - c1;               \
1268    }
1269
1270
1271#define SQRT3DIV2 0.86602540378443f
1272
1273
1274#define DFT_3(c0, c1, c2)                             \
1275    {                                                 \
1276        VEC_DATA_TYPE(DATA_TYPE, 2)                   \
1277        v0 = c1 + c2;                                 \
1278        VEC_DATA_TYPE(DATA_TYPE, 2)                   \
1279        v1   = c1 - c2;                               \
1280        c1.x = c0.x - 0.5f * v0.x + v1.y * SQRT3DIV2; \
1281        c1.y = c0.y - 0.5f * v0.y - v1.x * SQRT3DIV2; \
1282        c2.x = c0.x - 0.5f * v0.x - v1.y * SQRT3DIV2; \
1283        c2.y = c0.y - 0.5f * v0.y + v1.x * SQRT3DIV2; \
1284        c0   = c0 + v0;                               \
1285    }
1286
1287
1288#define DFT_4(c0, c1, c2, c3)       \
1289    {                               \
1290        VEC_DATA_TYPE(DATA_TYPE, 2) \
1291        v0, v1, v2, v3;             \
1292        v0   = c0 + c2;             \
1293        v1   = c1 + c3;             \
1294        v2   = c0 - c2;             \
1295        v3.x = c1.y - c3.y;         \
1296        v3.y = c3.x - c1.x;         \
1297        c0   = v0 + v1;             \
1298        c2   = v0 - v1;             \
1299        c1   = v2 + v3;             \
1300        c3   = v2 - v3;             \
1301    }
1302
1303
1304#define W5_A (DATA_TYPE)0.30901699437494f
1305#define W5_B (DATA_TYPE)0.95105651629515f
1306#define W5_C (DATA_TYPE)0.80901699437494f
1307#define W5_D (DATA_TYPE)0.58778525229247f
1308
1309
1310#define DFT_5(c0, c1, c2, c3, c4)                                  \
1311    {                                                              \
1312        VEC_DATA_TYPE(DATA_TYPE, 2)                                \
1313        v0, v1, v2, v3, v4;                                        \
1314        v0 = c0;                                                   \
1315        v1 = W5_A * (c1 + c4) - W5_C * (c2 + c3);                  \
1316        v2 = W5_C * (c1 + c4) - W5_A * (c2 + c3);                  \
1317        v3 = W5_D * (c1 - c4) - W5_B * (c2 - c3);                  \
1318        v4 = W5_B * (c1 - c4) + W5_D * (c2 - c3);                  \
1319        c0 = v0 + c1 + c2 + c3 + c4;                               \
1320        c1 = v0 + v1 + (VEC_DATA_TYPE(DATA_TYPE, 2))(v4.y, -v4.x); \
1321        c2 = v0 - v2 + (VEC_DATA_TYPE(DATA_TYPE, 2))(v3.y, -v3.x); \
1322        c3 = v0 - v2 + (VEC_DATA_TYPE(DATA_TYPE, 2))(-v3.y, v3.x); \
1323        c4 = v0 + v1 + (VEC_DATA_TYPE(DATA_TYPE, 2))(-v4.y, v4.x); \
1324    }
1325
1326
1327#define W7_A (DATA_TYPE)0.62348980185873f
1328#define W7_B (DATA_TYPE)0.78183148246802f
1329#define W7_C (DATA_TYPE)0.22252093395631f
1330#define W7_D (DATA_TYPE)0.97492791218182f
1331#define W7_E (DATA_TYPE)0.90096886790241f
1332#define W7_F (DATA_TYPE)0.43388373911755f
1333
1334
1335#define DFT_7(c0, c1, c2, c3, c4, c5, c6)                            \
1336    {                                                                \
1337        VEC_DATA_TYPE(DATA_TYPE, 2)                                  \
1338        v0, v1, v2, v3, v4, v5, v6;                                  \
1339        v0 = c0;                                                     \
1340        v1 = W7_A * (c1 + c6) - W7_C * (c2 + c5) - W7_E * (c3 + c4); \
1341        v2 = W7_C * (c1 + c6) + W7_E * (c2 + c5) - W7_A * (c3 + c4); \
1342        v3 = W7_E * (c1 + c6) - W7_A * (c2 + c5) + W7_C * (c3 + c4); \
1343        v4 = W7_B * (c1 - c6) + W7_D * (c2 - c5) + W7_F * (c3 - c4); \
1344        v5 = W7_D * (c1 - c6) - W7_F * (c2 - c5) - W7_B * (c3 - c4); \
1345        v6 = W7_F * (c1 - c6) - W7_B * (c2 - c5) + W7_D * (c3 - c4); \
1346        c0 = v0 + c1 + c2 + c3 + c4 + c5 + c6;                       \
1347        c1 = v0 + v1 + (VEC_DATA_TYPE(DATA_TYPE, 2))(v4.y, -v4.x);   \
1348        c2 = v0 - v2 + (VEC_DATA_TYPE(DATA_TYPE, 2))(v5.y, -v5.x);   \
1349        c3 = v0 - v3 + (VEC_DATA_TYPE(DATA_TYPE, 2))(v6.y, -v6.x);   \
1350        c4 = v0 - v3 + (VEC_DATA_TYPE(DATA_TYPE, 2))(-v6.y, v6.x);   \
1351        c5 = v0 - v2 + (VEC_DATA_TYPE(DATA_TYPE, 2))(-v5.y, v5.x);   \
1352        c6 = v0 + v1 + (VEC_DATA_TYPE(DATA_TYPE, 2))(-v4.y, v4.x);   \
1353    }
1354
1355
1356#define DFT_8(c0, c1, c2, c3, c4, c5, c6, c7) \
1357    {                                         \
1358        VEC_DATA_TYPE(DATA_TYPE, 2)           \
1359        v0, v1, v2, v3, v4, v5, v6, v7;       \
1360        VEC_DATA_TYPE(DATA_TYPE, 2)           \
1361        s0, s1, s2, s3, s4, s5, s6, s7;       \
1362        VEC_DATA_TYPE(DATA_TYPE, 2)           \
1363        t0, t1, t2;                           \
1364        v0   = c0 + c4;                       \
1365        v1   = c1 + c5;                       \
1366        v2   = c2 + c6;                       \
1367        v3   = c3 + c7;                       \
1368        v4   = c0 - c4;                       \
1369        v5   = c1 - c5;                       \
1370        v6   = c2 - c6;                       \
1371        v7   = c3 - c7;                       \
1372        s0   = v0 + v2;                       \
1373        s1   = v1 + v3;                       \
1374        s2   = v0 - v2;                       \
1375        s3   = v1 - v3;                       \
1376        s4.x = v4.x - v6.y;                   \
1377        s4.y = v4.y + v6.x;                   \
1378        s5.x = v5.x - v7.y;                   \
1379        s5.y = v5.y + v7.x;                   \
1380        s6.x = v4.x + v6.y;                   \
1381        s6.y = v4.y - v6.x;                   \
1382        s7.x = v5.x + v7.y;                   \
1383        s7.y = v5.y - v7.x;                   \
1384        t0.x = -s3.y;                         \
1385        t0.y = s3.x;                          \
1386        t1.x = M_SQRT1_2_F * (s5.x - s5.y);   \
1387        t1.y = M_SQRT1_2_F * (s5.x + s5.y);   \
1388        t2.x = -M_SQRT1_2_F * (s7.x + s7.y);  \
1389        t2.y = M_SQRT1_2_F * (s7.x - s7.y);   \
1390        c0   = s0 + s1;                       \
1391        c1   = s6 - t2;                       \
1392        c2   = s2 - t0;                       \
1393        c3   = s4 - t1;                       \
1394        c4   = s0 - s1;                       \
1395        c5   = s6 + t2;                       \
1396        c6   = s2 + t0;                       \
1397        c7   = s4 + t1;                       \
1398    }
1399
1400
1401__kernel void fft_radix_2_first_stage_axis_0(
1402    TENSOR3D_DECLARATION(input)
1403#ifndef IN_PLACE
1404    ,
1405    TENSOR3D_DECLARATION(output)
1406#endif
1407)
1408{
1409
1410    Tensor3D input = CONVERT_TO_TENSOR3D_STRUCT(input);
1411#ifdef IN_PLACE
1412    Tensor3D output = input;
1413#else
1414    Tensor3D output = CONVERT_TO_TENSOR3D_STRUCT(output);
1415#endif
1416
1417
1418    VEC_DATA_TYPE(DATA_TYPE, 4)
1419    data = vload4(0, (__global DATA_TYPE *)input.ptr);
1420
1421
1422    DFT_2(data.s01, data.s23);
1423
1424
1425    vstore4(data, 0, (__global DATA_TYPE *)output.ptr);
1426}
1427
1428
1429__kernel void fft_radix_2_first_stage_axis_1(
1430    TENSOR3D_DECLARATION(input)
1431#ifndef IN_PLACE
1432    ,
1433    TENSOR3D_DECLARATION(output)
1434#endif
1435)
1436{
1437
1438    Tensor3D input = CONVERT_TO_TENSOR3D_STRUCT(input);
1439#ifdef IN_PLACE
1440    Tensor3D output = input;
1441#else
1442    Tensor3D output = CONVERT_TO_TENSOR3D_STRUCT(output);
1443#endif
1444
1445
1446    VEC_DATA_TYPE(DATA_TYPE, 2)
1447    data1 = vload2(0, (__global DATA_TYPE *)input.ptr);
1448    VEC_DATA_TYPE(DATA_TYPE, 2)
1449    data2 = vload2(0, (__global DATA_TYPE *)tensor3D_offset(&input, 0, 1, 0));
1450
1451
1452    DFT_2(data1, data2);
1453
1454
1455    vstore2(data1, 0, (__global DATA_TYPE *)output.ptr);
1456    vstore2(data2, 0, (__global DATA_TYPE *)tensor3D_offset(&output, 0, 1, 0));
1457}
1458
1459
1460__kernel void fft_radix_3_first_stage_axis_0(
1461    TENSOR3D_DECLARATION(input)
1462#ifndef IN_PLACE
1463    ,
1464    TENSOR3D_DECLARATION(output)
1465#endif
1466)
1467{
1468
1469    Tensor3D input = CONVERT_TO_TENSOR3D_STRUCT(input);
1470#ifdef IN_PLACE
1471    Tensor3D output = input;
1472#else
1473    Tensor3D output = CONVERT_TO_TENSOR3D_STRUCT(output);
1474#endif
1475
1476
1477    VEC_DATA_TYPE(DATA_TYPE, 4)
1478    data0 = vload4(0, (__global DATA_TYPE *)input.ptr);
1479    VEC_DATA_TYPE(DATA_TYPE, 2)
1480    data1 = vload2(0, (__global DATA_TYPE *)tensor3D_offset(&input, 2, 0, 0));
1481
1482
1483    DFT_3(data0.s01, data0.s23, data1.s01);
1484
1485
1486    vstore4(data0, 0, (__global DATA_TYPE *)output.ptr);
1487    vstore2(data1, 0, (__global DATA_TYPE *)tensor3D_offset(&output, 2, 0, 0));
1488}
1489
1490
1491__kernel void fft_radix_3_first_stage_axis_1(
1492    TENSOR3D_DECLARATION(input)
1493#ifndef IN_PLACE
1494    ,
1495    TENSOR3D_DECLARATION(output)
1496#endif
1497)
1498{
1499
1500    Tensor3D input = CONVERT_TO_TENSOR3D_STRUCT(input);
1501#ifdef IN_PLACE
1502    Tensor3D output = input;
1503#else
1504    Tensor3D output = CONVERT_TO_TENSOR3D_STRUCT(output);
1505#endif
1506
1507
1508    VEC_DATA_TYPE(DATA_TYPE, 2)
1509    data0 = vload2(0, (__global DATA_TYPE *)input.ptr);
1510    VEC_DATA_TYPE(DATA_TYPE, 2)
1511    data1 = vload2(0, (__global DATA_TYPE *)tensor3D_offset(&input, 0, 1, 0));
1512    VEC_DATA_TYPE(DATA_TYPE, 2)
1513    data2 = vload2(0, (__global DATA_TYPE *)tensor3D_offset(&input, 0, 2, 0));
1514
1515
1516    DFT_3(data0, data1, data2);
1517
1518
1519    vstore2(data0, 0, (__global DATA_TYPE *)output.ptr);
1520    vstore2(data1, 0, (__global DATA_TYPE *)tensor3D_offset(&output, 0, 1, 0));
1521    vstore2(data2, 0, (__global DATA_TYPE *)tensor3D_offset(&output, 0, 2, 0));
1522}
1523
1524
1525__kernel void fft_radix_4_first_stage_axis_0(
1526    TENSOR3D_DECLARATION(input)
1527#ifndef IN_PLACE
1528    ,
1529    TENSOR3D_DECLARATION(output)
1530#endif
1531)
1532{
1533
1534    Tensor3D input = CONVERT_TO_TENSOR3D_STRUCT(input);
1535#ifdef IN_PLACE
1536    Tensor3D output = input;
1537#else
1538    Tensor3D output = CONVERT_TO_TENSOR3D_STRUCT(output);
1539#endif
1540
1541
1542    VEC_DATA_TYPE(DATA_TYPE, 8)
1543    data = vload8(0, (__global DATA_TYPE *)input.ptr);
1544
1545
1546    DFT_4(data.s01, data.s23, data.s45, data.s67);
1547
1548
1549    vstore8(data, 0, (__global DATA_TYPE *)output.ptr);
1550}
1551
1552
1553__kernel void fft_radix_4_first_stage_axis_1(
1554    TENSOR3D_DECLARATION(input)
1555#ifndef IN_PLACE
1556    ,
1557    TENSOR3D_DECLARATION(output)
1558#endif
1559)
1560{
1561
1562    Tensor3D input = CONVERT_TO_TENSOR3D_STRUCT(input);
1563#ifdef IN_PLACE
1564    Tensor3D output = input;
1565#else
1566    Tensor3D output = CONVERT_TO_TENSOR3D_STRUCT(output);
1567#endif
1568
1569
1570    VEC_DATA_TYPE(DATA_TYPE, 2)
1571    data0 = vload2(0, (__global DATA_TYPE *)input.ptr);
1572    VEC_DATA_TYPE(DATA_TYPE, 2)
1573    data1 = vload2(0, (__global DATA_TYPE *)tensor3D_offset(&input, 0, 1, 0));
1574    VEC_DATA_TYPE(DATA_TYPE, 2)
1575    data2 = vload2(0, (__global DATA_TYPE *)tensor3D_offset(&input, 0, 2, 0));
1576    VEC_DATA_TYPE(DATA_TYPE, 2)
1577    data3 = vload2(0, (__global DATA_TYPE *)tensor3D_offset(&input, 0, 3, 0));
1578
1579
1580    DFT_4(data0, data1, data2, data3);
1581
1582
1583    vstore2(data0, 0, (__global DATA_TYPE *)output.ptr);
1584    vstore2(data1, 0, (__global DATA_TYPE *)tensor3D_offset(&output, 0, 1, 0));
1585    vstore2(data2, 0, (__global DATA_TYPE *)tensor3D_offset(&output, 0, 2, 0));
1586    vstore2(data3, 0, (__global DATA_TYPE *)tensor3D_offset(&output, 0, 3, 0));
1587}
1588
1589
1590__kernel void fft_radix_5_first_stage_axis_0(
1591    TENSOR3D_DECLARATION(input)
1592#ifndef IN_PLACE
1593    ,
1594    TENSOR3D_DECLARATION(output)
1595#endif
1596)
1597{
1598
1599    Tensor3D input = CONVERT_TO_TENSOR3D_STRUCT(input);
1600#ifdef IN_PLACE
1601    Tensor3D output = input;
1602#else
1603    Tensor3D output = CONVERT_TO_TENSOR3D_STRUCT(output);
1604#endif
1605
1606
1607    VEC_DATA_TYPE(DATA_TYPE, 8)
1608    data0 = vload8(0, (__global DATA_TYPE *)input.ptr);
1609    VEC_DATA_TYPE(DATA_TYPE, 2)
1610    data1 = vload2(0, (__global DATA_TYPE *)tensor3D_offset(&input, 4, 0, 0));
1611
1612
1613    DFT_5(data0.s01, data0.s23, data0.s45, data0.s67, data1.s01);
1614
1615
1616    vstore8(data0, 0, (__global DATA_TYPE *)output.ptr);
1617    vstore2(data1, 0, (__global DATA_TYPE *)tensor3D_offset(&output, 4, 0, 0));
1618}
1619
1620
1621__kernel void fft_radix_5_first_stage_axis_1(
1622    TENSOR3D_DECLARATION(input)
1623#ifndef IN_PLACE
1624    ,
1625    TENSOR3D_DECLARATION(output)
1626#endif
1627)
1628{
1629
1630    Tensor3D input = CONVERT_TO_TENSOR3D_STRUCT(input);
1631#ifdef IN_PLACE
1632    Tensor3D output = input;
1633#else
1634    Tensor3D output = CONVERT_TO_TENSOR3D_STRUCT(output);
1635#endif
1636
1637
1638    VEC_DATA_TYPE(DATA_TYPE, 2)
1639    data0 = vload2(0, (__global DATA_TYPE *)input.ptr);
1640    VEC_DATA_TYPE(DATA_TYPE, 2)
1641    data1 = vload2(0, (__global DATA_TYPE *)tensor3D_offset(&input, 0, 1, 0));
1642    VEC_DATA_TYPE(DATA_TYPE, 2)
1643    data2 = vload2(0, (__global DATA_TYPE *)tensor3D_offset(&input, 0, 2, 0));
1644    VEC_DATA_TYPE(DATA_TYPE, 2)
1645    data3 = vload2(0, (__global DATA_TYPE *)tensor3D_offset(&input, 0, 3, 0));
1646    VEC_DATA_TYPE(DATA_TYPE, 2)
1647    data4 = vload2(0, (__global DATA_TYPE *)tensor3D_offset(&input, 0, 4, 0));
1648
1649
1650    DFT_5(data0, data1, data2, data3, data4);
1651
1652
1653    vstore2(data0, 0, (__global DATA_TYPE *)output.ptr);
1654    vstore2(data1, 0, (__global DATA_TYPE *)tensor3D_offset(&output, 0, 1, 0));
1655    vstore2(data2, 0, (__global DATA_TYPE *)tensor3D_offset(&output, 0, 2, 0));
1656    vstore2(data3, 0, (__global DATA_TYPE *)tensor3D_offset(&output, 0, 3, 0));
1657    vstore2(data4, 0, (__global DATA_TYPE *)tensor3D_offset(&output, 0, 4, 0));
1658}
1659
1660
1661__kernel void fft_radix_7_first_stage_axis_0(
1662    TENSOR3D_DECLARATION(input)
1663#ifndef IN_PLACE
1664    ,
1665    TENSOR3D_DECLARATION(output)
1666#endif
1667)
1668{
1669
1670    Tensor3D input = CONVERT_TO_TENSOR3D_STRUCT(input);
1671#ifdef IN_PLACE
1672    Tensor3D output = input;
1673#else
1674    Tensor3D output = CONVERT_TO_TENSOR3D_STRUCT(output);
1675#endif
1676
1677
1678    VEC_DATA_TYPE(DATA_TYPE, 8)
1679    data0 = vload8(0, (__global DATA_TYPE *)input.ptr);
1680    VEC_DATA_TYPE(DATA_TYPE, 4)
1681    data1 = vload4(0, (__global DATA_TYPE *)tensor3D_offset(&input, 4, 0, 0));
1682    VEC_DATA_TYPE(DATA_TYPE, 2)
1683    data2 = vload2(0, (__global DATA_TYPE *)tensor3D_offset(&input, 6, 0, 0));
1684
1685
1686    DFT_7(data0.s01, data0.s23, data0.s45, data0.s67, data1.s01, data1.s23, data2.s01);
1687
1688
1689    vstore8(data0, 0, (__global DATA_TYPE *)output.ptr);
1690    vstore4(data1, 0, (__global DATA_TYPE *)tensor3D_offset(&output, 4, 0, 0));
1691    vstore2(data2, 0, (__global DATA_TYPE *)tensor3D_offset(&output, 6, 0, 0));
1692}
1693
1694
1695__kernel void fft_radix_7_first_stage_axis_1(
1696    TENSOR3D_DECLARATION(input)
1697#ifndef IN_PLACE
1698    ,
1699    TENSOR3D_DECLARATION(output)
1700#endif
1701)
1702{
1703
1704    Tensor3D input = CONVERT_TO_TENSOR3D_STRUCT(input);
1705#ifdef IN_PLACE
1706    Tensor3D output = input;
1707#else
1708    Tensor3D output = CONVERT_TO_TENSOR3D_STRUCT(output);
1709#endif
1710
1711
1712    VEC_DATA_TYPE(DATA_TYPE, 2)
1713    data0 = vload2(0, (__global DATA_TYPE *)input.ptr);
1714    VEC_DATA_TYPE(DATA_TYPE, 2)
1715    data1 = vload2(0, (__global DATA_TYPE *)tensor3D_offset(&input, 0, 1, 0));
1716    VEC_DATA_TYPE(DATA_TYPE, 2)
1717    data2 = vload2(0, (__global DATA_TYPE *)tensor3D_offset(&input, 0, 2, 0));
1718    VEC_DATA_TYPE(DATA_TYPE, 2)
1719    data3 = vload2(0, (__global DATA_TYPE *)tensor3D_offset(&input, 0, 3, 0));
1720    VEC_DATA_TYPE(DATA_TYPE, 2)
1721    data4 = vload2(0, (__global DATA_TYPE *)tensor3D_offset(&input, 0, 4, 0));
1722    VEC_DATA_TYPE(DATA_TYPE, 2)
1723    data5 = vload2(0, (__global DATA_TYPE *)tensor3D_offset(&input, 0, 5, 0));
1724    VEC_DATA_TYPE(DATA_TYPE, 2)
1725    data6 = vload2(0, (__global DATA_TYPE *)tensor3D_offset(&input, 0, 6, 0));
1726
1727
1728    DFT_7(data0, data1, data2, data3, data4, data5, data6);
1729
1730
1731    vstore2(data0, 0, (__global DATA_TYPE *)output.ptr);
1732    vstore2(data1, 0, (__global DATA_TYPE *)tensor3D_offset(&output, 0, 1, 0));
1733    vstore2(data2, 0, (__global DATA_TYPE *)tensor3D_offset(&output, 0, 2, 0));
1734    vstore2(data3, 0, (__global DATA_TYPE *)tensor3D_offset(&output, 0, 3, 0));
1735    vstore2(data4, 0, (__global DATA_TYPE *)tensor3D_offset(&output, 0, 4, 0));
1736    vstore2(data5, 0, (__global DATA_TYPE *)tensor3D_offset(&output, 0, 5, 0));
1737    vstore2(data6, 0, (__global DATA_TYPE *)tensor3D_offset(&output, 0, 6, 0));
1738}
1739
1740
1741__kernel void fft_radix_8_first_stage_axis_0(
1742    TENSOR3D_DECLARATION(input)
1743#ifndef IN_PLACE
1744    ,
1745    TENSOR3D_DECLARATION(output)
1746#endif
1747)
1748{
1749
1750    Tensor3D input = CONVERT_TO_TENSOR3D_STRUCT(input);
1751#ifdef IN_PLACE
1752    Tensor3D output = input;
1753#else
1754    Tensor3D output = CONVERT_TO_TENSOR3D_STRUCT(output);
1755#endif
1756
1757
1758    VEC_DATA_TYPE(DATA_TYPE, 16)
1759    data = vload16(0, (__global DATA_TYPE *)input.ptr);
1760
1761
1762    DFT_8(data.s01, data.s23, data.s45, data.s67, data.s89, data.sAB, data.sCD, data.sEF);
1763
1764
1765    vstore16(data, 0, (__global DATA_TYPE *)output.ptr);
1766}
1767
1768
1769__kernel void fft_radix_8_first_stage_axis_1(
1770    TENSOR3D_DECLARATION(input)
1771#ifndef IN_PLACE
1772    ,
1773    TENSOR3D_DECLARATION(output)
1774#endif
1775)
1776{
1777
1778    Tensor3D input = CONVERT_TO_TENSOR3D_STRUCT(input);
1779#ifdef IN_PLACE
1780    Tensor3D output = input;
1781#else
1782    Tensor3D output = CONVERT_TO_TENSOR3D_STRUCT(output);
1783#endif
1784
1785
1786    VEC_DATA_TYPE(DATA_TYPE, 2)
1787    data0 = vload2(0, (__global DATA_TYPE *)input.ptr);
1788    VEC_DATA_TYPE(DATA_TYPE, 2)
1789    data1 = vload2(0, (__global DATA_TYPE *)tensor3D_offset(&input, 0, 1, 0));
1790    VEC_DATA_TYPE(DATA_TYPE, 2)
1791    data2 = vload2(0, (__global DATA_TYPE *)tensor3D_offset(&input, 0, 2, 0));
1792    VEC_DATA_TYPE(DATA_TYPE, 2)
1793    data3 = vload2(0, (__global DATA_TYPE *)tensor3D_offset(&input, 0, 3, 0));
1794    VEC_DATA_TYPE(DATA_TYPE, 2)
1795    data4 = vload2(0, (__global DATA_TYPE *)tensor3D_offset(&input, 0, 4, 0));
1796    VEC_DATA_TYPE(DATA_TYPE, 2)
1797    data5 = vload2(0, (__global DATA_TYPE *)tensor3D_offset(&input, 0, 5, 0));
1798    VEC_DATA_TYPE(DATA_TYPE, 2)
1799    data6 = vload2(0, (__global DATA_TYPE *)tensor3D_offset(&input, 0, 6, 0));
1800    VEC_DATA_TYPE(DATA_TYPE, 2)
1801    data7 = vload2(0, (__global DATA_TYPE *)tensor3D_offset(&input, 0, 7, 0));
1802
1803
1804    DFT_8(data0, data1, data2, data3, data4, data5, data6, data7);
1805
1806
1807    vstore2(data0, 0, (__global DATA_TYPE *)output.ptr);
1808    vstore2(data1, 0, (__global DATA_TYPE *)tensor3D_offset(&output, 0, 1, 0));
1809    vstore2(data2, 0, (__global DATA_TYPE *)tensor3D_offset(&output, 0, 2, 0));
1810    vstore2(data3, 0, (__global DATA_TYPE *)tensor3D_offset(&output, 0, 3, 0));
1811    vstore2(data4, 0, (__global DATA_TYPE *)tensor3D_offset(&output, 0, 4, 0));
1812    vstore2(data5, 0, (__global DATA_TYPE *)tensor3D_offset(&output, 0, 5, 0));
1813    vstore2(data6, 0, (__global DATA_TYPE *)tensor3D_offset(&output, 0, 6, 0));
1814    vstore2(data7, 0, (__global DATA_TYPE *)tensor3D_offset(&output, 0, 7, 0));
1815}
1816
1817
1818__kernel void fft_radix_2_axis_0(
1819    TENSOR3D_DECLARATION(input)
1820#ifndef IN_PLACE
1821    ,
1822    TENSOR3D_DECLARATION(output)
1823#endif
1824    ,
1825    uint Nx, uint Ni, float exp_const)
1826{
1827
1828    uint kx = get_global_id(0);
1829
1830
1831    uint nx = kx % Nx;
1832
1833
1834    uint n = nx + (kx / Nx) * Ni;
1835
1836
1837    Tensor3D input = CONVERT_TO_TENSOR3D_STRUCT_NO_STEP(input);
1838    input.ptr += n * input.stride_x + get_global_id(1) * input.stride_y + get_global_id(2) * input.stride_z;
1839#ifdef IN_PLACE
1840    Tensor3D output = input;
1841#else
1842    Tensor3D output = CONVERT_TO_TENSOR3D_STRUCT_NO_STEP(output);
1843    output.ptr += n * output.stride_x + get_global_id(1) * output.stride_y + get_global_id(2) * output.stride_z;
1844#endif
1845
1846
1847    VEC_DATA_TYPE(DATA_TYPE, 2)
1848    c0 = vload2(0, (__global DATA_TYPE *)input.ptr);
1849    VEC_DATA_TYPE(DATA_TYPE, 2)
1850    c1 = vload2(0, (__global DATA_TYPE *)tensor3D_offset(&input, Nx, 0, 0));
1851
1852
1853    DATA_TYPE phi = (DATA_TYPE)nx * (DATA_TYPE)exp_const;
1854
1855
1856    TWIDDLE_FACTOR_MULTIPLICATION(phi, c1);
1857
1858
1859    DFT_2(c0, c1);
1860
1861
1862    vstore2(c0, 0, (__global DATA_TYPE *)output.ptr);
1863    vstore2(c1, 0, (__global DATA_TYPE *)tensor3D_offset(&output, Nx, 0, 0));
1864}
1865
1866
1867__kernel void fft_radix_2_axis_1(
1868    TENSOR3D_DECLARATION(input)
1869#ifndef IN_PLACE
1870    ,
1871    TENSOR3D_DECLARATION(output)
1872#endif
1873    ,
1874    uint Nx, uint Ni, float exp_const)
1875{
1876
1877    uint kx = get_global_id(1);
1878
1879
1880    uint nx = kx % Nx;
1881
1882
1883    uint n = nx + (kx / Nx) * Ni;
1884
1885
1886    Tensor3D input = CONVERT_TO_TENSOR3D_STRUCT_NO_STEP(input);
1887    input.ptr += get_global_id(0) * input.stride_x + n * input.stride_y + get_global_id(2) * input.stride_z;
1888#ifdef IN_PLACE
1889    Tensor3D output = input;
1890#else
1891    Tensor3D output = CONVERT_TO_TENSOR3D_STRUCT_NO_STEP(output);
1892    output.ptr += get_global_id(0) * output.stride_x + n * output.stride_y + get_global_id(2) * output.stride_z;
1893#endif
1894
1895
1896    VEC_DATA_TYPE(DATA_TYPE, 2)
1897    c0 = vload2(0, (__global DATA_TYPE *)input.ptr);
1898    VEC_DATA_TYPE(DATA_TYPE, 2)
1899    c1 = vload2(0, (__global DATA_TYPE *)tensor3D_offset(&input, 0, Nx, 0));
1900
1901
1902    DATA_TYPE phi = (DATA_TYPE)nx * (DATA_TYPE)exp_const;
1903
1904
1905    TWIDDLE_FACTOR_MULTIPLICATION(phi, c1);
1906
1907
1908    DFT_2(c0, c1);
1909
1910
1911    vstore2(c0, 0, (__global DATA_TYPE *)output.ptr);
1912    vstore2(c1, 0, (__global DATA_TYPE *)tensor3D_offset(&output, 0, Nx, 0));
1913}
1914
1915
1916__kernel void fft_radix_3_axis_0(
1917    TENSOR3D_DECLARATION(input)
1918#ifndef IN_PLACE
1919    ,
1920    TENSOR3D_DECLARATION(output)
1921#endif
1922    ,
1923    uint Nx, uint Ni, float exp_const)
1924{
1925
1926    uint kx = get_global_id(0);
1927
1928
1929    uint nx = kx % Nx;
1930
1931
1932    uint n = nx + (kx / Nx) * Ni;
1933
1934
1935    Tensor3D input = CONVERT_TO_TENSOR3D_STRUCT_NO_STEP(input);
1936    input.ptr += n * input.stride_x + get_global_id(1) * input.stride_y + get_global_id(2) * input.stride_z;
1937#ifdef IN_PLACE
1938    Tensor3D output = input;
1939#else
1940    Tensor3D output = CONVERT_TO_TENSOR3D_STRUCT_NO_STEP(output);
1941    output.ptr += n * output.stride_x + get_global_id(1) * output.stride_y + get_global_id(2) * output.stride_z;
1942#endif
1943
1944
1945    VEC_DATA_TYPE(DATA_TYPE, 2)
1946    c0 = vload2(0, (__global DATA_TYPE *)input.ptr);
1947    VEC_DATA_TYPE(DATA_TYPE, 2)
1948    c1 = vload2(0, (__global DATA_TYPE *)tensor3D_offset(&input, Nx, 0, 0));
1949    VEC_DATA_TYPE(DATA_TYPE, 2)
1950    c2 = vload2(0, (__global DATA_TYPE *)tensor3D_offset(&input, 2 * Nx, 0, 0));
1951
1952
1953    DATA_TYPE phi = (DATA_TYPE)nx * (DATA_TYPE)exp_const;
1954
1955
1956    TWIDDLE_FACTOR_MULTIPLICATION(phi, c1);
1957    TWIDDLE_FACTOR_MULTIPLICATION(2 * phi, c2);
1958
1959
1960    DFT_3(c0, c1, c2);
1961
1962
1963    vstore2(c0, 0, (__global DATA_TYPE *)output.ptr);
1964    vstore2(c1, 0, (__global DATA_TYPE *)tensor3D_offset(&output, Nx, 0, 0));
1965    vstore2(c2, 0, (__global DATA_TYPE *)tensor3D_offset(&output, 2 * Nx, 0, 0));
1966}
1967
1968
1969__kernel void fft_radix_3_axis_1(
1970    TENSOR3D_DECLARATION(input)
1971#ifndef IN_PLACE
1972    ,
1973    TENSOR3D_DECLARATION(output)
1974#endif
1975    ,
1976    uint Nx, uint Ni, float exp_const)
1977{
1978
1979    uint kx = get_global_id(1);
1980
1981
1982    uint nx = kx % Nx;
1983
1984
1985    uint n = nx + (kx / Nx) * Ni;
1986
1987
1988    Tensor3D input = CONVERT_TO_TENSOR3D_STRUCT_NO_STEP(input);
1989    input.ptr += get_global_id(0) * input.stride_x + n * input.stride_y + get_global_id(2) * input.stride_z;
1990#ifdef IN_PLACE
1991    Tensor3D output = input;
1992#else
1993    Tensor3D output = CONVERT_TO_TENSOR3D_STRUCT_NO_STEP(output);
1994    output.ptr += get_global_id(0) * output.stride_x + n * output.stride_y + get_global_id(2) * output.stride_z;
1995#endif
1996
1997
1998    VEC_DATA_TYPE(DATA_TYPE, 2)
1999    c0 = vload2(0, (__global DATA_TYPE *)input.ptr);
2000    VEC_DATA_TYPE(DATA_TYPE, 2)
2001    c1 = vload2(0, (__global DATA_TYPE *)tensor3D_offset(&input, 0, Nx, 0));
2002    VEC_DATA_TYPE(DATA_TYPE, 2)
2003    c2 = vload2(0, (__global DATA_TYPE *)tensor3D_offset(&input, 0, 2 * Nx, 0));
2004
2005
2006    DATA_TYPE phi = (DATA_TYPE)nx * (DATA_TYPE)exp_const;
2007
2008
2009    TWIDDLE_FACTOR_MULTIPLICATION(phi, c1);
2010    TWIDDLE_FACTOR_MULTIPLICATION(2 * phi, c2);
2011
2012
2013    DFT_3(c0, c1, c2);
2014
2015
2016    vstore2(c0, 0, (__global DATA_TYPE *)output.ptr);
2017    vstore2(c1, 0, (__global DATA_TYPE *)tensor3D_offset(&output, 0, Nx, 0));
2018    vstore2(c2, 0, (__global DATA_TYPE *)tensor3D_offset(&output, 0, 2 * Nx, 0));
2019}
2020
2021
2022__kernel void fft_radix_4_axis_0(
2023    TENSOR3D_DECLARATION(input)
2024#ifndef IN_PLACE
2025    ,
2026    TENSOR3D_DECLARATION(output)
2027#endif
2028    ,
2029    uint Nx, uint Ni, float exp_const)
2030{
2031
2032    uint kx = get_global_id(0);
2033
2034
2035    uint nx = kx % Nx;
2036
2037
2038    uint n = nx + (kx / Nx) * Ni;
2039
2040
2041    Tensor3D input = CONVERT_TO_TENSOR3D_STRUCT_NO_STEP(input);
2042    input.ptr += n * input.stride_x + get_global_id(1) * input.stride_y + get_global_id(2) * input.stride_z;
2043#ifdef IN_PLACE
2044    Tensor3D output = input;
2045#else
2046    Tensor3D output = CONVERT_TO_TENSOR3D_STRUCT_NO_STEP(output);
2047    output.ptr += n * output.stride_x + get_global_id(1) * output.stride_y + get_global_id(2) * output.stride_z;
2048#endif
2049
2050
2051    VEC_DATA_TYPE(DATA_TYPE, 2)
2052    c0 = vload2(0, (__global DATA_TYPE *)input.ptr);
2053    VEC_DATA_TYPE(DATA_TYPE, 2)
2054    c1 = vload2(0, (__global DATA_TYPE *)tensor3D_offset(&input, Nx, 0, 0));
2055    VEC_DATA_TYPE(DATA_TYPE, 2)
2056    c2 = vload2(0, (__global DATA_TYPE *)tensor3D_offset(&input, 2 * Nx, 0, 0));
2057    VEC_DATA_TYPE(DATA_TYPE, 2)
2058    c3 = vload2(0, (__global DATA_TYPE *)tensor3D_offset(&input, 3 * Nx, 0, 0));
2059
2060
2061    DATA_TYPE phi = (DATA_TYPE)nx * (DATA_TYPE)exp_const;
2062
2063
2064    TWIDDLE_FACTOR_MULTIPLICATION(phi, c1);
2065    TWIDDLE_FACTOR_MULTIPLICATION(2 * phi, c2);
2066    TWIDDLE_FACTOR_MULTIPLICATION(3 * phi, c3);
2067
2068
2069    DFT_4(c0, c1, c2, c3);
2070
2071
2072    vstore2(c0, 0, (__global DATA_TYPE *)output.ptr);
2073    vstore2(c1, 0, (__global DATA_TYPE *)tensor3D_offset(&output, Nx, 0, 0));
2074    vstore2(c2, 0, (__global DATA_TYPE *)tensor3D_offset(&output, 2 * Nx, 0, 0));
2075    vstore2(c3, 0, (__global DATA_TYPE *)tensor3D_offset(&output, 3 * Nx, 0, 0));
2076}
2077
2078
2079__kernel void fft_radix_4_axis_1(
2080    TENSOR3D_DECLARATION(input)
2081#ifndef IN_PLACE
2082    ,
2083    TENSOR3D_DECLARATION(output)
2084#endif
2085    ,
2086    uint Nx, uint Ni, float exp_const)
2087{
2088
2089    uint kx = get_global_id(1);
2090
2091
2092    uint nx = kx % Nx;
2093
2094
2095    uint n = nx + (kx / Nx) * Ni;
2096
2097
2098    Tensor3D input = CONVERT_TO_TENSOR3D_STRUCT_NO_STEP(input);
2099    input.ptr += get_global_id(0) * input.stride_x + n * input.stride_y + get_global_id(2) * input.stride_z;
2100#ifdef IN_PLACE
2101    Tensor3D output = input;
2102#else
2103    Tensor3D output = CONVERT_TO_TENSOR3D_STRUCT_NO_STEP(output);
2104    output.ptr += get_global_id(0) * output.stride_x + n * output.stride_y + get_global_id(2) * output.stride_z;
2105#endif
2106
2107
2108    VEC_DATA_TYPE(DATA_TYPE, 2)
2109    c0 = vload2(0, (__global DATA_TYPE *)input.ptr);
2110    VEC_DATA_TYPE(DATA_TYPE, 2)
2111    c1 = vload2(0, (__global DATA_TYPE *)tensor3D_offset(&input, 0, Nx, 0));
2112    VEC_DATA_TYPE(DATA_TYPE, 2)
2113    c2 = vload2(0, (__global DATA_TYPE *)tensor3D_offset(&input, 0, 2 * Nx, 0));
2114    VEC_DATA_TYPE(DATA_TYPE, 2)
2115    c3 = vload2(0, (__global DATA_TYPE *)tensor3D_offset(&input, 0, 3 * Nx, 0));
2116
2117
2118    DATA_TYPE phi = (DATA_TYPE)nx * (DATA_TYPE)exp_const;
2119
2120
2121    TWIDDLE_FACTOR_MULTIPLICATION(phi, c1);
2122    TWIDDLE_FACTOR_MULTIPLICATION(2 * phi, c2);
2123    TWIDDLE_FACTOR_MULTIPLICATION(3 * phi, c3);
2124
2125
2126    DFT_4(c0, c1, c2, c3);
2127
2128
2129    vstore2(c0, 0, (__global DATA_TYPE *)output.ptr);
2130    vstore2(c1, 0, (__global DATA_TYPE *)tensor3D_offset(&output, 0, Nx, 0));
2131    vstore2(c2, 0, (__global DATA_TYPE *)tensor3D_offset(&output, 0, 2 * Nx, 0));
2132    vstore2(c3, 0, (__global DATA_TYPE *)tensor3D_offset(&output, 0, 3 * Nx, 0));
2133}
2134
2135
2136__kernel void fft_radix_5_axis_0(
2137    TENSOR3D_DECLARATION(input)
2138#ifndef IN_PLACE
2139    ,
2140    TENSOR3D_DECLARATION(output)
2141#endif
2142    ,
2143    uint Nx, uint Ni, float exp_const)
2144{
2145
2146    uint kx = get_global_id(0);
2147
2148
2149    uint nx = kx % Nx;
2150
2151
2152    uint n = nx + (kx / Nx) * Ni;
2153
2154
2155    Tensor3D input = CONVERT_TO_TENSOR3D_STRUCT_NO_STEP(input);
2156    input.ptr += n * input.stride_x + get_global_id(1) * input.stride_y + get_global_id(2) * input.stride_z;
2157#ifdef IN_PLACE
2158    Tensor3D output = input;
2159#else
2160    Tensor3D output = CONVERT_TO_TENSOR3D_STRUCT_NO_STEP(output);
2161    output.ptr += n * output.stride_x + get_global_id(1) * output.stride_y + get_global_id(2) * output.stride_z;
2162#endif
2163
2164
2165    VEC_DATA_TYPE(DATA_TYPE, 2)
2166    c0 = vload2(0, (__global DATA_TYPE *)input.ptr);
2167    VEC_DATA_TYPE(DATA_TYPE, 2)
2168    c1 = vload2(0, (__global DATA_TYPE *)tensor3D_offset(&input, Nx, 0, 0));
2169    VEC_DATA_TYPE(DATA_TYPE, 2)
2170    c2 = vload2(0, (__global DATA_TYPE *)tensor3D_offset(&input, 2 * Nx, 0, 0));
2171    VEC_DATA_TYPE(DATA_TYPE, 2)
2172    c3 = vload2(0, (__global DATA_TYPE *)tensor3D_offset(&input, 3 * Nx, 0, 0));
2173    VEC_DATA_TYPE(DATA_TYPE, 2)
2174    c4 = vload2(0, (__global DATA_TYPE *)tensor3D_offset(&input, 4 * Nx, 0, 0));
2175
2176
2177    DATA_TYPE phi = (DATA_TYPE)nx * (DATA_TYPE)exp_const;
2178
2179
2180    TWIDDLE_FACTOR_MULTIPLICATION(phi, c1);
2181    TWIDDLE_FACTOR_MULTIPLICATION(2 * phi, c2);
2182    TWIDDLE_FACTOR_MULTIPLICATION(3 * phi, c3);
2183    TWIDDLE_FACTOR_MULTIPLICATION(4 * phi, c4);
2184
2185
2186    DFT_5(c0, c1, c2, c3, c4);
2187
2188
2189    vstore2(c0, 0, (__global DATA_TYPE *)output.ptr);
2190    vstore2(c1, 0, (__global DATA_TYPE *)tensor3D_offset(&output, Nx, 0, 0));
2191    vstore2(c2, 0, (__global DATA_TYPE *)tensor3D_offset(&output, 2 * Nx, 0, 0));
2192    vstore2(c3, 0, (__global DATA_TYPE *)tensor3D_offset(&output, 3 * Nx, 0, 0));
2193    vstore2(c4, 0, (__global DATA_TYPE *)tensor3D_offset(&output, 4 * Nx, 0, 0));
2194}
2195
2196
2197__kernel void fft_radix_5_axis_1(
2198    TENSOR3D_DECLARATION(input)
2199#ifndef IN_PLACE
2200    ,
2201    TENSOR3D_DECLARATION(output)
2202#endif
2203    ,
2204    uint Nx, uint Ni, float exp_const)
2205{
2206
2207    uint kx = get_global_id(1);
2208
2209
2210    uint nx = kx % Nx;
2211
2212
2213    uint n = nx + (kx / Nx) * Ni;
2214
2215
2216    Tensor3D input = CONVERT_TO_TENSOR3D_STRUCT_NO_STEP(input);
2217    input.ptr += get_global_id(0) * input.stride_x + n * input.stride_y + get_global_id(2) * input.stride_z;
2218#ifdef IN_PLACE
2219    Tensor3D output = input;
2220#else
2221    Tensor3D output = CONVERT_TO_TENSOR3D_STRUCT_NO_STEP(output);
2222    output.ptr += get_global_id(0) * output.stride_x + n * output.stride_y + get_global_id(2) * output.stride_z;
2223#endif
2224
2225
2226    VEC_DATA_TYPE(DATA_TYPE, 2)
2227    c0 = vload2(0, (__global DATA_TYPE *)input.ptr);
2228    VEC_DATA_TYPE(DATA_TYPE, 2)
2229    c1 = vload2(0, (__global DATA_TYPE *)tensor3D_offset(&input, 0, Nx, 0));
2230    VEC_DATA_TYPE(DATA_TYPE, 2)
2231    c2 = vload2(0, (__global DATA_TYPE *)tensor3D_offset(&input, 0, 2 * Nx, 0));
2232    VEC_DATA_TYPE(DATA_TYPE, 2)
2233    c3 = vload2(0, (__global DATA_TYPE *)tensor3D_offset(&input, 0, 3 * Nx, 0));
2234    VEC_DATA_TYPE(DATA_TYPE, 2)
2235    c4 = vload2(0, (__global DATA_TYPE *)tensor3D_offset(&input, 0, 4 * Nx, 0));
2236
2237
2238    DATA_TYPE phi = (DATA_TYPE)nx * (DATA_TYPE)exp_const;
2239
2240
2241    TWIDDLE_FACTOR_MULTIPLICATION(phi, c1);
2242    TWIDDLE_FACTOR_MULTIPLICATION(2 * phi, c2);
2243    TWIDDLE_FACTOR_MULTIPLICATION(3 * phi, c3);
2244    TWIDDLE_FACTOR_MULTIPLICATION(4 * phi, c4);
2245
2246
2247    DFT_5(c0, c1, c2, c3, c4);
2248
2249
2250    vstore2(c0, 0, (__global DATA_TYPE *)output.ptr);
2251    vstore2(c1, 0, (__global DATA_TYPE *)tensor3D_offset(&output, 0, Nx, 0));
2252    vstore2(c2, 0, (__global DATA_TYPE *)tensor3D_offset(&output, 0, 2 * Nx, 0));
2253    vstore2(c3, 0, (__global DATA_TYPE *)tensor3D_offset(&output, 0, 3 * Nx, 0));
2254    vstore2(c4, 0, (__global DATA_TYPE *)tensor3D_offset(&output, 0, 4 * Nx, 0));
2255}
2256
2257
2258__kernel void fft_radix_7_axis_0(
2259    TENSOR3D_DECLARATION(input)
2260#ifndef IN_PLACE
2261    ,
2262    TENSOR3D_DECLARATION(output)
2263#endif
2264    ,
2265    uint Nx, uint Ni, float exp_const)
2266{
2267
2268    uint kx = get_global_id(0);
2269
2270
2271    uint nx = kx % Nx;
2272
2273
2274    uint n = nx + (kx / Nx) * Ni;
2275
2276
2277    Tensor3D input = CONVERT_TO_TENSOR3D_STRUCT_NO_STEP(input);
2278    input.ptr += n * input.stride_x + get_global_id(1) * input.stride_y + get_global_id(2) * input.stride_z;
2279#ifdef IN_PLACE
2280    Tensor3D output = input;
2281#else
2282    Tensor3D output = CONVERT_TO_TENSOR3D_STRUCT_NO_STEP(output);
2283    output.ptr += n * output.stride_x + get_global_id(1) * output.stride_y + get_global_id(2) * output.stride_z;
2284#endif
2285
2286
2287    VEC_DATA_TYPE(DATA_TYPE, 2)
2288    c0 = vload2(0, (__global DATA_TYPE *)input.ptr);
2289    VEC_DATA_TYPE(DATA_TYPE, 2)
2290    c1 = vload2(0, (__global DATA_TYPE *)tensor3D_offset(&input, Nx, 0, 0));
2291    VEC_DATA_TYPE(DATA_TYPE, 2)
2292    c2 = vload2(0, (__global DATA_TYPE *)tensor3D_offset(&input, 2 * Nx, 0, 0));
2293    VEC_DATA_TYPE(DATA_TYPE, 2)
2294    c3 = vload2(0, (__global DATA_TYPE *)tensor3D_offset(&input, 3 * Nx, 0, 0));
2295    VEC_DATA_TYPE(DATA_TYPE, 2)
2296    c4 = vload2(0, (__global DATA_TYPE *)tensor3D_offset(&input, 4 * Nx, 0, 0));
2297    VEC_DATA_TYPE(DATA_TYPE, 2)
2298    c5 = vload2(0, (__global DATA_TYPE *)tensor3D_offset(&input, 5 * Nx, 0, 0));
2299    VEC_DATA_TYPE(DATA_TYPE, 2)
2300    c6 = vload2(0, (__global DATA_TYPE *)tensor3D_offset(&input, 6 * Nx, 0, 0));
2301
2302
2303    DATA_TYPE phi = (DATA_TYPE)nx * (DATA_TYPE)exp_const;
2304
2305
2306    TWIDDLE_FACTOR_MULTIPLICATION(phi, c1);
2307    TWIDDLE_FACTOR_MULTIPLICATION(2 * phi, c2);
2308    TWIDDLE_FACTOR_MULTIPLICATION(3 * phi, c3);
2309    TWIDDLE_FACTOR_MULTIPLICATION(4 * phi, c4);
2310    TWIDDLE_FACTOR_MULTIPLICATION(5 * phi, c5);
2311    TWIDDLE_FACTOR_MULTIPLICATION(6 * phi, c6);
2312
2313
2314    DFT_7(c0, c1, c2, c3, c4, c5, c6);
2315
2316
2317    vstore2(c0, 0, (__global DATA_TYPE *)output.ptr);
2318    vstore2(c1, 0, (__global DATA_TYPE *)tensor3D_offset(&output, Nx, 0, 0));
2319    vstore2(c2, 0, (__global DATA_TYPE *)tensor3D_offset(&output, 2 * Nx, 0, 0));
2320    vstore2(c3, 0, (__global DATA_TYPE *)tensor3D_offset(&output, 3 * Nx, 0, 0));
2321    vstore2(c4, 0, (__global DATA_TYPE *)tensor3D_offset(&output, 4 * Nx, 0, 0));
2322    vstore2(c5, 0, (__global DATA_TYPE *)tensor3D_offset(&output, 5 * Nx, 0, 0));
2323    vstore2(c6, 0, (__global DATA_TYPE *)tensor3D_offset(&output, 6 * Nx, 0, 0));
2324}
2325
2326
2327__kernel void fft_radix_7_axis_1(
2328    TENSOR3D_DECLARATION(input)
2329#ifndef IN_PLACE
2330    ,
2331    TENSOR3D_DECLARATION(output)
2332#endif
2333    ,
2334    uint Nx, uint Ni, float exp_const)
2335{
2336
2337    uint kx = get_global_id(1);
2338
2339
2340    uint nx = kx % Nx;
2341
2342
2343    uint n = nx + (kx / Nx) * Ni;
2344
2345
2346    Tensor3D input = CONVERT_TO_TENSOR3D_STRUCT_NO_STEP(input);
2347    input.ptr += get_global_id(0) * input.stride_x + n * input.stride_y + get_global_id(2) * input.stride_z;
2348#ifdef IN_PLACE
2349    Tensor3D output = input;
2350#else
2351    Tensor3D output = CONVERT_TO_TENSOR3D_STRUCT_NO_STEP(output);
2352    output.ptr += get_global_id(0) * output.stride_x + n * output.stride_y + get_global_id(2) * output.stride_z;
2353#endif
2354
2355
2356    VEC_DATA_TYPE(DATA_TYPE, 2)
2357    c0 = vload2(0, (__global DATA_TYPE *)input.ptr);
2358    VEC_DATA_TYPE(DATA_TYPE, 2)
2359    c1 = vload2(0, (__global DATA_TYPE *)tensor3D_offset(&input, 0, Nx, 0));
2360    VEC_DATA_TYPE(DATA_TYPE, 2)
2361    c2 = vload2(0, (__global DATA_TYPE *)tensor3D_offset(&input, 0, 2 * Nx, 0));
2362    VEC_DATA_TYPE(DATA_TYPE, 2)
2363    c3 = vload2(0, (__global DATA_TYPE *)tensor3D_offset(&input, 0, 3 * Nx, 0));
2364    VEC_DATA_TYPE(DATA_TYPE, 2)
2365    c4 = vload2(0, (__global DATA_TYPE *)tensor3D_offset(&input, 0, 4 * Nx, 0));
2366    VEC_DATA_TYPE(DATA_TYPE, 2)
2367    c5 = vload2(0, (__global DATA_TYPE *)tensor3D_offset(&input, 0, 5 * Nx, 0));
2368    VEC_DATA_TYPE(DATA_TYPE, 2)
2369    c6 = vload2(0, (__global DATA_TYPE *)tensor3D_offset(&input, 0, 6 * Nx, 0));
2370
2371
2372    DATA_TYPE phi = (DATA_TYPE)nx * (DATA_TYPE)exp_const;
2373
2374
2375    TWIDDLE_FACTOR_MULTIPLICATION(phi, c1);
2376    TWIDDLE_FACTOR_MULTIPLICATION(2 * phi, c2);
2377    TWIDDLE_FACTOR_MULTIPLICATION(3 * phi, c3);
2378    TWIDDLE_FACTOR_MULTIPLICATION(4 * phi, c4);
2379    TWIDDLE_FACTOR_MULTIPLICATION(5 * phi, c5);
2380    TWIDDLE_FACTOR_MULTIPLICATION(6 * phi, c6);
2381
2382
2383    DFT_7(c0, c1, c2, c3, c4, c5, c6);
2384
2385
2386    vstore2(c0, 0, (__global DATA_TYPE *)output.ptr);
2387    vstore2(c1, 0, (__global DATA_TYPE *)tensor3D_offset(&output, 0, Nx, 0));
2388    vstore2(c2, 0, (__global DATA_TYPE *)tensor3D_offset(&output, 0, 2 * Nx, 0));
2389    vstore2(c3, 0, (__global DATA_TYPE *)tensor3D_offset(&output, 0, 3 * Nx, 0));
2390    vstore2(c4, 0, (__global DATA_TYPE *)tensor3D_offset(&output, 0, 4 * Nx, 0));
2391    vstore2(c5, 0, (__global DATA_TYPE *)tensor3D_offset(&output, 0, 5 * Nx, 0));
2392    vstore2(c6, 0, (__global DATA_TYPE *)tensor3D_offset(&output, 0, 6 * Nx, 0));
2393}
2394
2395
2396__kernel void fft_radix_8_axis_0(
2397    TENSOR3D_DECLARATION(input)
2398#ifndef IN_PLACE
2399    ,
2400    TENSOR3D_DECLARATION(output)
2401#endif
2402    ,
2403    uint Nx, uint Ni, float exp_const)
2404{
2405
2406    uint kx = get_global_id(0);
2407
2408
2409    uint nx = kx % Nx;
2410
2411
2412    uint n = nx + (kx / Nx) * Ni;
2413
2414
2415    Tensor3D input = CONVERT_TO_TENSOR3D_STRUCT_NO_STEP(input);
2416    input.ptr += n * input.stride_x + get_global_id(1) * input.stride_y + get_global_id(2) * input.stride_z;
2417#ifdef IN_PLACE
2418    Tensor3D output = input;
2419#else
2420    Tensor3D output = CONVERT_TO_TENSOR3D_STRUCT_NO_STEP(output);
2421    output.ptr += n * output.stride_x + get_global_id(1) * output.stride_y + get_global_id(2) * output.stride_z;
2422#endif
2423
2424
2425    VEC_DATA_TYPE(DATA_TYPE, 2)
2426    c0 = vload2(0, (__global DATA_TYPE *)input.ptr);
2427    VEC_DATA_TYPE(DATA_TYPE, 2)
2428    c1 = vload2(0, (__global DATA_TYPE *)tensor3D_offset(&input, Nx, 0, 0));
2429    VEC_DATA_TYPE(DATA_TYPE, 2)
2430    c2 = vload2(0, (__global DATA_TYPE *)tensor3D_offset(&input, 2 * Nx, 0, 0));
2431    VEC_DATA_TYPE(DATA_TYPE, 2)
2432    c3 = vload2(0, (__global DATA_TYPE *)tensor3D_offset(&input, 3 * Nx, 0, 0));
2433    VEC_DATA_TYPE(DATA_TYPE, 2)
2434    c4 = vload2(0, (__global DATA_TYPE *)tensor3D_offset(&input, 4 * Nx, 0, 0));
2435    VEC_DATA_TYPE(DATA_TYPE, 2)
2436    c5 = vload2(0, (__global DATA_TYPE *)tensor3D_offset(&input, 5 * Nx, 0, 0));
2437    VEC_DATA_TYPE(DATA_TYPE, 2)
2438    c6 = vload2(0, (__global DATA_TYPE *)tensor3D_offset(&input, 6 * Nx, 0, 0));
2439    VEC_DATA_TYPE(DATA_TYPE, 2)
2440    c7 = vload2(0, (__global DATA_TYPE *)tensor3D_offset(&input, 7 * Nx, 0, 0));
2441
2442
2443    DATA_TYPE phi = (DATA_TYPE)nx * (DATA_TYPE)exp_const;
2444
2445
2446    TWIDDLE_FACTOR_MULTIPLICATION(phi, c1);
2447    TWIDDLE_FACTOR_MULTIPLICATION(2 * phi, c2);
2448    TWIDDLE_FACTOR_MULTIPLICATION(3 * phi, c3);
2449    TWIDDLE_FACTOR_MULTIPLICATION(4 * phi, c4);
2450    TWIDDLE_FACTOR_MULTIPLICATION(5 * phi, c5);
2451    TWIDDLE_FACTOR_MULTIPLICATION(6 * phi, c6);
2452    TWIDDLE_FACTOR_MULTIPLICATION(7 * phi, c7);
2453
2454
2455    DFT_8(c0, c1, c2, c3, c4, c5, c6, c7);
2456
2457
2458    vstore2(c0, 0, (__global DATA_TYPE *)output.ptr);
2459    vstore2(c1, 0, (__global DATA_TYPE *)tensor3D_offset(&output, Nx, 0, 0));
2460    vstore2(c2, 0, (__global DATA_TYPE *)tensor3D_offset(&output, 2 * Nx, 0, 0));
2461    vstore2(c3, 0, (__global DATA_TYPE *)tensor3D_offset(&output, 3 * Nx, 0, 0));
2462    vstore2(c4, 0, (__global DATA_TYPE *)tensor3D_offset(&output, 4 * Nx, 0, 0));
2463    vstore2(c5, 0, (__global DATA_TYPE *)tensor3D_offset(&output, 5 * Nx, 0, 0));
2464    vstore2(c6, 0, (__global DATA_TYPE *)tensor3D_offset(&output, 6 * Nx, 0, 0));
2465    vstore2(c7, 0, (__global DATA_TYPE *)tensor3D_offset(&output, 7 * Nx, 0, 0));
2466}
2467
2468
2469__kernel void fft_radix_8_axis_1(
2470    TENSOR3D_DECLARATION(input)
2471#ifndef IN_PLACE
2472    ,
2473    TENSOR3D_DECLARATION(output)
2474#endif
2475    ,
2476    uint Nx, uint Ni, float exp_const)
2477{
2478
2479    uint kx = get_global_id(1);
2480
2481
2482    uint nx = kx % Nx;
2483
2484
2485    uint n = nx + (kx / Nx) * Ni;
2486
2487
2488    Tensor3D input = CONVERT_TO_TENSOR3D_STRUCT_NO_STEP(input);
2489    input.ptr += get_global_id(0) * input.stride_x + n * input.stride_y + get_global_id(2) * input.stride_z;
2490#ifdef IN_PLACE
2491    Tensor3D output = input;
2492#else
2493    Tensor3D output = CONVERT_TO_TENSOR3D_STRUCT_NO_STEP(output);
2494    output.ptr += get_global_id(0) * output.stride_x + n * output.stride_y + get_global_id(2) * output.stride_z;
2495#endif
2496
2497
2498    VEC_DATA_TYPE(DATA_TYPE, 2)
2499    c0 = vload2(0, (__global DATA_TYPE *)input.ptr);
2500    VEC_DATA_TYPE(DATA_TYPE, 2)
2501    c1 = vload2(0, (__global DATA_TYPE *)tensor3D_offset(&input, 0, Nx, 0));
2502    VEC_DATA_TYPE(DATA_TYPE, 2)
2503    c2 = vload2(0, (__global DATA_TYPE *)tensor3D_offset(&input, 0, 2 * Nx, 0));
2504    VEC_DATA_TYPE(DATA_TYPE, 2)
2505    c3 = vload2(0, (__global DATA_TYPE *)tensor3D_offset(&input, 0, 3 * Nx, 0));
2506    VEC_DATA_TYPE(DATA_TYPE, 2)
2507    c4 = vload2(0, (__global DATA_TYPE *)tensor3D_offset(&input, 0, 4 * Nx, 0));
2508    VEC_DATA_TYPE(DATA_TYPE, 2)
2509    c5 = vload2(0, (__global DATA_TYPE *)tensor3D_offset(&input, 0, 5 * Nx, 0));
2510    VEC_DATA_TYPE(DATA_TYPE, 2)
2511    c6 = vload2(0, (__global DATA_TYPE *)tensor3D_offset(&input, 0, 6 * Nx, 0));
2512    VEC_DATA_TYPE(DATA_TYPE, 2)
2513    c7 = vload2(0, (__global DATA_TYPE *)tensor3D_offset(&input, 0, 7 * Nx, 0));
2514
2515
2516    DATA_TYPE phi = (DATA_TYPE)nx * (DATA_TYPE)exp_const;
2517
2518
2519    TWIDDLE_FACTOR_MULTIPLICATION(phi, c1);
2520    TWIDDLE_FACTOR_MULTIPLICATION(2 * phi, c2);
2521    TWIDDLE_FACTOR_MULTIPLICATION(3 * phi, c3);
2522    TWIDDLE_FACTOR_MULTIPLICATION(4 * phi, c4);
2523    TWIDDLE_FACTOR_MULTIPLICATION(5 * phi, c5);
2524    TWIDDLE_FACTOR_MULTIPLICATION(6 * phi, c6);
2525    TWIDDLE_FACTOR_MULTIPLICATION(7 * phi, c7);
2526
2527
2528    DFT_8(c0, c1, c2, c3, c4, c5, c6, c7);
2529
2530
2531    vstore2(c0, 0, (__global DATA_TYPE *)output.ptr);
2532    vstore2(c1, 0, (__global DATA_TYPE *)tensor3D_offset(&output, 0, Nx, 0));
2533    vstore2(c2, 0, (__global DATA_TYPE *)tensor3D_offset(&output, 0, 2 * Nx, 0));
2534    vstore2(c3, 0, (__global DATA_TYPE *)tensor3D_offset(&output, 0, 3 * Nx, 0));
2535    vstore2(c4, 0, (__global DATA_TYPE *)tensor3D_offset(&output, 0, 4 * Nx, 0));
2536    vstore2(c5, 0, (__global DATA_TYPE *)tensor3D_offset(&output, 0, 5 * Nx, 0));
2537    vstore2(c6, 0, (__global DATA_TYPE *)tensor3D_offset(&output, 0, 6 * Nx, 0));
2538    vstore2(c7, 0, (__global DATA_TYPE *)tensor3D_offset(&output, 0, 7 * Nx, 0));
2539}
2540#endif  )"