xref: /aosp_15_r20/external/ComputeLibrary/cl_kernels/nhwc/scale.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#ifndef SRC_CORE_CL_CL_KERNELS_TILE_HELPERS
1248*c217d954SCole Faust#define SRC_CORE_CL_CL_KERNELS_TILE_HELPERS
1249*c217d954SCole Faust
1250*c217d954SCole Faust
1251*c217d954SCole Faust
1252*c217d954SCole Faust
1253*c217d954SCole Faust#define TILE_VECTOR_SIZE1 1
1254*c217d954SCole Faust#define TILE_VECTOR_SIZE2 2
1255*c217d954SCole Faust#define TILE_VECTOR_SIZE3 3
1256*c217d954SCole Faust#define TILE_VECTOR_SIZE4 4
1257*c217d954SCole Faust#define TILE_VECTOR_SIZE5 8
1258*c217d954SCole Faust#define TILE_VECTOR_SIZE6 8
1259*c217d954SCole Faust#define TILE_VECTOR_SIZE7 8
1260*c217d954SCole Faust#define TILE_VECTOR_SIZE8 8
1261*c217d954SCole Faust#define TILE_VECTOR_SIZE9 16
1262*c217d954SCole Faust#define TILE_VECTOR_SIZE10 16
1263*c217d954SCole Faust#define TILE_VECTOR_SIZE11 16
1264*c217d954SCole Faust#define TILE_VECTOR_SIZE12 16
1265*c217d954SCole Faust#define TILE_VECTOR_SIZE13 16
1266*c217d954SCole Faust#define TILE_VECTOR_SIZE14 16
1267*c217d954SCole Faust#define TILE_VECTOR_SIZE15 16
1268*c217d954SCole Faust#define TILE_VECTOR_SIZE16 16
1269*c217d954SCole Faust
1270*c217d954SCole Faust#define TILE_VECTOR_TYPE1(DATA_TYPE) DATA_TYPE##1
1271*c217d954SCole Faust#define TILE_VECTOR_TYPE2(DATA_TYPE) DATA_TYPE##2
1272*c217d954SCole Faust#define TILE_VECTOR_TYPE3(DATA_TYPE) DATA_TYPE##3
1273*c217d954SCole Faust#define TILE_VECTOR_TYPE4(DATA_TYPE) DATA_TYPE##4
1274*c217d954SCole Faust#define TILE_VECTOR_TYPE5(DATA_TYPE) DATA_TYPE##8
1275*c217d954SCole Faust#define TILE_VECTOR_TYPE6(DATA_TYPE) DATA_TYPE##8
1276*c217d954SCole Faust#define TILE_VECTOR_TYPE7(DATA_TYPE) DATA_TYPE##8
1277*c217d954SCole Faust#define TILE_VECTOR_TYPE8(DATA_TYPE) DATA_TYPE##8
1278*c217d954SCole Faust#define TILE_VECTOR_TYPE9(DATA_TYPE) DATA_TYPE##16
1279*c217d954SCole Faust#define TILE_VECTOR_TYPE10(DATA_TYPE) DATA_TYPE##16
1280*c217d954SCole Faust#define TILE_VECTOR_TYPE11(DATA_TYPE) DATA_TYPE##16
1281*c217d954SCole Faust#define TILE_VECTOR_TYPE12(DATA_TYPE) DATA_TYPE##16
1282*c217d954SCole Faust#define TILE_VECTOR_TYPE13(DATA_TYPE) DATA_TYPE##16
1283*c217d954SCole Faust#define TILE_VECTOR_TYPE14(DATA_TYPE) DATA_TYPE##16
1284*c217d954SCole Faust#define TILE_VECTOR_TYPE15(DATA_TYPE) DATA_TYPE##16
1285*c217d954SCole Faust#define TILE_VECTOR_TYPE16(DATA_TYPE) DATA_TYPE##16
1286*c217d954SCole Faust
1287*c217d954SCole Faust
1288*c217d954SCole Faust#define TILE(DATA_TYPE, H, W, BASENAME) TILE_STR(DATA_TYPE, H, W, BASENAME)
1289*c217d954SCole Faust#define TILE_STR(DATA_TYPE, H, W, BASENAME) \
1290*c217d954SCole Faust    union {                                 \
1291*c217d954SCole Faust        DATA_TYPE                      s[TILE_VECTOR_SIZE##W];                  \
1292*c217d954SCole Faust        TILE_VECTOR_TYPE##W(DATA_TYPE) v;                     \
1293*c217d954SCole Faust    } BASENAME[H]
1294*c217d954SCole Faust
1295*c217d954SCole Faust#define TENSOR4D_IMAGE(name)          \
1296*c217d954SCole Faust    __read_only image2d_t name##_img, \
1297*c217d954SCole Faust    __global uchar *name##_ptr,       \
1298*c217d954SCole Faust    uint            name##_stride_x,  \
1299*c217d954SCole Faust    uint            name##_step_x,    \
1300*c217d954SCole Faust    uint            name##_stride_y,  \
1301*c217d954SCole Faust    uint            name##_step_y,    \
1302*c217d954SCole Faust    uint            name##_stride_z,  \
1303*c217d954SCole Faust    uint            name##_step_z,    \
1304*c217d954SCole Faust    uint            name##_stride_w,  \
1305*c217d954SCole Faust    uint            name##_step_w,    \
1306*c217d954SCole Faust    uint            name##_offset_first_element_in_bytes
1307*c217d954SCole Faust
1308*c217d954SCole Faust#define TENSOR4D_BUFFER(name)    \
1309*c217d954SCole Faust    __global uchar *name##_ptr,  \
1310*c217d954SCole Faust    uint        name##_stride_x, \
1311*c217d954SCole Faust    uint        name##_step_x,   \
1312*c217d954SCole Faust    uint        name##_stride_y, \
1313*c217d954SCole Faust    uint        name##_step_y,   \
1314*c217d954SCole Faust    uint        name##_stride_z, \
1315*c217d954SCole Faust    uint        name##_step_z,   \
1316*c217d954SCole Faust    uint        name##_stride_w, \
1317*c217d954SCole Faust    uint        name##_step_w,   \
1318*c217d954SCole Faust    uint        name##_offset_first_element_in_bytes
1319*c217d954SCole Faust
1320*c217d954SCole Faust#define TENSOR4D_STR(name, type) TENSOR4D_##type(name)
1321*c217d954SCole Faust#define TENSOR4D(name, type) TENSOR4D_STR(name, type)
1322*c217d954SCole Faust
1323*c217d954SCole Faust#define TENSOR4D_T_IMAGE(name)          \
1324*c217d954SCole Faust    __read_only image2d_t name##_img, \
1325*c217d954SCole Faust    __global uchar *name##_ptr,       \
1326*c217d954SCole Faust    uint        name##_stride_y, \
1327*c217d954SCole Faust    uint        name##_stride_z, \
1328*c217d954SCole Faust    uint        name##_stride_w, \
1329*c217d954SCole Faust    uint        name##_c,   \
1330*c217d954SCole Faust    uint        name##_w,   \
1331*c217d954SCole Faust    uint        name##_h,   \
1332*c217d954SCole Faust    uint        name##_n,   \
1333*c217d954SCole Faust    uint        name##_offset_first_element_in_bytes
1334*c217d954SCole Faust
1335*c217d954SCole Faust#define TENSOR4D_T_BUFFER(name)    \
1336*c217d954SCole Faust    __global uchar *name##_ptr,  \
1337*c217d954SCole Faust    uint        name##_stride_y, \
1338*c217d954SCole Faust    uint        name##_stride_z, \
1339*c217d954SCole Faust    uint        name##_stride_w, \
1340*c217d954SCole Faust    uint        name##_c,   \
1341*c217d954SCole Faust    uint        name##_w,   \
1342*c217d954SCole Faust    uint        name##_h,   \
1343*c217d954SCole Faust    uint        name##_n,   \
1344*c217d954SCole Faust    uint        name##_offset_first_element_in_bytes
1345*c217d954SCole Faust
1346*c217d954SCole Faust#define TENSOR4D_T_STR(name, type) TENSOR4D_T_##type(name)
1347*c217d954SCole Faust
1348*c217d954SCole Faust
1349*c217d954SCole Faust#define TENSOR4D_T(name, type) TENSOR4D_T_STR(name, type)
1350*c217d954SCole Faust
1351*c217d954SCole Faust#define TENSOR4D_RO_T_IMAGE(name)          \
1352*c217d954SCole Faust    __read_only image2d_t name##_img, \
1353*c217d954SCole Faust    TENSOR4D_T_BUFFER(name)
1354*c217d954SCole Faust
1355*c217d954SCole Faust#define TENSOR4D_RO_T_BUFFER(name) TENSOR4D_T_BUFFER(name)
1356*c217d954SCole Faust
1357*c217d954SCole Faust#define TENSOR4D_RO_T_STR(name, type) TENSOR4D_RO_T_##type(name)
1358*c217d954SCole Faust
1359*c217d954SCole Faust
1360*c217d954SCole Faust#define TENSOR4D_RO_T(name, type) TENSOR4D_RO_T_STR(name, type)
1361*c217d954SCole Faust
1362*c217d954SCole Faust#define TENSOR4D_WO_T_IMAGE(name)          \
1363*c217d954SCole Faust    __write_only image2d_t name##_img, \
1364*c217d954SCole Faust    TENSOR4D_T_BUFFER(name)
1365*c217d954SCole Faust
1366*c217d954SCole Faust#define TENSOR4D_WO_T_BUFFER(name) TENSOR4D_T_BUFFER(name)
1367*c217d954SCole Faust
1368*c217d954SCole Faust#define TENSOR4D_WO_T_STR(name, type) TENSOR4D_WO_T_##type(name)
1369*c217d954SCole Faust
1370*c217d954SCole Faust
1371*c217d954SCole Faust#define TENSOR4D_WO_T(name, type) TENSOR4D_WO_T_STR(name, type)
1372*c217d954SCole Faust
1373*c217d954SCole Faust#define TENSOR3D_T_IMAGE(name)          \
1374*c217d954SCole Faust    __read_only image2d_t name##_img, \
1375*c217d954SCole Faust    __global uchar *name##_ptr,       \
1376*c217d954SCole Faust    uint        name##_stride_y, \
1377*c217d954SCole Faust    uint        name##_stride_z, \
1378*c217d954SCole Faust    uint        name##_w,   \
1379*c217d954SCole Faust    uint        name##_h,   \
1380*c217d954SCole Faust    uint        name##_n,   \
1381*c217d954SCole Faust    uint        name##_offset_first_element_in_bytes
1382*c217d954SCole Faust
1383*c217d954SCole Faust#define TENSOR3D_T_BUFFER(name)    \
1384*c217d954SCole Faust    __global uchar *name##_ptr,  \
1385*c217d954SCole Faust    uint        name##_stride_y, \
1386*c217d954SCole Faust    uint        name##_stride_z, \
1387*c217d954SCole Faust    uint        name##_w,   \
1388*c217d954SCole Faust    uint        name##_h,   \
1389*c217d954SCole Faust    uint        name##_n,   \
1390*c217d954SCole Faust    uint        name##_offset_first_element_in_bytes
1391*c217d954SCole Faust
1392*c217d954SCole Faust#define TENSOR3D_T_STR(name, type) TENSOR3D_T_##type(name)
1393*c217d954SCole Faust#define TENSOR3D_T(name, type) TENSOR3D_T_STR(name, type)
1394*c217d954SCole Faust
1395*c217d954SCole Faust#if !defined(UNROLL_WITH_PRAGMA)
1396*c217d954SCole Faust#define UNROLL_INCR(idx, step, macro) idx += (step); (macro)
1397*c217d954SCole Faust
1398*c217d954SCole Faust#define LOOP_UNROLLING_1(idx, step, macro) (macro)
1399*c217d954SCole Faust#define LOOP_UNROLLING_2(idx, step, macro) LOOP_UNROLLING_1(idx, step, macro); UNROLL_INCR(idx, step, macro)
1400*c217d954SCole Faust#define LOOP_UNROLLING_3(idx, step, macro) LOOP_UNROLLING_2(idx, step, macro); UNROLL_INCR(idx, step, macro)
1401*c217d954SCole Faust#define LOOP_UNROLLING_4(idx, step, macro) LOOP_UNROLLING_3(idx, step, macro); UNROLL_INCR(idx, step, macro)
1402*c217d954SCole Faust#define LOOP_UNROLLING_5(idx, step, macro) LOOP_UNROLLING_4(idx, step, macro); UNROLL_INCR(idx, step, macro)
1403*c217d954SCole Faust#define LOOP_UNROLLING_6(idx, step, macro) LOOP_UNROLLING_5(idx, step, macro); UNROLL_INCR(idx, step, macro)
1404*c217d954SCole Faust#define LOOP_UNROLLING_7(idx, step, macro) LOOP_UNROLLING_6(idx, step, macro); UNROLL_INCR(idx, step, macro)
1405*c217d954SCole Faust#define LOOP_UNROLLING_8(idx, step, macro) LOOP_UNROLLING_7(idx, step, macro); UNROLL_INCR(idx, step, macro)
1406*c217d954SCole Faust#define LOOP_UNROLLING_9(idx, step, macro) LOOP_UNROLLING_8(idx, step, macro); UNROLL_INCR(idx, step, macro)
1407*c217d954SCole Faust#define LOOP_UNROLLING_10(idx, step, macro) LOOP_UNROLLING_9(idx, step, macro); UNROLL_INCR(idx, step, macro)
1408*c217d954SCole Faust#define LOOP_UNROLLING_11(idx, step, macro) LOOP_UNROLLING_10(idx, step, macro); UNROLL_INCR(idx, step, macro)
1409*c217d954SCole Faust#define LOOP_UNROLLING_12(idx, step, macro) LOOP_UNROLLING_11(idx, step, macro); UNROLL_INCR(idx, step, macro)
1410*c217d954SCole Faust#define LOOP_UNROLLING_13(idx, step, macro) LOOP_UNROLLING_12(idx, step, macro); UNROLL_INCR(idx, step, macro)
1411*c217d954SCole Faust#define LOOP_UNROLLING_14(idx, step, macro) LOOP_UNROLLING_13(idx, step, macro); UNROLL_INCR(idx, step, macro)
1412*c217d954SCole Faust#define LOOP_UNROLLING_15(idx, step, macro) LOOP_UNROLLING_14(idx, step, macro); UNROLL_INCR(idx, step, macro)
1413*c217d954SCole Faust#define LOOP_UNROLLING_16(idx, step, macro) LOOP_UNROLLING_15(idx, step, macro); UNROLL_INCR(idx, step, macro)
1414*c217d954SCole Faust#define LOOP_UNROLLING_17(idx, step, macro) LOOP_UNROLLING_16(idx, step, macro); UNROLL_INCR(idx, step, macro)
1415*c217d954SCole Faust#define LOOP_UNROLLING_18(idx, step, macro) LOOP_UNROLLING_17(idx, step, macro); UNROLL_INCR(idx, step, macro)
1416*c217d954SCole Faust#define LOOP_UNROLLING_19(idx, step, macro) LOOP_UNROLLING_18(idx, step, macro); UNROLL_INCR(idx, step, macro)
1417*c217d954SCole Faust#define LOOP_UNROLLING_20(idx, step, macro) LOOP_UNROLLING_19(idx, step, macro); UNROLL_INCR(idx, step, macro)
1418*c217d954SCole Faust#define LOOP_UNROLLING_21(idx, step, macro) LOOP_UNROLLING_20(idx, step, macro); UNROLL_INCR(idx, step, macro)
1419*c217d954SCole Faust#define LOOP_UNROLLING_22(idx, step, macro) LOOP_UNROLLING_21(idx, step, macro); UNROLL_INCR(idx, step, macro)
1420*c217d954SCole Faust#define LOOP_UNROLLING_23(idx, step, macro) LOOP_UNROLLING_22(idx, step, macro); UNROLL_INCR(idx, step, macro)
1421*c217d954SCole Faust#define LOOP_UNROLLING_24(idx, step, macro) LOOP_UNROLLING_23(idx, step, macro); UNROLL_INCR(idx, step, macro)
1422*c217d954SCole Faust#define LOOP_UNROLLING_25(idx, step, macro) LOOP_UNROLLING_24(idx, step, macro); UNROLL_INCR(idx, step, macro)
1423*c217d954SCole Faust#define LOOP_UNROLLING_26(idx, step, macro) LOOP_UNROLLING_25(idx, step, macro); UNROLL_INCR(idx, step, macro)
1424*c217d954SCole Faust#define LOOP_UNROLLING_27(idx, step, macro) LOOP_UNROLLING_26(idx, step, macro); UNROLL_INCR(idx, step, macro)
1425*c217d954SCole Faust#define LOOP_UNROLLING_28(idx, step, macro) LOOP_UNROLLING_27(idx, step, macro); UNROLL_INCR(idx, step, macro)
1426*c217d954SCole Faust#define LOOP_UNROLLING_29(idx, step, macro) LOOP_UNROLLING_28(idx, step, macro); UNROLL_INCR(idx, step, macro)
1427*c217d954SCole Faust#define LOOP_UNROLLING_30(idx, step, macro) LOOP_UNROLLING_29(idx, step, macro); UNROLL_INCR(idx, step, macro)
1428*c217d954SCole Faust#define LOOP_UNROLLING_31(idx, step, macro) LOOP_UNROLLING_30(idx, step, macro); UNROLL_INCR(idx, step, macro)
1429*c217d954SCole Faust#define LOOP_UNROLLING_32(idx, step, macro) LOOP_UNROLLING_31(idx, step, macro); UNROLL_INCR(idx, step, macro)
1430*c217d954SCole Faust#define LOOP_UNROLLING_33(idx, step, macro) LOOP_UNROLLING_32(idx, step, macro); UNROLL_INCR(idx, step, macro)
1431*c217d954SCole Faust#define LOOP_UNROLLING_34(idx, step, macro) LOOP_UNROLLING_33(idx, step, macro); UNROLL_INCR(idx, step, macro)
1432*c217d954SCole Faust#define LOOP_UNROLLING_35(idx, step, macro) LOOP_UNROLLING_34(idx, step, macro); UNROLL_INCR(idx, step, macro)
1433*c217d954SCole Faust#define LOOP_UNROLLING_36(idx, step, macro) LOOP_UNROLLING_35(idx, step, macro); UNROLL_INCR(idx, step, macro)
1434*c217d954SCole Faust#define LOOP_UNROLLING_37(idx, step, macro) LOOP_UNROLLING_36(idx, step, macro); UNROLL_INCR(idx, step, macro)
1435*c217d954SCole Faust#define LOOP_UNROLLING_38(idx, step, macro) LOOP_UNROLLING_37(idx, step, macro); UNROLL_INCR(idx, step, macro)
1436*c217d954SCole Faust#define LOOP_UNROLLING_39(idx, step, macro) LOOP_UNROLLING_38(idx, step, macro); UNROLL_INCR(idx, step, macro)
1437*c217d954SCole Faust#define LOOP_UNROLLING_40(idx, step, macro) LOOP_UNROLLING_39(idx, step, macro); UNROLL_INCR(idx, step, macro)
1438*c217d954SCole Faust#define LOOP_UNROLLING_41(idx, step, macro) LOOP_UNROLLING_40(idx, step, macro); UNROLL_INCR(idx, step, macro)
1439*c217d954SCole Faust#define LOOP_UNROLLING_42(idx, step, macro) LOOP_UNROLLING_41(idx, step, macro); UNROLL_INCR(idx, step, macro)
1440*c217d954SCole Faust#define LOOP_UNROLLING_43(idx, step, macro) LOOP_UNROLLING_42(idx, step, macro); UNROLL_INCR(idx, step, macro)
1441*c217d954SCole Faust#define LOOP_UNROLLING_44(idx, step, macro) LOOP_UNROLLING_43(idx, step, macro); UNROLL_INCR(idx, step, macro)
1442*c217d954SCole Faust#define LOOP_UNROLLING_45(idx, step, macro) LOOP_UNROLLING_44(idx, step, macro); UNROLL_INCR(idx, step, macro)
1443*c217d954SCole Faust#define LOOP_UNROLLING_46(idx, step, macro) LOOP_UNROLLING_45(idx, step, macro); UNROLL_INCR(idx, step, macro)
1444*c217d954SCole Faust#define LOOP_UNROLLING_47(idx, step, macro) LOOP_UNROLLING_46(idx, step, macro); UNROLL_INCR(idx, step, macro)
1445*c217d954SCole Faust#define LOOP_UNROLLING_48(idx, step, macro) LOOP_UNROLLING_47(idx, step, macro); UNROLL_INCR(idx, step, macro)
1446*c217d954SCole Faust#define LOOP_UNROLLING_49(idx, step, macro) LOOP_UNROLLING_48(idx, step, macro); UNROLL_INCR(idx, step, macro)
1447*c217d954SCole Faust#define LOOP_UNROLLING_50(idx, step, macro) LOOP_UNROLLING_49(idx, step, macro); UNROLL_INCR(idx, step, macro)
1448*c217d954SCole Faust#define LOOP_UNROLLING_51(idx, step, macro) LOOP_UNROLLING_50(idx, step, macro); UNROLL_INCR(idx, step, macro)
1449*c217d954SCole Faust#define LOOP_UNROLLING_52(idx, step, macro) LOOP_UNROLLING_51(idx, step, macro); UNROLL_INCR(idx, step, macro)
1450*c217d954SCole Faust#define LOOP_UNROLLING_53(idx, step, macro) LOOP_UNROLLING_52(idx, step, macro); UNROLL_INCR(idx, step, macro)
1451*c217d954SCole Faust#define LOOP_UNROLLING_54(idx, step, macro) LOOP_UNROLLING_53(idx, step, macro); UNROLL_INCR(idx, step, macro)
1452*c217d954SCole Faust#define LOOP_UNROLLING_55(idx, step, macro) LOOP_UNROLLING_54(idx, step, macro); UNROLL_INCR(idx, step, macro)
1453*c217d954SCole Faust#define LOOP_UNROLLING_56(idx, step, macro) LOOP_UNROLLING_55(idx, step, macro); UNROLL_INCR(idx, step, macro)
1454*c217d954SCole Faust#define LOOP_UNROLLING_57(idx, step, macro) LOOP_UNROLLING_56(idx, step, macro); UNROLL_INCR(idx, step, macro)
1455*c217d954SCole Faust#define LOOP_UNROLLING_58(idx, step, macro) LOOP_UNROLLING_57(idx, step, macro); UNROLL_INCR(idx, step, macro)
1456*c217d954SCole Faust#define LOOP_UNROLLING_59(idx, step, macro) LOOP_UNROLLING_58(idx, step, macro); UNROLL_INCR(idx, step, macro)
1457*c217d954SCole Faust#define LOOP_UNROLLING_60(idx, step, macro) LOOP_UNROLLING_59(idx, step, macro); UNROLL_INCR(idx, step, macro)
1458*c217d954SCole Faust#define LOOP_UNROLLING_61(idx, step, macro) LOOP_UNROLLING_60(idx, step, macro); UNROLL_INCR(idx, step, macro)
1459*c217d954SCole Faust#define LOOP_UNROLLING_62(idx, step, macro) LOOP_UNROLLING_61(idx, step, macro); UNROLL_INCR(idx, step, macro)
1460*c217d954SCole Faust#define LOOP_UNROLLING_63(idx, step, macro) LOOP_UNROLLING_62(idx, step, macro); UNROLL_INCR(idx, step, macro)
1461*c217d954SCole Faust#define LOOP_UNROLLING_64(idx, step, macro) LOOP_UNROLLING_63(idx, step, macro); UNROLL_INCR(idx, step, macro)
1462*c217d954SCole Faust#define LOOP_UNROLLING_65(idx, step, macro) LOOP_UNROLLING_64(idx, step, macro); UNROLL_INCR(idx, step, macro)
1463*c217d954SCole Faust#define LOOP_UNROLLING_66(idx, step, macro) LOOP_UNROLLING_65(idx, step, macro); UNROLL_INCR(idx, step, macro)
1464*c217d954SCole Faust#define LOOP_UNROLLING_67(idx, step, macro) LOOP_UNROLLING_66(idx, step, macro); UNROLL_INCR(idx, step, macro)
1465*c217d954SCole Faust#define LOOP_UNROLLING_68(idx, step, macro) LOOP_UNROLLING_67(idx, step, macro); UNROLL_INCR(idx, step, macro)
1466*c217d954SCole Faust#define LOOP_UNROLLING_69(idx, step, macro) LOOP_UNROLLING_68(idx, step, macro); UNROLL_INCR(idx, step, macro)
1467*c217d954SCole Faust#define LOOP_UNROLLING_70(idx, step, macro) LOOP_UNROLLING_69(idx, step, macro); UNROLL_INCR(idx, step, macro)
1468*c217d954SCole Faust#define LOOP_UNROLLING_71(idx, step, macro) LOOP_UNROLLING_70(idx, step, macro); UNROLL_INCR(idx, step, macro)
1469*c217d954SCole Faust#define LOOP_UNROLLING_72(idx, step, macro) LOOP_UNROLLING_71(idx, step, macro); UNROLL_INCR(idx, step, macro)
1470*c217d954SCole Faust#define LOOP_UNROLLING_73(idx, step, macro) LOOP_UNROLLING_72(idx, step, macro); UNROLL_INCR(idx, step, macro)
1471*c217d954SCole Faust#define LOOP_UNROLLING_74(idx, step, macro) LOOP_UNROLLING_73(idx, step, macro); UNROLL_INCR(idx, step, macro)
1472*c217d954SCole Faust#define LOOP_UNROLLING_75(idx, step, macro) LOOP_UNROLLING_74(idx, step, macro); UNROLL_INCR(idx, step, macro)
1473*c217d954SCole Faust#define LOOP_UNROLLING_76(idx, step, macro) LOOP_UNROLLING_75(idx, step, macro); UNROLL_INCR(idx, step, macro)
1474*c217d954SCole Faust#define LOOP_UNROLLING_77(idx, step, macro) LOOP_UNROLLING_76(idx, step, macro); UNROLL_INCR(idx, step, macro)
1475*c217d954SCole Faust#define LOOP_UNROLLING_78(idx, step, macro) LOOP_UNROLLING_77(idx, step, macro); UNROLL_INCR(idx, step, macro)
1476*c217d954SCole Faust#define LOOP_UNROLLING_79(idx, step, macro) LOOP_UNROLLING_78(idx, step, macro); UNROLL_INCR(idx, step, macro)
1477*c217d954SCole Faust#define LOOP_UNROLLING_80(idx, step, macro) LOOP_UNROLLING_79(idx, step, macro); UNROLL_INCR(idx, step, macro)
1478*c217d954SCole Faust#define LOOP_UNROLLING_81(idx, step, macro) LOOP_UNROLLING_80(idx, step, macro); UNROLL_INCR(idx, step, macro)
1479*c217d954SCole Faust#define LOOP_UNROLLING_82(idx, step, macro) LOOP_UNROLLING_81(idx, step, macro); UNROLL_INCR(idx, step, macro)
1480*c217d954SCole Faust#define LOOP_UNROLLING_83(idx, step, macro) LOOP_UNROLLING_82(idx, step, macro); UNROLL_INCR(idx, step, macro)
1481*c217d954SCole Faust#define LOOP_UNROLLING_84(idx, step, macro) LOOP_UNROLLING_83(idx, step, macro); UNROLL_INCR(idx, step, macro)
1482*c217d954SCole Faust#define LOOP_UNROLLING_85(idx, step, macro) LOOP_UNROLLING_84(idx, step, macro); UNROLL_INCR(idx, step, macro)
1483*c217d954SCole Faust#define LOOP_UNROLLING_86(idx, step, macro) LOOP_UNROLLING_85(idx, step, macro); UNROLL_INCR(idx, step, macro)
1484*c217d954SCole Faust#define LOOP_UNROLLING_87(idx, step, macro) LOOP_UNROLLING_86(idx, step, macro); UNROLL_INCR(idx, step, macro)
1485*c217d954SCole Faust#define LOOP_UNROLLING_88(idx, step, macro) LOOP_UNROLLING_87(idx, step, macro); UNROLL_INCR(idx, step, macro)
1486*c217d954SCole Faust#define LOOP_UNROLLING_89(idx, step, macro) LOOP_UNROLLING_88(idx, step, macro); UNROLL_INCR(idx, step, macro)
1487*c217d954SCole Faust#define LOOP_UNROLLING_90(idx, step, macro) LOOP_UNROLLING_89(idx, step, macro); UNROLL_INCR(idx, step, macro)
1488*c217d954SCole Faust#define LOOP_UNROLLING_91(idx, step, macro) LOOP_UNROLLING_90(idx, step, macro); UNROLL_INCR(idx, step, macro)
1489*c217d954SCole Faust#define LOOP_UNROLLING_92(idx, step, macro) LOOP_UNROLLING_91(idx, step, macro); UNROLL_INCR(idx, step, macro)
1490*c217d954SCole Faust#define LOOP_UNROLLING_93(idx, step, macro) LOOP_UNROLLING_92(idx, step, macro); UNROLL_INCR(idx, step, macro)
1491*c217d954SCole Faust#define LOOP_UNROLLING_94(idx, step, macro) LOOP_UNROLLING_93(idx, step, macro); UNROLL_INCR(idx, step, macro)
1492*c217d954SCole Faust#define LOOP_UNROLLING_95(idx, step, macro) LOOP_UNROLLING_94(idx, step, macro); UNROLL_INCR(idx, step, macro)
1493*c217d954SCole Faust#define LOOP_UNROLLING_96(idx, step, macro) LOOP_UNROLLING_95(idx, step, macro); UNROLL_INCR(idx, step, macro)
1494*c217d954SCole Faust#define LOOP_UNROLLING_97(idx, step, macro) LOOP_UNROLLING_96(idx, step, macro); UNROLL_INCR(idx, step, macro)
1495*c217d954SCole Faust#define LOOP_UNROLLING_98(idx, step, macro) LOOP_UNROLLING_97(idx, step, macro); UNROLL_INCR(idx, step, macro)
1496*c217d954SCole Faust#define LOOP_UNROLLING_99(idx, step, macro) LOOP_UNROLLING_98(idx, step, macro); UNROLL_INCR(idx, step, macro)
1497*c217d954SCole Faust#define LOOP_UNROLLING_100(idx, step, macro) LOOP_UNROLLING_99(idx, step, macro); UNROLL_INCR(idx, step, macro)
1498*c217d954SCole Faust#define LOOP_UNROLLING_101(idx, step, macro) LOOP_UNROLLING_100(idx, step, macro); UNROLL_INCR(idx, step, macro)
1499*c217d954SCole Faust#define LOOP_UNROLLING_102(idx, step, macro) LOOP_UNROLLING_101(idx, step, macro); UNROLL_INCR(idx, step, macro)
1500*c217d954SCole Faust#define LOOP_UNROLLING_103(idx, step, macro) LOOP_UNROLLING_102(idx, step, macro); UNROLL_INCR(idx, step, macro)
1501*c217d954SCole Faust#define LOOP_UNROLLING_104(idx, step, macro) LOOP_UNROLLING_103(idx, step, macro); UNROLL_INCR(idx, step, macro)
1502*c217d954SCole Faust#define LOOP_UNROLLING_105(idx, step, macro) LOOP_UNROLLING_104(idx, step, macro); UNROLL_INCR(idx, step, macro)
1503*c217d954SCole Faust#define LOOP_UNROLLING_106(idx, step, macro) LOOP_UNROLLING_105(idx, step, macro); UNROLL_INCR(idx, step, macro)
1504*c217d954SCole Faust#define LOOP_UNROLLING_107(idx, step, macro) LOOP_UNROLLING_106(idx, step, macro); UNROLL_INCR(idx, step, macro)
1505*c217d954SCole Faust#define LOOP_UNROLLING_108(idx, step, macro) LOOP_UNROLLING_107(idx, step, macro); UNROLL_INCR(idx, step, macro)
1506*c217d954SCole Faust#define LOOP_UNROLLING_109(idx, step, macro) LOOP_UNROLLING_108(idx, step, macro); UNROLL_INCR(idx, step, macro)
1507*c217d954SCole Faust#define LOOP_UNROLLING_110(idx, step, macro) LOOP_UNROLLING_109(idx, step, macro); UNROLL_INCR(idx, step, macro)
1508*c217d954SCole Faust#define LOOP_UNROLLING_111(idx, step, macro) LOOP_UNROLLING_110(idx, step, macro); UNROLL_INCR(idx, step, macro)
1509*c217d954SCole Faust#define LOOP_UNROLLING_112(idx, step, macro) LOOP_UNROLLING_111(idx, step, macro); UNROLL_INCR(idx, step, macro)
1510*c217d954SCole Faust#define LOOP_UNROLLING_113(idx, step, macro) LOOP_UNROLLING_112(idx, step, macro); UNROLL_INCR(idx, step, macro)
1511*c217d954SCole Faust#define LOOP_UNROLLING_114(idx, step, macro) LOOP_UNROLLING_113(idx, step, macro); UNROLL_INCR(idx, step, macro)
1512*c217d954SCole Faust#define LOOP_UNROLLING_115(idx, step, macro) LOOP_UNROLLING_114(idx, step, macro); UNROLL_INCR(idx, step, macro)
1513*c217d954SCole Faust#define LOOP_UNROLLING_116(idx, step, macro) LOOP_UNROLLING_115(idx, step, macro); UNROLL_INCR(idx, step, macro)
1514*c217d954SCole Faust#define LOOP_UNROLLING_117(idx, step, macro) LOOP_UNROLLING_116(idx, step, macro); UNROLL_INCR(idx, step, macro)
1515*c217d954SCole Faust#define LOOP_UNROLLING_118(idx, step, macro) LOOP_UNROLLING_117(idx, step, macro); UNROLL_INCR(idx, step, macro)
1516*c217d954SCole Faust#define LOOP_UNROLLING_119(idx, step, macro) LOOP_UNROLLING_118(idx, step, macro); UNROLL_INCR(idx, step, macro)
1517*c217d954SCole Faust#define LOOP_UNROLLING_120(idx, step, macro) LOOP_UNROLLING_119(idx, step, macro); UNROLL_INCR(idx, step, macro)
1518*c217d954SCole Faust#define LOOP_UNROLLING_121(idx, step, macro) LOOP_UNROLLING_120(idx, step, macro); UNROLL_INCR(idx, step, macro)
1519*c217d954SCole Faust#define LOOP_UNROLLING_122(idx, step, macro) LOOP_UNROLLING_121(idx, step, macro); UNROLL_INCR(idx, step, macro)
1520*c217d954SCole Faust#define LOOP_UNROLLING_123(idx, step, macro) LOOP_UNROLLING_122(idx, step, macro); UNROLL_INCR(idx, step, macro)
1521*c217d954SCole Faust#define LOOP_UNROLLING_124(idx, step, macro) LOOP_UNROLLING_123(idx, step, macro); UNROLL_INCR(idx, step, macro)
1522*c217d954SCole Faust#define LOOP_UNROLLING_125(idx, step, macro) LOOP_UNROLLING_124(idx, step, macro); UNROLL_INCR(idx, step, macro)
1523*c217d954SCole Faust#define LOOP_UNROLLING_126(idx, step, macro) LOOP_UNROLLING_125(idx, step, macro); UNROLL_INCR(idx, step, macro)
1524*c217d954SCole Faust#define LOOP_UNROLLING_127(idx, step, macro) LOOP_UNROLLING_126(idx, step, macro); UNROLL_INCR(idx, step, macro)
1525*c217d954SCole Faust#define LOOP_UNROLLING_128(idx, step, macro) LOOP_UNROLLING_127(idx, step, macro); UNROLL_INCR(idx, step, macro)
1526*c217d954SCole Faust
1527*c217d954SCole Faust#define LOOP_UNROLLING_STR(type, idx, start, step, num, macro) \
1528*c217d954SCole Faust    {                                                          \
1529*c217d954SCole Faust        type idx = start;                                      \
1530*c217d954SCole Faust        LOOP_UNROLLING_##num(idx, step, macro);                \
1531*c217d954SCole Faust    }
1532*c217d954SCole Faust#else
1533*c217d954SCole Faust#define LOOP_UNROLLING_STR(type, idx, start, step, num, macro) \
1534*c217d954SCole Faust    {                                                          \
1535*c217d954SCole Faust        _Pragma("unroll")                                      \
1536*c217d954SCole Faust        for(type idx = start; idx < (num * step); idx += step) \
1537*c217d954SCole Faust        {                                                      \
1538*c217d954SCole Faust            (macro);                                           \
1539*c217d954SCole Faust        }                                                      \
1540*c217d954SCole Faust    }
1541*c217d954SCole Faust#endif
1542*c217d954SCole Faust#define LOOP_UNROLLING(type, idx, start, step, num, macro) LOOP_UNROLLING_STR(type, idx, start, step, num, macro)
1543*c217d954SCole Faust
1544*c217d954SCole Faust
1545*c217d954SCole Faust#define GET_SPATIAL_IDX(IDX, N0, PARTIAL_N0) (max((int)(get_global_id(IDX) * N0 - (N0 - PARTIAL_N0) % N0), 0))
1546*c217d954SCole Faust
1547*c217d954SCole Faust
1548*c217d954SCole Faust#define DOT_PRODUCT_INTEGER8(A_DATA_TYPE, B_DATA_TYPE, C_DATA_TYPE, K0, a, b, c) DOT_PRODUCT_INTEGER8_STR(A_DATA_TYPE, B_DATA_TYPE, C_DATA_TYPE, K0, a, b, c)
1549*c217d954SCole Faust#define DOT_PRODUCT_INTEGER8_STR(A_DATA_TYPE, B_DATA_TYPE, C_DATA_TYPE, K0, a, b, c) DOT_PRODUCT##K0##_INTEGER8(A_DATA_TYPE, B_DATA_TYPE, C_DATA_TYPE, a, b, c)
1550*c217d954SCole Faust#define DOT_PRODUCT1_INTEGER8(A_DATA_TYPE, B_DATA_TYPE, C_DATA_TYPE, a, b, c) \
1551*c217d954SCole Faust    ({                                                \
1552*c217d954SCole Faust        c += (C_DATA_TYPE)(a) * (C_DATA_TYPE)(b);     \
1553*c217d954SCole Faust    })
1554*c217d954SCole Faust#if defined(ARM_COMPUTE_OPENCL_DOT8_ENABLED) && defined(cl_khr_integer_dot_product)
1555*c217d954SCole Faust#define DOT_PRODUCT2_INTEGER8(A_DATA_TYPE, B_DATA_TYPE, C_DATA_TYPE, a, b, c) c += dot((A_DATA_TYPE##4)((a).s01, (A_DATA_TYPE##2)(0)), (B_DATA_TYPE##4)(((b).s01), (B_DATA_TYPE##2)(0)));
1556*c217d954SCole Faust#define DOT_PRODUCT3_INTEGER8(A_DATA_TYPE, B_DATA_TYPE, C_DATA_TYPE, a, b, c) c += dot((A_DATA_TYPE##4)((a).s012, (A_DATA_TYPE)0), (B_DATA_TYPE##4)(((b).s012), (B_DATA_TYPE)0));
1557*c217d954SCole Faust#define DOT_PRODUCT4_INTEGER8(A_DATA_TYPE, B_DATA_TYPE, C_DATA_TYPE, a, b, c) c += dot((a), (b));
1558*c217d954SCole Faust#elif defined(ARM_COMPUTE_OPENCL_DOT8_ACC_ENABLED) && defined(cl_arm_integer_dot_product_accumulate_int8)
1559*c217d954SCole Faust#define DOT_PRODUCT2_INTEGER8(A_DATA_TYPE, B_DATA_TYPE, C_DATA_TYPE, a, b, c) c = arm_dot_acc((A_DATA_TYPE##4)((a).s01, (A_DATA_TYPE##2)(0)), (B_DATA_TYPE##4)(((b).s01), (B_DATA_TYPE##2)(0)), (c));
1560*c217d954SCole Faust#define DOT_PRODUCT3_INTEGER8(A_DATA_TYPE, B_DATA_TYPE, C_DATA_TYPE, a, b, c) c = arm_dot_acc((A_DATA_TYPE##4)((a).s012, (A_DATA_TYPE)0), (B_DATA_TYPE##4)(((b).s012), (B_DATA_TYPE)0), (c));
1561*c217d954SCole Faust#define DOT_PRODUCT4_INTEGER8(A_DATA_TYPE, B_DATA_TYPE, C_DATA_TYPE, a, b, c) c = arm_dot_acc((a), (b), (c));
1562*c217d954SCole Faust#elif defined(ARM_COMPUTE_OPENCL_DOT8_ENABLED) && defined(cl_arm_integer_dot_product_int8)
1563*c217d954SCole Faust#define DOT_PRODUCT2_INTEGER8(A_DATA_TYPE, B_DATA_TYPE, C_DATA_TYPE, a, b, c) c += arm_dot((A_DATA_TYPE##4)((a).s01, (A_DATA_TYPE##2)(0)), (B_DATA_TYPE##4)(((b).s01), (B_DATA_TYPE##2)(0)));
1564*c217d954SCole Faust#define DOT_PRODUCT3_INTEGER8(A_DATA_TYPE, B_DATA_TYPE, C_DATA_TYPE, a, b, c) c += arm_dot((A_DATA_TYPE##4)((a).s012, (A_DATA_TYPE)0), (B_DATA_TYPE##4)(((b).s012), (B_DATA_TYPE)0));
1565*c217d954SCole Faust#define DOT_PRODUCT4_INTEGER8(A_DATA_TYPE, B_DATA_TYPE, C_DATA_TYPE, a, b, c) c += arm_dot((a), (b));
1566*c217d954SCole Faust#else
1567*c217d954SCole Faust#define DOT_PRODUCT2_INTEGER8(A_DATA_TYPE, B_DATA_TYPE, C_DATA_TYPE, a, b, c)   \
1568*c217d954SCole Faust    ({                                                  \
1569*c217d954SCole Faust        c += (C_DATA_TYPE)(a).s0 * (C_DATA_TYPE)(b).s0; \
1570*c217d954SCole Faust        c += (C_DATA_TYPE)(a).s1 * (C_DATA_TYPE)(b).s1; \
1571*c217d954SCole Faust    })
1572*c217d954SCole Faust#define DOT_PRODUCT3_INTEGER8(A_DATA_TYPE, B_DATA_TYPE, C_DATA_TYPE, a, b, c)   \
1573*c217d954SCole Faust    ({                                                  \
1574*c217d954SCole Faust        DOT_PRODUCT2_INTEGER8(A_DATA_TYPE, B_DATA_TYPE, C_DATA_TYPE, a, b, c);  \
1575*c217d954SCole Faust        c += (C_DATA_TYPE)(a).s2 * (C_DATA_TYPE)(b).s2; \
1576*c217d954SCole Faust    })
1577*c217d954SCole Faust#define DOT_PRODUCT4_INTEGER8(A_DATA_TYPE, B_DATA_TYPE, C_DATA_TYPE, x, y, val)   \
1578*c217d954SCole Faust    ({                                                    \
1579*c217d954SCole Faust        val += (C_DATA_TYPE)(x).s0 * (C_DATA_TYPE)(y).s0; \
1580*c217d954SCole Faust        val += (C_DATA_TYPE)(x).s1 * (C_DATA_TYPE)(y).s1; \
1581*c217d954SCole Faust        val += (C_DATA_TYPE)(x).s2 * (C_DATA_TYPE)(y).s2; \
1582*c217d954SCole Faust        val += (C_DATA_TYPE)(x).s3 * (C_DATA_TYPE)(y).s3; \
1583*c217d954SCole Faust    })
1584*c217d954SCole Faust#endif
1585*c217d954SCole Faust#define DOT_PRODUCT5_INTEGER8(A_DATA_TYPE, B_DATA_TYPE, C_DATA_TYPE, a, b, c) \
1586*c217d954SCole Faust    ({                                                \
1587*c217d954SCole Faust        DOT_PRODUCT4_INTEGER8(A_DATA_TYPE, B_DATA_TYPE, C_DATA_TYPE, ((a).s0123), ((b).s0123), c);     \
1588*c217d954SCole Faust        DOT_PRODUCT1_INTEGER8(A_DATA_TYPE, B_DATA_TYPE, C_DATA_TYPE, ((a).s4), ((b).s4), c);     \
1589*c217d954SCole Faust    })
1590*c217d954SCole Faust#define DOT_PRODUCT6_INTEGER8(A_DATA_TYPE, B_DATA_TYPE, C_DATA_TYPE, a, b, c) \
1591*c217d954SCole Faust    ({                                                \
1592*c217d954SCole Faust        DOT_PRODUCT4_INTEGER8(A_DATA_TYPE, B_DATA_TYPE, C_DATA_TYPE, ((a).s0123), ((b).s0123), c);     \
1593*c217d954SCole Faust        DOT_PRODUCT2_INTEGER8(A_DATA_TYPE, B_DATA_TYPE, C_DATA_TYPE, ((a).s45), ((b).s45), c);     \
1594*c217d954SCole Faust    })
1595*c217d954SCole Faust#define DOT_PRODUCT7_INTEGER8(A_DATA_TYPE, B_DATA_TYPE, C_DATA_TYPE, a, b, c) \
1596*c217d954SCole Faust    ({                                                \
1597*c217d954SCole Faust        DOT_PRODUCT4_INTEGER8(A_DATA_TYPE, B_DATA_TYPE, C_DATA_TYPE, ((a).s0123), ((b).s0123), c);     \
1598*c217d954SCole Faust        DOT_PRODUCT3_INTEGER8(A_DATA_TYPE, B_DATA_TYPE, C_DATA_TYPE, ((a).s456), ((b).s456), c);     \
1599*c217d954SCole Faust    })
1600*c217d954SCole Faust#define DOT_PRODUCT8_INTEGER8(A_DATA_TYPE, B_DATA_TYPE, C_DATA_TYPE, a, b, c) \
1601*c217d954SCole Faust    ({                                                \
1602*c217d954SCole Faust        DOT_PRODUCT4_INTEGER8(A_DATA_TYPE, B_DATA_TYPE, C_DATA_TYPE, ((a).lo), ((b).lo), c);     \
1603*c217d954SCole Faust        DOT_PRODUCT4_INTEGER8(A_DATA_TYPE, B_DATA_TYPE, C_DATA_TYPE, ((a).hi), ((b).hi), c);     \
1604*c217d954SCole Faust    })
1605*c217d954SCole Faust#define DOT_PRODUCT9_INTEGER8(A_DATA_TYPE, B_DATA_TYPE, C_DATA_TYPE, a, b, c) \
1606*c217d954SCole Faust    ({                                                \
1607*c217d954SCole Faust        DOT_PRODUCT8_INTEGER8(A_DATA_TYPE, B_DATA_TYPE, C_DATA_TYPE, ((a).s01234567), ((b).s01234567), c);     \
1608*c217d954SCole Faust        DOT_PRODUCT1_INTEGER8(A_DATA_TYPE, B_DATA_TYPE, C_DATA_TYPE, ((a).s8), ((b).s8), c);     \
1609*c217d954SCole Faust    })
1610*c217d954SCole Faust#define DOT_PRODUCT10_INTEGER8(A_DATA_TYPE, B_DATA_TYPE, C_DATA_TYPE, a, b, c) \
1611*c217d954SCole Faust    ({                                                \
1612*c217d954SCole Faust        DOT_PRODUCT8_INTEGER8(A_DATA_TYPE, B_DATA_TYPE, C_DATA_TYPE, ((a).s01234567), ((b).s01234567), c);     \
1613*c217d954SCole Faust        DOT_PRODUCT2_INTEGER8(A_DATA_TYPE, B_DATA_TYPE, C_DATA_TYPE, ((a).s89), ((b).s89), c);     \
1614*c217d954SCole Faust    })
1615*c217d954SCole Faust#define DOT_PRODUCT11_INTEGER8(A_DATA_TYPE, B_DATA_TYPE, C_DATA_TYPE, a, b, c) \
1616*c217d954SCole Faust    ({                                                \
1617*c217d954SCole Faust        DOT_PRODUCT8_INTEGER8(A_DATA_TYPE, B_DATA_TYPE, C_DATA_TYPE, ((a).s01234567), ((b).s01234567), c);     \
1618*c217d954SCole Faust        DOT_PRODUCT3_INTEGER8(A_DATA_TYPE, B_DATA_TYPE, C_DATA_TYPE, ((a).s89A), ((b).s89A), c);     \
1619*c217d954SCole Faust    })
1620*c217d954SCole Faust#define DOT_PRODUCT12_INTEGER8(A_DATA_TYPE, B_DATA_TYPE, C_DATA_TYPE, a, b, c) \
1621*c217d954SCole Faust    ({                                                \
1622*c217d954SCole Faust        DOT_PRODUCT8_INTEGER8(A_DATA_TYPE, B_DATA_TYPE, C_DATA_TYPE, ((a).s01234567), ((b).s01234567), c);     \
1623*c217d954SCole Faust        DOT_PRODUCT4_INTEGER8(A_DATA_TYPE, B_DATA_TYPE, C_DATA_TYPE, ((a).s89AB), ((b).s89AB), c);     \
1624*c217d954SCole Faust    })
1625*c217d954SCole Faust#define DOT_PRODUCT13_INTEGER8(A_DATA_TYPE, B_DATA_TYPE, C_DATA_TYPE, a, b, c) \
1626*c217d954SCole Faust    ({                                                \
1627*c217d954SCole Faust        DOT_PRODUCT8_INTEGER8(A_DATA_TYPE, B_DATA_TYPE, C_DATA_TYPE, ((a).s01234567), ((b).s01234567), c);     \
1628*c217d954SCole Faust        DOT_PRODUCT5_INTEGER8(A_DATA_TYPE, B_DATA_TYPE, C_DATA_TYPE, ((a).s89ABC), ((b).s89ABC), c);     \
1629*c217d954SCole Faust    })
1630*c217d954SCole Faust#define DOT_PRODUCT14_INTEGER8(A_DATA_TYPE, B_DATA_TYPE, C_DATA_TYPE, a, b, c) \
1631*c217d954SCole Faust    ({                                                \
1632*c217d954SCole Faust        DOT_PRODUCT8_INTEGER8(A_DATA_TYPE, B_DATA_TYPE, C_DATA_TYPE, ((a).s01234567), ((b).s01234567), c);     \
1633*c217d954SCole Faust        DOT_PRODUCT6_INTEGER8(A_DATA_TYPE, B_DATA_TYPE, C_DATA_TYPE, ((a).s89ABCD), ((b).s89ABCD), c);     \
1634*c217d954SCole Faust    })
1635*c217d954SCole Faust#define DOT_PRODUCT15_INTEGER8(A_DATA_TYPE, B_DATA_TYPE, C_DATA_TYPE, a, b, c) \
1636*c217d954SCole Faust    ({                                                \
1637*c217d954SCole Faust        DOT_PRODUCT8_INTEGER8(A_DATA_TYPE, B_DATA_TYPE, C_DATA_TYPE, ((a).s01234567), ((b).s01234567), c);     \
1638*c217d954SCole Faust        DOT_PRODUCT7_INTEGER8(A_DATA_TYPE, B_DATA_TYPE, C_DATA_TYPE, ((a).s89ABCDE), ((b).s89ABCDE), c);     \
1639*c217d954SCole Faust    })
1640*c217d954SCole Faust#define DOT_PRODUCT16_INTEGER8(A_DATA_TYPE, B_DATA_TYPE, C_DATA_TYPE, a, b, c) \
1641*c217d954SCole Faust    ({                                                 \
1642*c217d954SCole Faust        DOT_PRODUCT8_INTEGER8(A_DATA_TYPE, B_DATA_TYPE, C_DATA_TYPE, ((a).lo), ((b).lo), c);      \
1643*c217d954SCole Faust        DOT_PRODUCT8_INTEGER8(A_DATA_TYPE, B_DATA_TYPE, C_DATA_TYPE, ((a).hi), ((b).hi), c);      \
1644*c217d954SCole Faust    })
1645*c217d954SCole Faust
1646*c217d954SCole Faust
1647*c217d954SCole Faust#define REDUCE_INTEGER8(A_DATA_TYPE, B_DATA_TYPE, C_DATA_TYPE, K0, a, c) REDUCE_INTEGER8_STR(A_DATA_TYPE, B_DATA_TYPE, C_DATA_TYPE, K0, a, c)
1648*c217d954SCole Faust#define REDUCE_INTEGER8_STR(A_DATA_TYPE, B_DATA_TYPE, C_DATA_TYPE, K0, a, c) DOT_PRODUCT_INTEGER8(A_DATA_TYPE, B_DATA_TYPE, C_DATA_TYPE, K0, a, (TILE_VECTOR_TYPE##K0(B_DATA_TYPE))1, c)
1649*c217d954SCole Faust
1650*c217d954SCole Faust
1651*c217d954SCole Faust#define V_LOAD(DATA_TYPE, WIDTH, TENSOR_TYPE, TENSOR, X, Y, STRIDE_Y) V_LOAD_STR(DATA_TYPE, WIDTH, TENSOR_TYPE, TENSOR, X, Y, STRIDE_Y)
1652*c217d954SCole Faust#define V_LOAD_STR(DATA_TYPE, WIDTH, TENSOR_TYPE, TENSOR, X, Y, STRIDE_Y) V_LOAD_##TENSOR_TYPE(DATA_TYPE, WIDTH, TENSOR, X, Y, STRIDE_Y)
1653*c217d954SCole Faust#define V_LOAD_BUFFER(DATA_TYPE, WIDTH, TENSOR, X, Y, STRIDE_Y) \
1654*c217d954SCole Faust    VLOAD(WIDTH)                                                \
1655*c217d954SCole Faust    (0, (__global DATA_TYPE *)(TENSOR##_ptr + TENSOR##_offset_first_element_in_bytes + (X) * sizeof(DATA_TYPE) + (Y) * (STRIDE_Y)))
1656*c217d954SCole Faust#define V_LOAD_IMAGE(DATA_TYPE, WIDTH, TENSOR, X, Y, STRIDE_Y) READ_IMAGE2D(DATA_TYPE, CONVERT_VECTOR_SIZE_TO_PIXEL_UNIT(WIDTH), TENSOR##_img, (X) / 4, (Y))
1657*c217d954SCole Faust
1658*c217d954SCole Faust
1659*c217d954SCole Faust#define V_STORE(DATA_TYPE, WIDTH, TENSOR_TYPE, TENSOR, X, Y, STRIDE_Y, VALUES) V_STORE_STR(DATA_TYPE, WIDTH, TENSOR_TYPE, TENSOR, X, Y, STRIDE_Y, VALUES)
1660*c217d954SCole Faust#define V_STORE_STR(DATA_TYPE, WIDTH, TENSOR_TYPE, TENSOR, X, Y, STRIDE_Y, VALUES) V_STORE_##TENSOR_TYPE(DATA_TYPE, WIDTH, TENSOR, X, Y, STRIDE_Y, VALUES)
1661*c217d954SCole Faust#define V_STORE_BUFFER(DATA_TYPE, WIDTH, TENSOR, X, Y, STRIDE_Y, VALUES) \
1662*c217d954SCole Faust    VSTORE(WIDTH)                                                \
1663*c217d954SCole Faust    (VALUES, 0, (__global DATA_TYPE *)(TENSOR##_ptr + TENSOR##_offset_first_element_in_bytes + (X) * sizeof(DATA_TYPE) + (Y) * (STRIDE_Y)))
1664*c217d954SCole Faust#define V_STORE_IMAGE(DATA_TYPE, WIDTH, TENSOR, X, Y, STRIDE_Y, VALUES) WRITE_IMAGE2D(DATA_TYPE, CONVERT_VECTOR_SIZE_TO_PIXEL_UNIT(WIDTH), TENSOR##_img, (X) / 4, (Y), VALUES)
1665*c217d954SCole Faust
1666*c217d954SCole Faust
1667*c217d954SCole Faust#define T_LOAD(DATA_TYPE, HEIGHT, WIDTH, TENSOR_TYPE, TENSOR, X, Y, YI_MULTIPLIER, STRIDE_Y, dst)                      \
1668*c217d954SCole Faust    ({                                                                                                                 \
1669*c217d954SCole Faust        LOOP_UNROLLING(int, _i, 0, 1, HEIGHT,                                                                          \
1670*c217d954SCole Faust        {                                                                                                              \
1671*c217d954SCole Faust            dst[_i].v = V_LOAD(DATA_TYPE, WIDTH, TENSOR_TYPE, TENSOR, X, ((Y) + _i * (int)(YI_MULTIPLIER)), STRIDE_Y); \
1672*c217d954SCole Faust        })                                                                                                             \
1673*c217d954SCole Faust    })
1674*c217d954SCole Faust
1675*c217d954SCole Faust
1676*c217d954SCole Faust#define T_LOAD_INDIRECT(DATA_TYPE, HEIGHT, WIDTH, TENSOR_TYPE, TENSOR, X, STRIDE_Y, indirect_y, dst)    \
1677*c217d954SCole Faust    ({                                                                                                  \
1678*c217d954SCole Faust        LOOP_UNROLLING(int, _i, 0, 1, HEIGHT,                                                           \
1679*c217d954SCole Faust        {                                                                                               \
1680*c217d954SCole Faust            dst[_i].v = V_LOAD(DATA_TYPE, WIDTH, TENSOR_TYPE, TENSOR, X, (indirect_y[_i].v), STRIDE_Y); \
1681*c217d954SCole Faust        })                                                                                              \
1682*c217d954SCole Faust    })
1683*c217d954SCole Faust
1684*c217d954SCole Faust
1685*c217d954SCole Faust#define T_LOAD_INDIRECT_WIDTH_SELECT(DATA_TYPE, HEIGHT, WIDTH0, WIDTH1, TENSOR_TYPE, TENSOR, X, STRIDE_Y, WIDTH1_CONDITION, dst, indirect_y)                                                      \
1686*c217d954SCole Faust    ({                                                                                                                                                                                             \
1687*c217d954SCole Faust        if(WIDTH1_CONDITION)                                                                                                                                                                       \
1688*c217d954SCole Faust        {                                                                                                                                                                                          \
1689*c217d954SCole Faust            LOOP_UNROLLING(int, _i, 0, 1, HEIGHT,                                                                                                                                                  \
1690*c217d954SCole Faust            {                                                                                                                                                                                      \
1691*c217d954SCole Faust                VLOAD_PARTIAL(WIDTH0, WIDTH1)                                                         \
1692*c217d954SCole Faust                (dst[HEIGHT - 1 - _i].v, 0, (__global DATA_TYPE *)(TENSOR##_ptr + TENSOR##_offset_first_element_in_bytes + (X) * sizeof(DATA_TYPE) + (indirect_y[HEIGHT - 1 - _i].v) * STRIDE_Y));               \
1693*c217d954SCole Faust            })                                                                                                                                                                                     \
1694*c217d954SCole Faust        }                                                                                                                                                                                          \
1695*c217d954SCole Faust        else                                                                                                                                                                                       \
1696*c217d954SCole Faust        {                                                                                                                                                                                          \
1697*c217d954SCole Faust            LOOP_UNROLLING(int, _i, 0, 1, HEIGHT,                                                                                                                                                  \
1698*c217d954SCole Faust            {                                                                                                                                                                                      \
1699*c217d954SCole Faust                dst[HEIGHT - 1 - _i].v = V_LOAD(DATA_TYPE, WIDTH0, TENSOR_TYPE, TENSOR, X, (indirect_y[HEIGHT - 1 - _i].v), STRIDE_Y); \
1700*c217d954SCole Faust            })                                                                                                                                                                                     \
1701*c217d954SCole Faust        }                                                                                                                                                                                          \
1702*c217d954SCole Faust    })
1703*c217d954SCole Faust
1704*c217d954SCole Faust#define T_LOAD_NHWC(DATA_TYPE, TILE_HEIGHT, TILE_WIDTH, TILE_CHANNELS, TENSOR_TYPE, TENSOR, B, Y, X, C, TENSOR_WIDTH, TENSOR_HEIGHT, STRIDE_Y, dst)   \
1705*c217d954SCole Faust    ({                                                                                                                                                \
1706*c217d954SCole Faust        LOOP_UNROLLING(int, _yk, 0, 1, TILE_HEIGHT,                                                                                                   \
1707*c217d954SCole Faust        {                                                                                                                                             \
1708*c217d954SCole Faust            LOOP_UNROLLING(int, _xk, 0, 1, TILE_WIDTH,                                                                                                \
1709*c217d954SCole Faust            {                                                                                                                                         \
1710*c217d954SCole Faust                int _src_y = (X) + _xk + ((Y) + _yk) * (TENSOR_WIDTH);                                                                                \
1711*c217d954SCole Faust                _src_y    += (B) * (int)(TENSOR_WIDTH) * (int)(TENSOR_HEIGHT);                                                                        \
1712*c217d954SCole Faust                int _src_valid_y = (((X) + _xk) >= 0 && ((X) + _xk) < (int)(TENSOR_WIDTH) && ((Y) + _yk) >= 0 && ((Y) + _yk) < (int)(TENSOR_HEIGHT)); \
1713*c217d954SCole Faust                if(_src_valid_y != 0)                                                                                                                 \
1714*c217d954SCole Faust                {                                                                                                                                     \
1715*c217d954SCole Faust                    dst[_xk + _yk * (TILE_WIDTH)].v = V_LOAD(DATA_TYPE, TILE_CHANNELS, TENSOR_TYPE, TENSOR, C, _src_y, STRIDE_Y);                     \
1716*c217d954SCole Faust                }                                                                                                                                     \
1717*c217d954SCole Faust            })                                                                                                                                        \
1718*c217d954SCole Faust        })                                                                                                                                            \
1719*c217d954SCole Faust    })
1720*c217d954SCole Faust
1721*c217d954SCole Faust
1722*c217d954SCole Faust#define T_LOAD_NHWC_WITH_DILATION(DATA_TYPE, TILE_HEIGHT, TILE_WIDTH, TILE_CHANNELS, TENSOR_TYPE, TENSOR, B, Y, X, C, TENSOR_WIDTH, TENSOR_HEIGHT, DILATION_X, DILATION_Y, BOUNDARY_CHECK, dst)         \
1723*c217d954SCole Faust    ({ \
1724*c217d954SCole Faust        LOOP_UNROLLING(int, _yk, 0, 1, TILE_HEIGHT, \
1725*c217d954SCole Faust        { \
1726*c217d954SCole Faust            LOOP_UNROLLING(int, _xk, 0, 1, TILE_WIDTH, \
1727*c217d954SCole Faust            { \
1728*c217d954SCole Faust                int _src_y = (X) + _xk * (DILATION_X); \
1729*c217d954SCole Faust                int _src_z = ((Y) + _yk * (DILATION_Y)); \
1730*c217d954SCole Faust                int _src_w    = (B); \
1731*c217d954SCole Faust                bool _src_valid_y = (((X) + _xk * (DILATION_X)) >= 0) && (((X) + _xk * (DILATION_X)) < (int)(TENSOR_WIDTH)) && (((Y) + _yk * (DILATION_Y)) >= 0) && (((Y) + _yk * (DILATION_Y)) < (int)(TENSOR_HEIGHT)); \
1732*c217d954SCole Faust                if(!(BOUNDARY_CHECK)) \
1733*c217d954SCole Faust                { \
1734*c217d954SCole Faust                    dst[_xk + _yk * (TILE_WIDTH)].v = VLOAD(TILE_CHANNELS)                                                \
1735*c217d954SCole Faust                    (0, (__global DATA_TYPE *)(TENSOR##_ptr + TENSOR##_offset_first_element_in_bytes + (C) * sizeof(DATA_TYPE) + (_src_y) * (TENSOR##_stride_y) + (_src_z) * (TENSOR##_stride_z) + (_src_w) * (TENSOR##_stride_w))); \
1736*c217d954SCole Faust                } \
1737*c217d954SCole Faust                else \
1738*c217d954SCole Faust                { \
1739*c217d954SCole Faust                    if(_src_valid_y) \
1740*c217d954SCole Faust                    { \
1741*c217d954SCole Faust                        dst[_xk + _yk * (TILE_WIDTH)].v = VLOAD(TILE_CHANNELS)                                                \
1742*c217d954SCole Faust                    (0, (__global DATA_TYPE *)(TENSOR##_ptr + TENSOR##_offset_first_element_in_bytes + (C) * sizeof(DATA_TYPE) + (_src_y) * (TENSOR##_stride_y) + (_src_z) * (TENSOR##_stride_z) + (_src_w) * (TENSOR##_stride_w))); \
1743*c217d954SCole Faust                    }                                                                                                                                                                                                 \
1744*c217d954SCole Faust                } \
1745*c217d954SCole Faust            })                                                                                                                                                                                                             \
1746*c217d954SCole Faust        })                                                                                                                                                                                                             \
1747*c217d954SCole Faust    })
1748*c217d954SCole Faust
1749*c217d954SCole Faust
1750*c217d954SCole Faust#define T_LOAD_NHWC_INDIRECT(DATA_TYPE, TILE_AREA, TILE_CHANNELS, TENSOR_TYPE, TENSOR, B, Y, X, C, TENSOR_WIDTH, TENSOR_HEIGHT, STRIDE_Y, xi, yi, dst)                \
1751*c217d954SCole Faust    ({                                                                                                                                                                \
1752*c217d954SCole Faust        LOOP_UNROLLING(int, _i, 0, 1, TILE_AREA,                                                                                                                      \
1753*c217d954SCole Faust        {                                                                                                                                                             \
1754*c217d954SCole Faust            int _src_y = (X) + xi[_i].v + ((Y) + yi[_i].v) * (TENSOR_WIDTH);                                                                                          \
1755*c217d954SCole Faust            _src_y += (B) * (int)(TENSOR_WIDTH) * (int)(TENSOR_HEIGHT);                                                                                               \
1756*c217d954SCole Faust            int _src_valid_y = (((X) + xi[_i].v) >= 0 && ((X) + xi[_i].v) < (int)(TENSOR_WIDTH) && ((Y) + yi[_i].v) >= 0 && ((Y) + yi[_i].v) < (int)(TENSOR_HEIGHT)); \
1757*c217d954SCole Faust            if(_src_valid_y != 0)                                                                                                                                     \
1758*c217d954SCole Faust            {                                                                                                                                                         \
1759*c217d954SCole Faust                dst[_i].v = V_LOAD(DATA_TYPE, TILE_CHANNELS, TENSOR_TYPE, TENSOR, C, _src_y, STRIDE_Y);                                                               \
1760*c217d954SCole Faust            }                                                                                                                                                         \
1761*c217d954SCole Faust        })                                                                                                                                                            \
1762*c217d954SCole Faust    })
1763*c217d954SCole Faust
1764*c217d954SCole Faust
1765*c217d954SCole Faust#define T_LOAD2D_INDIRECT(DATA_TYPE, TILE_AREA, TILE_CHANNELS, TENSOR_TYPE, TENSOR, C, STRIDE_Y, yi, dst) T_LOAD2D_INDIRECT_STR(DATA_TYPE, TILE_AREA, TILE_CHANNELS, TENSOR_TYPE, TENSOR, C, STRIDE_Y, yi, dst)
1766*c217d954SCole Faust#define T_LOAD2D_INDIRECT_STR(DATA_TYPE, TILE_AREA, TILE_CHANNELS, TENSOR_TYPE, TENSOR, C, STRIDE_Y, yi, dst) T_LOAD2D_INDIRECT_##TENSOR_TYPE(DATA_TYPE, TILE_AREA, TILE_CHANNELS, TENSOR_TYPE, TENSOR, C, STRIDE_Y, yi, dst)
1767*c217d954SCole Faust#define T_LOAD2D_INDIRECT_BUFFER(DATA_TYPE, TILE_AREA, TILE_CHANNELS, TENSOR_TYPE, TENSOR, C, STRIDE_Y, yi, dst) \
1768*c217d954SCole Faust    ({ \
1769*c217d954SCole Faust        LOOP_UNROLLING(int, _i, 0, 1, TILE_AREA, \
1770*c217d954SCole Faust        { \
1771*c217d954SCole Faust            if(yi[0].s[_i] >= 0) \
1772*c217d954SCole Faust            { \
1773*c217d954SCole Faust                dst[_i].v = V_LOAD(DATA_TYPE, TILE_CHANNELS, TENSOR_TYPE, TENSOR, C, yi[0].s[_i], STRIDE_Y); \
1774*c217d954SCole Faust            } \
1775*c217d954SCole Faust        }) \
1776*c217d954SCole Faust    })
1777*c217d954SCole Faust
1778*c217d954SCole Faust#define T_LOAD2D_INDIRECT_IMAGE(DATA_TYPE, TILE_AREA, TILE_CHANNELS, TENSOR_TYPE, TENSOR, C, STRIDE_Y, yi, dst) \
1779*c217d954SCole Faust    ({ \
1780*c217d954SCole Faust        LOOP_UNROLLING(int, _i, 0, 1, TILE_AREA, \
1781*c217d954SCole Faust        { \
1782*c217d954SCole Faust            dst[_i].v = V_LOAD(DATA_TYPE, TILE_CHANNELS, TENSOR_TYPE, TENSOR, C, yi[0].s[_i], STRIDE_Y); \
1783*c217d954SCole Faust        }) \
1784*c217d954SCole Faust    })
1785*c217d954SCole Faust
1786*c217d954SCole Faust
1787*c217d954SCole Faust#define T_LOAD_NDHWC_INDIRECT(DATA_TYPE, TILE_AREA, TILE_CHANNELS, TENSOR_TYPE, TENSOR, B, Z, Y, X, C, TENSOR_WIDTH, TENSOR_HEIGHT, TENSOR_DEPTH, STRIDE_Y, xi, yi, zi, dst) \
1788*c217d954SCole Faust    ({                                                                                                                                                                \
1789*c217d954SCole Faust        LOOP_UNROLLING(int, _i, 0, 1, TILE_AREA,                                                                                                                      \
1790*c217d954SCole Faust        {                                                                                                                                                             \
1791*c217d954SCole Faust            int _src_y = (X) + xi[_i].v + ((Y) + yi[_i].v) * (TENSOR_WIDTH) + ((Z) + zi[_i].v) * (TENSOR_WIDTH * TENSOR_HEIGHT);                                      \
1792*c217d954SCole Faust            _src_y += (B) * (int)(TENSOR_WIDTH) * (int)(TENSOR_HEIGHT) * (int)(TENSOR_DEPTH);                                                                         \
1793*c217d954SCole Faust            int _src_valid_y = (((X) + xi[_i].v) >= 0 && ((X) + xi[_i].v) < (int)(TENSOR_WIDTH) && ((Y) + yi[_i].v) >= 0 && ((Y) + yi[_i].v) < (int)(TENSOR_HEIGHT)   \
1794*c217d954SCole Faust                             && ((Z) + zi[_i].v) >= 0 && ((Z) + zi[_i].v) < (int)(TENSOR_DEPTH));                                                                     \
1795*c217d954SCole Faust            if(_src_valid_y != 0)                                                                                                                                     \
1796*c217d954SCole Faust            {                                                                                                                                                         \
1797*c217d954SCole Faust                dst[_i].v = V_LOAD(DATA_TYPE, TILE_CHANNELS, TENSOR_TYPE, TENSOR, C, _src_y, STRIDE_Y);                                                               \
1798*c217d954SCole Faust            }                                                                                                                                                         \
1799*c217d954SCole Faust        })                                                                                                                                                            \
1800*c217d954SCole Faust    })
1801*c217d954SCole Faust
1802*c217d954SCole Faust
1803*c217d954SCole Faust#define T_STORE_INDIRECT_WIDTH_SELECT(DATA_TYPE, HEIGHT, WIDTH0, WIDTH1, TENSOR_TYPE, TENSOR, X, STRIDE_Y, WIDTH1_CONDITION, src, indirect_y)                                                      \
1804*c217d954SCole Faust    ({                                                                                                                                                                                             \
1805*c217d954SCole Faust        if(WIDTH1_CONDITION)                                                                                                                                                                       \
1806*c217d954SCole Faust        {                                                                                                                                                                                          \
1807*c217d954SCole Faust            LOOP_UNROLLING(int, _i, 0, 1, HEIGHT,                                                                                                                                                  \
1808*c217d954SCole Faust            {                                                                                                                                                                                      \
1809*c217d954SCole Faust                VSTORE_PARTIAL(WIDTH0, WIDTH1)                                                                                                                                                     \
1810*c217d954SCole Faust                (CONVERT(src[HEIGHT - 1 - _i].v, VEC_DATA_TYPE(DATA_TYPE, WIDTH0)), 0, (__global DATA_TYPE *)(TENSOR##_ptr + TENSOR##_offset_first_element_in_bytes + (X) * sizeof(DATA_TYPE) + (indirect_y[HEIGHT - 1 - _i].v) * STRIDE_Y)); \
1811*c217d954SCole Faust            })                                                                                                                                                                                     \
1812*c217d954SCole Faust        }                                                                                                                                                                                          \
1813*c217d954SCole Faust        else                                                                                                                                                                                       \
1814*c217d954SCole Faust        {                                                                                                                                                                                          \
1815*c217d954SCole Faust            LOOP_UNROLLING(int, _i, 0, 1, HEIGHT,                                                                                                                                                  \
1816*c217d954SCole Faust            {                                                                                                                                                                                      \
1817*c217d954SCole Faust                VSTORE(WIDTH0)                                                                                                                                                                     \
1818*c217d954SCole Faust                (CONVERT(src[HEIGHT - 1 - _i].v, VEC_DATA_TYPE(DATA_TYPE, WIDTH0)), 0, (__global DATA_TYPE *)(TENSOR##_ptr + TENSOR##_offset_first_element_in_bytes + (X) * sizeof(DATA_TYPE) + (indirect_y[HEIGHT - 1 - _i].v) * STRIDE_Y)); \
1819*c217d954SCole Faust            })                                                                                                                                                                                     \
1820*c217d954SCole Faust        }                                                                                                                                                                                          \
1821*c217d954SCole Faust    })
1822*c217d954SCole Faust
1823*c217d954SCole Faust
1824*c217d954SCole Faust#define T_OFFSET_CORRECTION(ACC_DATA_TYPE, M0, N0, K0, SRC_OFFSET, WEI_OFFSET, lhs, rhs, dst)        \
1825*c217d954SCole Faust    ({                                                                                               \
1826*c217d954SCole Faust        LOOP_UNROLLING(int, _m0, 0, 1, M0,                                                           \
1827*c217d954SCole Faust        {                                                                                            \
1828*c217d954SCole Faust            ACC_DATA_TYPE _tm = 0;                                                                   \
1829*c217d954SCole Faust            LOOP_UNROLLING(int, _k0, 0, 1, K0,                                                       \
1830*c217d954SCole Faust            {                                                                                        \
1831*c217d954SCole Faust                _tm += ((ACC_DATA_TYPE)lhs[_m0].s[_k0] * (ACC_DATA_TYPE)WEI_OFFSET);                 \
1832*c217d954SCole Faust            })                                                                                       \
1833*c217d954SCole Faust            LOOP_UNROLLING(int, _n0, 0, 1, N0,                                                       \
1834*c217d954SCole Faust            {                                                                                        \
1835*c217d954SCole Faust                dst[_m0].s[_n0] += _tm;                                                              \
1836*c217d954SCole Faust                LOOP_UNROLLING(int, _k0, 0, 1, K0,                                                   \
1837*c217d954SCole Faust                {                                                                                    \
1838*c217d954SCole Faust                    dst[_m0].s[_n0] += ((ACC_DATA_TYPE)rhs[_n0].s[_k0] * (ACC_DATA_TYPE)SRC_OFFSET); \
1839*c217d954SCole Faust                })                                                                                   \
1840*c217d954SCole Faust            })                                                                                       \
1841*c217d954SCole Faust        })                                                                                          \
1842*c217d954SCole Faust    })
1843*c217d954SCole Faust
1844*c217d954SCole Faust
1845*c217d954SCole Faust#define T_QUANTIZE8(SRC_DATA_TYPE, DST_DATA_TYPE, QUANTIZATION_TYPE, M0, N0, DST_OFFSET, DST_SHIFT, DST_MULTIPLIER, src, dst_multipliers, dst_shifts, dst) T_QUANTIZE8_STR(SRC_DATA_TYPE, DST_DATA_TYPE, QUANTIZATION_TYPE, M0, N0, DST_OFFSET, DST_SHIFT, DST_MULTIPLIER, src, dst_multipliers, dst_shifts, dst)
1846*c217d954SCole Faust#define T_QUANTIZE8_STR(SRC_DATA_TYPE, DST_DATA_TYPE, QUANTIZATION_TYPE, M0, N0, DST_OFFSET, DST_SHIFT, DST_MULTIPLIER, src, dst_multipliers, dst_shifts, dst) T_QUANTIZE8_##QUANTIZATION_TYPE(SRC_DATA_TYPE, DST_DATA_TYPE, M0, N0, DST_OFFSET, DST_SHIFT, DST_MULTIPLIER, src, dst_multipliers, dst_shifts, dst)
1847*c217d954SCole Faust
1848*c217d954SCole Faust
1849*c217d954SCole Faust#define T_QUANTIZE8_PER_TENSOR(SRC_DATA_TYPE, DST_DATA_TYPE, M0, N0, DST_OFFSET, DST_SHIFT, DST_MULTIPLIER, src, dst_multipliers, dst_shifts, dst)                          \
1850*c217d954SCole Faust    ({ \
1851*c217d954SCole Faust        LOOP_UNROLLING(int, _m0, 0, 1, M0, \
1852*c217d954SCole Faust        { \
1853*c217d954SCole Faust            LOOP_UNROLLING(int, _n0, 0, 1, N0, \
1854*c217d954SCole Faust            { \
1855*c217d954SCole Faust                SRC_DATA_TYPE _tmp = 0; \
1856*c217d954SCole Faust                SRC_DATA_TYPE _src = src[_m0].s[_n0]; \
1857*c217d954SCole Faust                _src *= select((SRC_DATA_TYPE)1, ((SRC_DATA_TYPE)1 << (SRC_DATA_TYPE)(-DST_SHIFT)), ((SRC_DATA_TYPE)DST_SHIFT < (SRC_DATA_TYPE)0)); \
1858*c217d954SCole Faust                SRC_DATA_TYPE overflow = _src == DST_MULTIPLIER && _src == INT_MIN; \
1859*c217d954SCole Faust                long a_64 = (long)(_src); \
1860*c217d954SCole Faust                long b_64 = (long)(DST_MULTIPLIER); \
1861*c217d954SCole Faust                long ab_64 = a_64 * b_64; \
1862*c217d954SCole Faust                long mask1 = 1 << 30; \
1863*c217d954SCole Faust                long mask2 = 1 - (1 << 30); \
1864*c217d954SCole Faust                long is_positive_or_zero = ab_64 >= 0; \
1865*c217d954SCole Faust                long nudge = select(mask2, mask1, is_positive_or_zero); \
1866*c217d954SCole Faust                SRC_DATA_TYPE ab_x2_high32 = CONVERT((ab_64 + nudge) / (long)(1ll << 31), SRC_DATA_TYPE); \
1867*c217d954SCole Faust                _tmp = select(ab_x2_high32, (SRC_DATA_TYPE)INT_MAX, overflow); \
1868*c217d954SCole Faust                if(DST_SHIFT >= 0) \
1869*c217d954SCole Faust                { \
1870*c217d954SCole Faust                    long mask = ((((int)1) << DST_SHIFT) - (long)1); \
1871*c217d954SCole Faust                    long threshold = _tmp < (int)0 ? (mask >> 1) + (long)1 : (mask >> 1) + 0; \
1872*c217d954SCole Faust                    _tmp = (_tmp & mask) > threshold ? (_tmp >> DST_SHIFT) + (int)1 : (_tmp >> DST_SHIFT); \
1873*c217d954SCole Faust                } \
1874*c217d954SCole Faust                _tmp += DST_OFFSET; \
1875*c217d954SCole Faust                dst[_m0].s[_n0] = CONVERT_SAT(_tmp, DST_DATA_TYPE);                                                                            \
1876*c217d954SCole Faust            })                                                                                                                                          \
1877*c217d954SCole Faust        })                                                                                                                                          \
1878*c217d954SCole Faust    })
1879*c217d954SCole Faust
1880*c217d954SCole Faust
1881*c217d954SCole Faust#define T_QUANTIZE8_PER_CHANNEL(SRC_DATA_TYPE, DST_DATA_TYPE, M0, N0, DST_OFFSET, DST_SHIFT, DST_MULTIPLIER, src, dst_multipliers, dst_shifts, dst)                          \
1882*c217d954SCole Faust    ({ \
1883*c217d954SCole Faust        LOOP_UNROLLING(int, _m0, 0, 1, M0, \
1884*c217d954SCole Faust        { \
1885*c217d954SCole Faust            LOOP_UNROLLING(int, _n0, 0, 1, N0, \
1886*c217d954SCole Faust            { \
1887*c217d954SCole Faust                SRC_DATA_TYPE _tmp = 0; \
1888*c217d954SCole Faust                SRC_DATA_TYPE _tmp2 = 0; \
1889*c217d954SCole Faust                SRC_DATA_TYPE _src = src[_m0].s[_n0]; \
1890*c217d954SCole Faust                SRC_DATA_TYPE _dst_multiplier = dst_multipliers[0].s[_n0]; \
1891*c217d954SCole Faust                SRC_DATA_TYPE _dst_shift = dst_shifts[0].s[_n0]; \
1892*c217d954SCole Faust                _src *= select((SRC_DATA_TYPE)1, ((SRC_DATA_TYPE)1 << (SRC_DATA_TYPE)(-_dst_shift)), ((SRC_DATA_TYPE)_dst_shift < (SRC_DATA_TYPE)0)); \
1893*c217d954SCole Faust                SRC_DATA_TYPE overflow = _src == _dst_multiplier && _src == INT_MIN; \
1894*c217d954SCole Faust                long a_64 = (long)(_src); \
1895*c217d954SCole Faust                long b_64 = (long)(_dst_multiplier); \
1896*c217d954SCole Faust                long ab_64 = a_64 * b_64; \
1897*c217d954SCole Faust                long mask1 = 1 << 30; \
1898*c217d954SCole Faust                long mask2 = 1 - (1 << 30); \
1899*c217d954SCole Faust                long is_positive_or_zero = ab_64 >= 0; \
1900*c217d954SCole Faust                long nudge = select(mask2, mask1, is_positive_or_zero); \
1901*c217d954SCole Faust                SRC_DATA_TYPE ab_x2_high32 = CONVERT((ab_64 + nudge) / (long)(1ll << 31), SRC_DATA_TYPE); \
1902*c217d954SCole Faust                _tmp = select(ab_x2_high32, (SRC_DATA_TYPE)INT_MAX, overflow); \
1903*c217d954SCole Faust                long mask = ((((int)1) << _dst_shift) - (int)1); \
1904*c217d954SCole Faust                long threshold = (mask >> 1) + any(_tmp); \
1905*c217d954SCole Faust                _tmp2 = _tmp >> _dst_shift; \
1906*c217d954SCole Faust                _tmp2 += select(0, 1, (_tmp & mask) > threshold); \
1907*c217d954SCole Faust                _tmp = select(_tmp, _tmp2, _dst_shift >= 0); \
1908*c217d954SCole Faust                _tmp += DST_OFFSET; \
1909*c217d954SCole Faust                dst[_m0].s[_n0] = CONVERT_SAT(_tmp, DST_DATA_TYPE);                                                                            \
1910*c217d954SCole Faust            })                                                                                                                                          \
1911*c217d954SCole Faust        })                                                                                                                                         \
1912*c217d954SCole Faust    })
1913*c217d954SCole Faust
1914*c217d954SCole Faust
1915*c217d954SCole Faust#define T_QUANTIZE8_ASYMMETRIC(SRC_DATA_TYPE, DST_DATA_TYPE, M0, N0, DST_OFFSET, DST_SHIFT, DST_MULTIPLIER, src, dst)                          \
1916*c217d954SCole Faust    ({ \
1917*c217d954SCole Faust        LOOP_UNROLLING(int, _m0, 0, 1, M0, \
1918*c217d954SCole Faust        { \
1919*c217d954SCole Faust            LOOP_UNROLLING(int, _n0, 0, 1, N0, \
1920*c217d954SCole Faust            { \
1921*c217d954SCole Faust                SRC_DATA_TYPE _tmp = 0; \
1922*c217d954SCole Faust                SRC_DATA_TYPE _src = src[_m0].s[_n0]; \
1923*c217d954SCole Faust                _src *= select((SRC_DATA_TYPE)1, ((SRC_DATA_TYPE)1 << (SRC_DATA_TYPE)(-DST_SHIFT)), ((SRC_DATA_TYPE)DST_SHIFT < (SRC_DATA_TYPE)0)); \
1924*c217d954SCole Faust                SRC_DATA_TYPE overflow = _src == DST_MULTIPLIER && _src == INT_MIN; \
1925*c217d954SCole Faust                long a_64 = (long)(_src); \
1926*c217d954SCole Faust                long b_64 = (long)(DST_MULTIPLIER); \
1927*c217d954SCole Faust                long ab_64 = a_64 * b_64; \
1928*c217d954SCole Faust                long mask1 = 1 << 30; \
1929*c217d954SCole Faust                long mask2 = 1 - (1 << 30); \
1930*c217d954SCole Faust                long is_positive_or_zero = ab_64 >= 0; \
1931*c217d954SCole Faust                long nudge = select(mask2, mask1, is_positive_or_zero); \
1932*c217d954SCole Faust                SRC_DATA_TYPE ab_x2_high32 = CONVERT((ab_64 + nudge) / (long)(1ll << 31), SRC_DATA_TYPE); \
1933*c217d954SCole Faust                _tmp = select(ab_x2_high32, (SRC_DATA_TYPE)INT_MAX, overflow); \
1934*c217d954SCole Faust                if(DST_SHIFT >= 0) \
1935*c217d954SCole Faust                { \
1936*c217d954SCole Faust                    long mask = ((((int)1) << DST_SHIFT) - (int)1); \
1937*c217d954SCole Faust                    long threshold = _tmp < (int)0 ? (mask >> 1) + (long)1 : (mask >> 1) + 0; \
1938*c217d954SCole Faust                    _tmp = (_tmp & mask) > threshold ? (_tmp >> DST_SHIFT) + (int)1 : (_tmp >> DST_SHIFT); \
1939*c217d954SCole Faust                } \
1940*c217d954SCole Faust                _tmp += DST_OFFSET; \
1941*c217d954SCole Faust                dst[_m0].s[_n0] = CONVERT_SAT(_tmp, DST_DATA_TYPE);                                                                            \
1942*c217d954SCole Faust            })                                                                                                                                          \
1943*c217d954SCole Faust        })                                                                                                                                          \
1944*c217d954SCole Faust    })
1945*c217d954SCole Faust
1946*c217d954SCole Faust
1947*c217d954SCole Faust#define T_ROWSET_MASK(DATA_TYPE, M0, N0, VALUE_TO_SET, a, mask)                                                                                            \
1948*c217d954SCole Faust    ({                                                                                                                                                     \
1949*c217d954SCole Faust        LOOP_UNROLLING(int, _m0, 0, 1, M0,                                                                                                                 \
1950*c217d954SCole Faust        {                                                                                                                                                  \
1951*c217d954SCole Faust            LOOP_UNROLLING(int, _n0, 0, 1, N0,                                                                                                             \
1952*c217d954SCole Faust            {                                                                                                                                              \
1953*c217d954SCole Faust                a[_m0].s[_n0] = select((DATA_TYPE)(a[_m0].s[_n0]), (DATA_TYPE)(VALUE_TO_SET), (SELECT_DATA_TYPE(DATA_TYPE))(mask[_m0].v == (DATA_TYPE)0)); \
1954*c217d954SCole Faust            })                                                                                                                                             \
1955*c217d954SCole Faust        })                                                                                                                                                 \
1956*c217d954SCole Faust    })
1957*c217d954SCole Faust
1958*c217d954SCole Faust
1959*c217d954SCole Faust#define T_ACTIVATION(DATA_TYPE, M0, N0, ACTIVATION_TYPE, A_VAL, B_VAL, src, dst)               \
1960*c217d954SCole Faust    ({                                                                                         \
1961*c217d954SCole Faust        LOOP_UNROLLING(int, _m0, 0, 1, M0,                                                     \
1962*c217d954SCole Faust        {                                                                                      \
1963*c217d954SCole Faust            dst[_m0].v = ACTIVATION(ACTIVATION_TYPE, DATA_TYPE, N0, src[_m0].v, A_VAL, B_VAL); \
1964*c217d954SCole Faust        })                                                                                     \
1965*c217d954SCole Faust    })
1966*c217d954SCole Faust
1967*c217d954SCole Faust
1968*c217d954SCole Faust#define relu_op_quantized(DATA_TYPE, VEC_SIZE, ZERO_VALUE, A_VAL, B_VAL, x) (max((DATA_TYPE)ZERO_VALUE, x))
1969*c217d954SCole Faust
1970*c217d954SCole Faust#define brelu_op_quantized(DATA_TYPE, VEC_SIZE, ZERO_VALUE, A_VAL, B_VAL, x) (min((DATA_TYPE)A_VAL, max((DATA_TYPE)ZERO_VALUE, x)))
1971*c217d954SCole Faust
1972*c217d954SCole Faust#define lu_brelu_op_quantized(DATA_TYPE, VEC_SIZE, ZERO_VALUE, A_VAL, B_VAL, x) (min(max(x, (DATA_TYPE)B_VAL), (DATA_TYPE)A_VAL))
1973*c217d954SCole Faust
1974*c217d954SCole Faust#define hard_swish_op_quantized(DATA_TYPE, VEC_SIZE, ZERO_VALUE, A_VAL, B_VAL, x) (x * ((min(max((DATA_TYPE)(x + (DATA_TYPE)3.f), (DATA_TYPE)0.f), (DATA_TYPE)6.f)) * (DATA_TYPE)0.166666667f))
1975*c217d954SCole Faust
1976*c217d954SCole Faust#define identity_op_quantized(DATA_TYPE, VEC_SIZE, ZERO_VALUE, A_VAL, B_VAL, x) (x)
1977*c217d954SCole Faust
1978*c217d954SCole Faust#define ACT_OP_QUANTIZED(op, DATA_TYPE, VEC_SIZE, ZERO_VALUE, A_VAL, B_VAL, x) op##_op_quantized(DATA_TYPE, VEC_SIZE, ZERO_VALUE, A_VAL, B_VAL, x)
1979*c217d954SCole Faust#define ACTIVATION_QUANTIZED(op, DATA_TYPE, VEC_SIZE, ZERO_VALUE, A_VAL, B_VAL, x) ACT_OP_QUANTIZED(op, DATA_TYPE, VEC_SIZE, ZERO_VALUE, A_VAL, B_VAL, x)
1980*c217d954SCole Faust
1981*c217d954SCole Faust#define V_ADD(A_VAL, B_VAL) ((A_VAL) + (B_VAL))
1982*c217d954SCole Faust#define V_SUB(A_VAL, B_VAL) ((A_VAL) - (B_VAL))
1983*c217d954SCole Faust#define V_DIV(A_VAL, B_VAL) ((A_VAL) / (B_VAL))
1984*c217d954SCole Faust#define V_MUL(A_VAL, B_VAL) ((A_VAL) * (B_VAL))
1985*c217d954SCole Faust
1986*c217d954SCole Faust
1987*c217d954SCole Faust#define T_ACTIVATION_QUANTIZED(DATA_TYPE, M0, N0, ACTIVATION_TYPE, ZERO_VALUE, A_VAL, B_VAL, src, dst)               \
1988*c217d954SCole Faust    ({ \
1989*c217d954SCole Faust        LOOP_UNROLLING(int, _m0, 0, 1, M0, \
1990*c217d954SCole Faust        { \
1991*c217d954SCole Faust            dst[_m0].v = ACTIVATION_QUANTIZED(ACTIVATION_TYPE, DATA_TYPE, N0, ZERO_VALUE, A_VAL, B_VAL, src[_m0].v); \
1992*c217d954SCole Faust        })                                                                                          \
1993*c217d954SCole Faust    })
1994*c217d954SCole Faust
1995*c217d954SCole Faust
1996*c217d954SCole Faust#define T_ADD(DATA_TYPE, M0, N0, lhs, rhs, dst) \
1997*c217d954SCole Faust    ({                                                            \
1998*c217d954SCole Faust        LOOP_UNROLLING(int, _m0, 0, 1, M0,                        \
1999*c217d954SCole Faust        {                                                         \
2000*c217d954SCole Faust            dst[_m0].v = lhs[_m0].v + rhs[_m0].v; \
2001*c217d954SCole Faust        })                                                        \
2002*c217d954SCole Faust    })
2003*c217d954SCole Faust
2004*c217d954SCole Faust
2005*c217d954SCole Faust#define T_ADD_CONSTANT(DATA_TYPE, M0, N0, lhs, rhs_constant, dst) \
2006*c217d954SCole Faust    ({                                                            \
2007*c217d954SCole Faust        LOOP_UNROLLING(int, _m0, 0, 1, M0,                        \
2008*c217d954SCole Faust        {                                                         \
2009*c217d954SCole Faust            dst[_m0].v = lhs[_m0].v + (DATA_TYPE)rhs_constant;               \
2010*c217d954SCole Faust        })                                                        \
2011*c217d954SCole Faust    })
2012*c217d954SCole Faust
2013*c217d954SCole Faust#define T_ELTWISE_BROADCAST_ADD_X(DST_DATA_TYPE, M0, N0, lhs, rhs, dst) T_ELTWISE_BROADCAST_X(V_ADD, DST_DATA_TYPE, M0, N0, lhs, rhs, dst)
2014*c217d954SCole Faust#define T_ELTWISE_BROADCAST_LHS_X_ADD(DST_DATA_TYPE, M0, N0, lhs, rhs, dst) T_ELTWISE_BROADCAST_LHS_X(V_ADD, DST_DATA_TYPE, M0, N0, lhs, rhs, dst)
2015*c217d954SCole Faust#define T_ELTWISE_BROADCAST_RHS_X_ADD(DST_DATA_TYPE, M0, N0, lhs, rhs, dst) T_ELTWISE_BROADCAST_X(V_ADD, DST_DATA_TYPE, M0, N0, lhs, rhs, dst)
2016*c217d954SCole Faust
2017*c217d954SCole Faust#define T_ELTWISE_BROADCAST_LHS_X_SUB(DST_DATA_TYPE, M0, N0, lhs, rhs, dst) T_ELTWISE_BROADCAST_LHS_X(V_SUB, DST_DATA_TYPE, M0, N0, lhs, rhs, dst)
2018*c217d954SCole Faust#define T_ELTWISE_BROADCAST_RHS_X_SUB(DST_DATA_TYPE, M0, N0, lhs, rhs, dst) T_ELTWISE_BROADCAST_X(V_SUB, DST_DATA_TYPE, M0, N0, lhs, rhs, dst)
2019*c217d954SCole Faust
2020*c217d954SCole Faust#define T_ELTWISE_BROADCAST_DIV_X(DST_DATA_TYPE, M0, N0, lhs, rhs, dst) T_ELTWISE_BROADCAST_X(V_DIV, DST_DATA_TYPE, M0, N0, lhs, rhs, dst)
2021*c217d954SCole Faust
2022*c217d954SCole Faust#define T_ELTWISE_BROADCAST_LHS_X_MUL(DST_DATA_TYPE, M0, N0, lhs, rhs, dst) T_ELTWISE_BROADCAST_LHS_X(V_MUL, DST_DATA_TYPE, M0, N0, lhs, rhs, dst)
2023*c217d954SCole Faust#define T_ELTWISE_BROADCAST_RHS_X_MUL(DST_DATA_TYPE, M0, N0, lhs, rhs, dst) T_ELTWISE_BROADCAST_X(V_MUL, DST_DATA_TYPE, M0, N0, lhs, rhs, dst)
2024*c217d954SCole Faust
2025*c217d954SCole Faust
2026*c217d954SCole Faust#define T_SCALE_CONSTANT(DATA_TYPE, M0, N0, lhs, rhs_constant, dst) \
2027*c217d954SCole Faust    ({                                                            \
2028*c217d954SCole Faust        LOOP_UNROLLING(int, _m0, 0, 1, M0,                        \
2029*c217d954SCole Faust        {                                                         \
2030*c217d954SCole Faust            dst[_m0].v = lhs[_m0].v * (DATA_TYPE)rhs_constant; \
2031*c217d954SCole Faust        })                                                        \
2032*c217d954SCole Faust    })
2033*c217d954SCole Faust
2034*c217d954SCole Faust
2035*c217d954SCole Faust#define T_ELTWISE_BROADCAST_X(T_ELWISE_OP, DST_DATA_TYPE, M0, N0, lhs, rhs, dst) \
2036*c217d954SCole Faust    ({                                                      \
2037*c217d954SCole Faust        LOOP_UNROLLING(int, _m0, 0, 1, M0,                  \
2038*c217d954SCole Faust        {                                                   \
2039*c217d954SCole Faust            dst[_m0].v = T_ELWISE_OP(CONVERT(lhs[_m0].v, VEC_DATA_TYPE(DST_DATA_TYPE, N0)), CONVERT(rhs[0].v, VEC_DATA_TYPE(DST_DATA_TYPE, N0)));             \
2040*c217d954SCole Faust        })                                                  \
2041*c217d954SCole Faust    })
2042*c217d954SCole Faust
2043*c217d954SCole Faust
2044*c217d954SCole Faust#define T_ELTWISE_BROADCAST_LHS_X(T_ELWISE_OP, DST_DATA_TYPE, M0, N0, lhs, rhs, dst) \
2045*c217d954SCole Faust    ({                                                      \
2046*c217d954SCole Faust        LOOP_UNROLLING(int, _m0, 0, 1, M0,                  \
2047*c217d954SCole Faust        {                                                   \
2048*c217d954SCole Faust            dst[_m0].v = T_ELWISE_OP(CONVERT(lhs[0].v, VEC_DATA_TYPE(DST_DATA_TYPE, N0)), CONVERT(rhs[_m0].v, VEC_DATA_TYPE(DST_DATA_TYPE, N0)));             \
2049*c217d954SCole Faust        })                                                  \
2050*c217d954SCole Faust    })
2051*c217d954SCole Faust
2052*c217d954SCole Faust#define T_ELTWISE_ADD(DST_DATA_TYPE, M0, N0, lhs, rhs, dst) T_ELTWISE(V_ADD, DST_DATA_TYPE, M0, N0, lhs, rhs, dst)
2053*c217d954SCole Faust#define T_ELTWISE_SUB(DST_DATA_TYPE, M0, N0, lhs, rhs, dst) T_ELTWISE(V_SUB, DST_DATA_TYPE, M0, N0, lhs, rhs, dst)
2054*c217d954SCole Faust#define T_ELTWISE_DIV(DST_DATA_TYPE, M0, N0, lhs, rhs, dst) T_ELTWISE(V_DIV, DST_DATA_TYPE, M0, N0, lhs, rhs, dst)
2055*c217d954SCole Faust#define T_ELTWISE_MUL(DST_DATA_TYPE, M0, N0, lhs, rhs, dst) T_ELTWISE(V_MUL, DST_DATA_TYPE, M0, N0, lhs, rhs, dst)
2056*c217d954SCole Faust
2057*c217d954SCole Faust
2058*c217d954SCole Faust#define T_ELTWISE(T_ELWISE_OP, DST_DATA_TYPE, M0, N0, lhs, rhs, dst) \
2059*c217d954SCole Faust    ({                                                      \
2060*c217d954SCole Faust        LOOP_UNROLLING(int, _m0, 0, 1, M0,                  \
2061*c217d954SCole Faust        {                                                   \
2062*c217d954SCole Faust            dst[_m0].v = T_ELWISE_OP(CONVERT(lhs[_m0].v, VEC_DATA_TYPE(DST_DATA_TYPE, N0)), CONVERT(rhs[_m0].v, VEC_DATA_TYPE(DST_DATA_TYPE, N0)));             \
2063*c217d954SCole Faust        })                                                  \
2064*c217d954SCole Faust    })
2065*c217d954SCole Faust
2066*c217d954SCole Faust
2067*c217d954SCole Faust#define T_FLOOR(DST_DATA_TYPE, M0, N0, src, dst) \
2068*c217d954SCole Faust    ({                                                      \
2069*c217d954SCole Faust        LOOP_UNROLLING(int, _m0, 0, 1, M0,                  \
2070*c217d954SCole Faust        {                                                   \
2071*c217d954SCole Faust            dst[_m0].v = floor(CONVERT(src[_m0].v, VEC_DATA_TYPE(DST_DATA_TYPE, N0)));             \
2072*c217d954SCole Faust        })                                                  \
2073*c217d954SCole Faust    })
2074*c217d954SCole Faust
2075*c217d954SCole Faust
2076*c217d954SCole Faust#define T_MMUL(LHS_DATA_TYPE, RHS_DATA_TYPE, DST_DATA_TYPE, M0, N0, K0, LHS_LAYOUT, RHS_LAYOUT, lhs, rhs, dst) T_MMUL_##LHS_LAYOUT##_##RHS_LAYOUT(LHS_DATA_TYPE, RHS_DATA_TYPE, DST_DATA_TYPE, M0, N0, K0, lhs, rhs, dst)
2077*c217d954SCole Faust#define T_MMUL_NT_T(LHS_DATA_TYPE, RHS_DATA_TYPE, DST_DATA_TYPE, M0, N0, K0, lhs, rhs, dst) T_MMUL_NT_T_##LHS_DATA_TYPE##_##RHS_DATA_TYPE##_##DST_DATA_TYPE(LHS_DATA_TYPE, RHS_DATA_TYPE, DST_DATA_TYPE, M0, N0, K0, lhs, rhs, dst)
2078*c217d954SCole Faust#define T_MMUL_NT_T_float_float_float(LHS_DATA_TYPE, RHS_DATA_TYPE, DST_DATA_TYPE, M0, N0, K0, lhs, rhs, dst) T_MMUL_NT_T_FLOAT(LHS_DATA_TYPE, RHS_DATA_TYPE, DST_DATA_TYPE, M0, N0, K0, lhs, rhs, dst)
2079*c217d954SCole Faust#define T_MMUL_NT_T_half_half_float(LHS_DATA_TYPE, RHS_DATA_TYPE, DST_DATA_TYPE, M0, N0, K0, lhs, rhs, dst) T_MMUL_NT_T_FLOAT(LHS_DATA_TYPE, RHS_DATA_TYPE, DST_DATA_TYPE, M0, N0, K0, lhs, rhs, dst)
2080*c217d954SCole Faust#define T_MMUL_NT_T_half_half_half(LHS_DATA_TYPE, RHS_DATA_TYPE, DST_DATA_TYPE, M0, N0, K0, lhs, rhs, dst) T_MMUL_NT_T_FLOAT(LHS_DATA_TYPE, RHS_DATA_TYPE, DST_DATA_TYPE, M0, N0, K0, lhs, rhs, dst)
2081*c217d954SCole Faust#define T_MMUL_NT_T_char_char_int(LHS_DATA_TYPE, RHS_DATA_TYPE, DST_DATA_TYPE, M0, N0, K0, lhs, rhs, dst) T_MMUL_NT_T_INTEGER8(LHS_DATA_TYPE, RHS_DATA_TYPE, DST_DATA_TYPE, M0, N0, K0, lhs, rhs, dst)
2082*c217d954SCole Faust#define T_MMUL_NT_T_uchar_uchar_uint(LHS_DATA_TYPE, RHS_DATA_TYPE, DST_DATA_TYPE, M0, N0, K0, lhs, rhs, dst) T_MMUL_NT_T_INTEGER8(LHS_DATA_TYPE, RHS_DATA_TYPE, DST_DATA_TYPE, M0, N0, K0, lhs, rhs, dst)
2083*c217d954SCole Faust#define T_MMUL_NT_T_uchar_uchar_int(LHS_DATA_TYPE, RHS_DATA_TYPE, DST_DATA_TYPE, M0, N0, K0, lhs, rhs, dst) T_MMUL_NT_T_INTEGER8(LHS_DATA_TYPE, RHS_DATA_TYPE, DST_DATA_TYPE, M0, N0, K0, lhs, rhs, dst)
2084*c217d954SCole Faust#define T_MMUL_NT_T_FLOAT(LHS_DATA_TYPE, RHS_DATA_TYPE, DST_DATA_TYPE, M0, N0, K0, lhs, rhs, dst)                       \
2085*c217d954SCole Faust    {                                                                                     \
2086*c217d954SCole Faust        LOOP_UNROLLING(int, _m, 0, 1, M0,                                                 \
2087*c217d954SCole Faust        {                                                                                 \
2088*c217d954SCole Faust            LOOP_UNROLLING(int, _n, 0, 1, N0,                                             \
2089*c217d954SCole Faust            {                                                                             \
2090*c217d954SCole Faust                LOOP_UNROLLING(int, _k, 0, 1, K0,                                         \
2091*c217d954SCole Faust                {                                                                         \
2092*c217d954SCole Faust                    dst[_m].s[_n] = fma((DST_DATA_TYPE)(lhs[_m].s[_k]), (DST_DATA_TYPE)(rhs[_n].s[_k]), dst[_m].s[_n]); \
2093*c217d954SCole Faust                })                                                                        \
2094*c217d954SCole Faust            })                                                                            \
2095*c217d954SCole Faust        })                                                                                \
2096*c217d954SCole Faust    }
2097*c217d954SCole Faust
2098*c217d954SCole Faust#define T_MMUL_NT_T_INTEGER8(LHS_DATA_TYPE, RHS_DATA_TYPE, DST_DATA_TYPE, M0, N0, K0, lhs, rhs, dst)                            \
2099*c217d954SCole Faust    ({ \
2100*c217d954SCole Faust        LOOP_UNROLLING(int, _m, 0, 1, M0, \
2101*c217d954SCole Faust        { \
2102*c217d954SCole Faust            LOOP_UNROLLING(int, _n, 0, 1, N0, \
2103*c217d954SCole Faust            { \
2104*c217d954SCole Faust                DOT_PRODUCT_INTEGER8(LHS_DATA_TYPE, RHS_DATA_TYPE, DST_DATA_TYPE, K0, (lhs[_m].v), (rhs[_n].v), dst[_m].s[_n]); \
2105*c217d954SCole Faust            })                                                                                             \
2106*c217d954SCole Faust        })                                                                                             \
2107*c217d954SCole Faust    })
2108*c217d954SCole Faust
2109*c217d954SCole Faust#endif
2110*c217d954SCole Faust
2111*c217d954SCole Faust#if defined(SCALE_NEAREST_NEIGHBOUR)
2112*c217d954SCole Faust
2113*c217d954SCole Faust
2114*c217d954SCole Faust
2115*c217d954SCole Faust__kernel void scale_nearest_neighbour_nhwc(
2116*c217d954SCole Faust    TENSOR4D_RO_T(src, SRC_TENSOR_TYPE),
2117*c217d954SCole Faust    TENSOR4D_WO_T(dst, DST_TENSOR_TYPE),
2118*c217d954SCole Faust    const float scale_x,
2119*c217d954SCole Faust    const float scale_y)
2120*c217d954SCole Faust{
2121*c217d954SCole Faust    const int cout = GET_SPATIAL_IDX(0, N0, PARTIAL_N0);
2122*c217d954SCole Faust    const int xo   = GET_SPATIAL_IDX(1, 1, 0);
2123*c217d954SCole Faust#if defined(BATCHED_EXECUTION)
2124*c217d954SCole Faust    const int yo   = GET_SPATIAL_IDX(2, 1, 0) % dst_h;
2125*c217d954SCole Faust    const int bout = GET_SPATIAL_IDX(2, 1, 0) / dst_h;
2126*c217d954SCole Faust#else
2127*c217d954SCole Faust    const int yo   = GET_SPATIAL_IDX(2, 1, 0);
2128*c217d954SCole Faust    const int bout = 0;
2129*c217d954SCole Faust#endif
2130*c217d954SCole Faust
2131*c217d954SCole Faust#ifdef SAMPLING_POLICY_TOP_LEFT
2132*c217d954SCole Faust    float xi_f = (xo * scale_x);
2133*c217d954SCole Faust    float yi_f = (yo * scale_y);
2134*c217d954SCole Faust#elif SAMPLING_POLICY_CENTER
2135*c217d954SCole Faust    float     xi_f = ((xo + 0.5f) * scale_x);
2136*c217d954SCole Faust    float     yi_f = ((yo + 0.5f) * scale_y);
2137*c217d954SCole Faust#else
2138*c217d954SCole Faust#error("Unsupported sampling policy");
2139*c217d954SCole Faust#endif
2140*c217d954SCole Faust
2141*c217d954SCole Faust#ifdef ALIGN_CORNERS
2142*c217d954SCole Faust    xi_f = round(xi_f);
2143*c217d954SCole Faust    yi_f = round(yi_f);
2144*c217d954SCole Faust#endif
2145*c217d954SCole Faust
2146*c217d954SCole Faust    const int xi0 = clamp((int)xi_f, 0, (int)src_w - 1);
2147*c217d954SCole Faust    const int yi0 = clamp((int)yi_f, 0, (int)src_h - 1);
2148*c217d954SCole Faust
2149*c217d954SCole Faust    TILE(SRC_DATA_TYPE, 1, N0, in00);
2150*c217d954SCole Faust
2151*c217d954SCole Faust    T_LOAD_NHWC_WITH_DILATION(SRC_DATA_TYPE, 1, 1, N0, SRC_TENSOR_TYPE, src, bout, yi0, xi0, cout, src_w, src_h, 1, 1, false, in00);
2152*c217d954SCole Faust
2153*c217d954SCole Faust    TILE(uint, 1, 1, dst_indirect_y);
2154*c217d954SCole Faust
2155*c217d954SCole Faust
2156*c217d954SCole Faust    dst_indirect_y[0].v = xo + (yo * (int)(dst_w)) + bout * (int)(dst_w * dst_h);
2157*c217d954SCole Faust
2158*c217d954SCole Faust    bool x_cond = PARTIAL_N0 != 0 && get_global_id(0) == 0;
2159*c217d954SCole Faust
2160*c217d954SCole Faust    T_STORE_INDIRECT_WIDTH_SELECT(DST_DATA_TYPE, 1, N0, PARTIAL_N0, DST_TENSOR_TYPE, dst, cout, dst_stride_y, x_cond, in00, dst_indirect_y);
2161*c217d954SCole Faust}
2162*c217d954SCole Faust#endif
2163*c217d954SCole Faust
2164*c217d954SCole Faust#if defined(SCALE_BILINEAR)
2165*c217d954SCole Faust
2166*c217d954SCole Faust
2167*c217d954SCole Faust
2168*c217d954SCole Faust__kernel void scale_bilinear_nhwc(
2169*c217d954SCole Faust    TENSOR4D_RO_T(src, SRC_TENSOR_TYPE),
2170*c217d954SCole Faust    TENSOR4D_WO_T(dst, DST_TENSOR_TYPE),
2171*c217d954SCole Faust    const float scale_x,
2172*c217d954SCole Faust    const float scale_y)
2173*c217d954SCole Faust{
2174*c217d954SCole Faust    const int cout = GET_SPATIAL_IDX(0, N0, PARTIAL_N0);
2175*c217d954SCole Faust    const int xo   = GET_SPATIAL_IDX(1, 1, 0);
2176*c217d954SCole Faust#if defined(BATCHED_EXECUTION)
2177*c217d954SCole Faust    const int yo   = GET_SPATIAL_IDX(2, 1, 0) % dst_h;
2178*c217d954SCole Faust    const int bout = GET_SPATIAL_IDX(2, 1, 0) / dst_h;
2179*c217d954SCole Faust#else
2180*c217d954SCole Faust    const int yo   = GET_SPATIAL_IDX(2, 1, 0);
2181*c217d954SCole Faust    const int bout = 0;
2182*c217d954SCole Faust#endif
2183*c217d954SCole Faust
2184*c217d954SCole Faust#ifdef SAMPLING_POLICY_TOP_LEFT
2185*c217d954SCole Faust    float xi_f = (xo * scale_x);
2186*c217d954SCole Faust    float yi_f = (yo * scale_y);
2187*c217d954SCole Faust#elif SAMPLING_POLICY_CENTER
2188*c217d954SCole Faust    float     xi_f = ((xo + 0.5f) * scale_x - 0.5f);
2189*c217d954SCole Faust    float     yi_f = ((yo + 0.5f) * scale_y - 0.5f);
2190*c217d954SCole Faust#else
2191*c217d954SCole Faust#error("Unsupported sampling policy");
2192*c217d954SCole Faust#endif
2193*c217d954SCole Faust
2194*c217d954SCole Faust    const int xi = (int)floor(xi_f);
2195*c217d954SCole Faust    const int yi = (int)floor(yi_f);
2196*c217d954SCole Faust
2197*c217d954SCole Faust    TILE(SRC_DATA_TYPE, 1, N0, in00);
2198*c217d954SCole Faust    TILE(SRC_DATA_TYPE, 1, N0, in01);
2199*c217d954SCole Faust    TILE(SRC_DATA_TYPE, 1, N0, in10);
2200*c217d954SCole Faust    TILE(SRC_DATA_TYPE, 1, N0, in11);
2201*c217d954SCole Faust
2202*c217d954SCole Faust
2203*c217d954SCole Faust    in00[0].v = CONSTANT_VALUE;
2204*c217d954SCole Faust    in01[0].v = CONSTANT_VALUE;
2205*c217d954SCole Faust    in10[0].v = CONSTANT_VALUE;
2206*c217d954SCole Faust    in11[0].v = CONSTANT_VALUE;
2207*c217d954SCole Faust
2208*c217d954SCole Faust#ifndef BORDER_MODE_REPLICATE
2209*c217d954SCole Faust    T_LOAD_NHWC_WITH_DILATION(SRC_DATA_TYPE, 1, 1, N0, SRC_TENSOR_TYPE, src, bout, yi, xi, cout, src_w, src_h, 1, 1, true, in00);
2210*c217d954SCole Faust    T_LOAD_NHWC_WITH_DILATION(SRC_DATA_TYPE, 1, 1, N0, SRC_TENSOR_TYPE, src, bout, yi, xi + 1, cout, src_w, src_h, 1, 1, true, in01);
2211*c217d954SCole Faust    T_LOAD_NHWC_WITH_DILATION(SRC_DATA_TYPE, 1, 1, N0, SRC_TENSOR_TYPE, src, bout, yi + 1, xi, cout, src_w, src_h, 1, 1, true, in10);
2212*c217d954SCole Faust    T_LOAD_NHWC_WITH_DILATION(SRC_DATA_TYPE, 1, 1, N0, SRC_TENSOR_TYPE, src, bout, yi + 1, xi + 1, cout, src_w, src_h, 1, 1, true, in11);
2213*c217d954SCole Faust#else
2214*c217d954SCole Faust    const int xi0  = clamp(xi, 0, (int)src_w - 1);
2215*c217d954SCole Faust    const int yi0  = clamp(yi, 0, (int)src_h - 1);
2216*c217d954SCole Faust    const int xi1  = clamp(xi + 1, 0, (int)src_w - 1);
2217*c217d954SCole Faust    const int yi1  = clamp(yi + 1, 0, (int)src_h - 1);
2218*c217d954SCole Faust
2219*c217d954SCole Faust    T_LOAD_NHWC_WITH_DILATION(SRC_DATA_TYPE, 1, 1, N0, SRC_TENSOR_TYPE, src, bout, yi0, xi0, cout, src_w, src_h, 1, 1, false, in00);
2220*c217d954SCole Faust    T_LOAD_NHWC_WITH_DILATION(SRC_DATA_TYPE, 1, 1, N0, SRC_TENSOR_TYPE, src, bout, yi0, xi1, cout, src_w, src_h, 1, 1, false, in01);
2221*c217d954SCole Faust    T_LOAD_NHWC_WITH_DILATION(SRC_DATA_TYPE, 1, 1, N0, SRC_TENSOR_TYPE, src, bout, yi1, xi0, cout, src_w, src_h, 1, 1, false, in10);
2222*c217d954SCole Faust    T_LOAD_NHWC_WITH_DILATION(SRC_DATA_TYPE, 1, 1, N0, SRC_TENSOR_TYPE, src, bout, yi1, xi1, cout, src_w, src_h, 1, 1, false, in11);
2223*c217d954SCole Faust#endif
2224*c217d954SCole Faust
2225*c217d954SCole Faust    TILE(DST_DATA_TYPE, 1, N0, out);
2226*c217d954SCole Faust
2227*c217d954SCole Faust#if defined(IS_FLOATING_POINT)
2228*c217d954SCole Faust    const SRC_DATA_TYPE a  = (SRC_DATA_TYPE)(xi_f - (float)xi);
2229*c217d954SCole Faust    const SRC_DATA_TYPE b  = (SRC_DATA_TYPE)(1.f - a);
2230*c217d954SCole Faust    const SRC_DATA_TYPE a1 = (SRC_DATA_TYPE)(yi_f - (float)yi);
2231*c217d954SCole Faust    const SRC_DATA_TYPE b1 = (SRC_DATA_TYPE)(1.f - a1);
2232*c217d954SCole Faust
2233*c217d954SCole Faust
2234*c217d954SCole Faust    out[0].v = ((in00[0].v * b * b1) + (in01[0].v * a * b1) + (in10[0].v * b * a1) + (in11[0].v * a * a1));
2235*c217d954SCole Faust#else
2236*c217d954SCole Faust
2237*c217d954SCole Faust    const float a  = (xi_f - (float)xi);
2238*c217d954SCole Faust    const float b  = (1.f - a);
2239*c217d954SCole Faust    const float a1 = (yi_f - (float)yi);
2240*c217d954SCole Faust    const float b1 = (1.f - a1);
2241*c217d954SCole Faust
2242*c217d954SCole Faust    out[0].v = CONVERT_SAT((CONVERT(in00[0].v, VEC_DATA_TYPE(float, N0)) * b * b1) +
2243*c217d954SCole Faust                           (CONVERT(in01[0].v, VEC_DATA_TYPE(float, N0)) * a * b1) +
2244*c217d954SCole Faust                           (CONVERT(in10[0].v, VEC_DATA_TYPE(float, N0)) * b * a1) +
2245*c217d954SCole Faust                           (CONVERT(in11[0].v, VEC_DATA_TYPE(float, N0)) * a * a1),
2246*c217d954SCole Faust                           VEC_DATA_TYPE(DST_DATA_TYPE, N0));
2247*c217d954SCole Faust#endif
2248*c217d954SCole Faust
2249*c217d954SCole Faust    TILE(uint, 1, 1, dst_indirect_y);
2250*c217d954SCole Faust
2251*c217d954SCole Faust
2252*c217d954SCole Faust    dst_indirect_y[0].v = xo + (yo * (int)(dst_w)) + bout * (int)(dst_w * dst_h);
2253*c217d954SCole Faust
2254*c217d954SCole Faust    bool x_cond = PARTIAL_N0 != 0 && get_global_id(0) == 0;
2255*c217d954SCole Faust
2256*c217d954SCole Faust    T_STORE_INDIRECT_WIDTH_SELECT(DST_DATA_TYPE, 1, N0, PARTIAL_N0, DST_TENSOR_TYPE, dst, cout, dst_stride_y, x_cond, out, dst_indirect_y);
2257*c217d954SCole Faust}
2258*c217d954SCole Faust#endif  )"