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