26 #if defined(SRC_DIM_Z)
53 __kernel
void winograd_filter_transform_2x2_3x3_nchw(
62 #if defined(WINOGRAD_FILTER_TRANSFORM_HORIZONTAL)
64 w0 = vload3(0, (__global DATA_TYPE *)(src_addr));
65 #elif defined(WINOGRAD_FILTER_TRANSFORM_VERTICAL)
67 w0 = (
VEC_DATA_TYPE(DATA_TYPE, 3))(*((__global DATA_TYPE *)(src_addr + 0 * src_stride_y)),
68 *((__global DATA_TYPE *)(src_addr + 1 * src_stride_y)),
69 *((__global DATA_TYPE *)(src_addr + 2 * src_stride_y)));
70 #else // defined(WINOGRAD_FILTER_TRANSFORM_VERTICAL)
72 w0 = vload3(0, (__global DATA_TYPE *)(src_addr + 0 * src_stride_y));
74 w1 = vload3(0, (__global DATA_TYPE *)(src_addr + 1 * src_stride_y));
76 w2 = vload3(0, (__global DATA_TYPE *)(src_addr + 2 * src_stride_y));
77 #endif // defined(WINOGRAD_FILTER_TRANSFORM_HORIZONTAL)
83 out0.s1 = (w0.s0 + w0.s1 + w0.s2) * 0.5f;
84 out0.s2 = (w0.s0 + w0.s2 - w0.s1) * 0.5f;
87 #if !defined(WINOGRAD_FILTER_TRANSFORM_HORIZONTAL) && !defined(WINOGRAD_FILTER_TRANSFORM_VERTICAL)
91 out1.s0 = (w0.s0 + w1.s0 + w2.s0) * 0.5f;
92 out1.s1 = (w0.s0 + w1.s0 + w2.s0 + w0.s1 + w1.s1 + w2.s1 + w0.s2 + w1.s2 + w2.s2) * 0.25f;
93 out1.s2 = (w0.s0 + w1.s0 + w2.s0 + w0.s2 + w1.s2 + w2.s2 - w0.s1 - w1.s1 - w2.s1) * 0.25f;
94 out1.s3 = (w0.s2 + w1.s2 + w2.s2) * 0.5f;
99 out2.s0 = (w0.s0 + w2.s0 - w1.s0) * 0.5f;
100 out2.s1 = (w0.s0 + w2.s0 + w0.s1 + w2.s1 + w0.s2 + w2.s2 - w1.s0 - w1.s1 - w1.s2) * 0.25f;
101 out2.s2 = (w0.s0 + w2.s0 + w1.s1 + w0.s2 + w2.s2 - w1.s0 - w0.s1 - w2.s1 - w1.s2) * 0.25f;
102 out2.s3 = (w0.s2 + w2.s2 - w1.s2) * 0.5f;
108 out3.s1 = (w2.s0 + w2.s1 + w2.s2) * 0.5f;
109 out3.s2 = (w2.s0 + w2.s2 - w2.s1) * 0.5f;
111 #endif // !defined(WINOGRAD_FILTER_TRANSFORM_HORIZONTAL) && !defined(WINOGRAD_FILTER_TRANSFORM_VERTICAL)
113 int z = get_global_id(2);
114 int x0 = z / SRC_DIM_Z;
115 int y0 = z % SRC_DIM_Z;
118 __global uchar *dst_addr = dst_ptr + dst_offset_first_element_in_bytes + x0 * dst_stride_x + y0 * dst_stride_y;
123 *(__global DATA_TYPE *)(dst_addr + 0 * dst_stride_z) = out0.s0;
124 *(__global DATA_TYPE *)(dst_addr + 1 * dst_stride_z) = out0.s1;
125 *(__global DATA_TYPE *)(dst_addr + 2 * dst_stride_z) = out0.s2;
126 *(__global DATA_TYPE *)(dst_addr + 3 * dst_stride_z) = out0.s3;
128 #if !defined(WINOGRAD_FILTER_TRANSFORM_HORIZONTAL) && !defined(WINOGRAD_FILTER_TRANSFORM_VERTICAL)
129 *(__global DATA_TYPE *)(dst_addr + 4 * dst_stride_z) = out1.s0;
130 *(__global DATA_TYPE *)(dst_addr + 5 * dst_stride_z) = out1.s1;
131 *(__global DATA_TYPE *)(dst_addr + 6 * dst_stride_z) = out1.s2;
132 *(__global DATA_TYPE *)(dst_addr + 7 * dst_stride_z) = out1.s3;
133 *(__global DATA_TYPE *)(dst_addr + 8 * dst_stride_z) = out2.s0;
134 *(__global DATA_TYPE *)(dst_addr + 9 * dst_stride_z) = out2.s1;
135 *(__global DATA_TYPE *)(dst_addr + 10 * dst_stride_z) = out2.s2;
136 *(__global DATA_TYPE *)(dst_addr + 11 * dst_stride_z) = out2.s3;
137 *(__global DATA_TYPE *)(dst_addr + 12 * dst_stride_z) = out3.s0;
138 *(__global DATA_TYPE *)(dst_addr + 13 * dst_stride_z) = out3.s1;
139 *(__global DATA_TYPE *)(dst_addr + 14 * dst_stride_z) = out3.s2;
140 *(__global DATA_TYPE *)(dst_addr + 15 * dst_stride_z) = out3.s3;
141 #endif // !defined(WINOGRAD_FILTER_TRANSFORM_HORIZONTAL) && !defined(WINOGRAD_FILTER_TRANSFORM_VERTICAL)
170 __kernel
void winograd_filter_transform_4x4_3x3_nchw(
179 #if defined(WINOGRAD_FILTER_TRANSFORM_HORIZONTAL)
181 w0 = vload3(0, (__global DATA_TYPE *)(src_addr));
182 #elif defined(WINOGRAD_FILTER_TRANSFORM_VERTICAL)
184 w0 = (
VEC_DATA_TYPE(DATA_TYPE, 3))(*((__global DATA_TYPE *)(src_addr + 0 * src_stride_y)),
185 *((__global DATA_TYPE *)(src_addr + 1 * src_stride_y)),
186 *((__global DATA_TYPE *)(src_addr + 2 * src_stride_y)));
187 #else // defined(WINOGRAD_FILTER_TRANSFORM_VERTICAL)
189 w0 = vload3(0, (__global DATA_TYPE *)(src_addr + 0 * src_stride_y));
191 w1 = vload3(0, (__global DATA_TYPE *)(src_addr + 1 * src_stride_y));
193 w2 = vload3(0, (__global DATA_TYPE *)(src_addr + 2 * src_stride_y));
194 #endif // defined(WINOGRAD_FILTER_TRANSFORM_HORIZONTAL)
199 out0.s0 = (w0.s0) / 16.f;
200 out0.s1 = (-w0.s0 - w0.s1 - w0.s2) / 24.f;
201 out0.s2 = (-w0.s0 + w0.s1 - w0.s2) / 24.f;
202 out0.s3 = (w0.s0 + 2.f * w0.s1 + 4.f * w0.s2) / 96.f;
203 out0.s4 = (w0.s0 - 2.f * w0.s1 + 4.f * w0.s2) / 96.f;
204 out0.s5 = (w0.s2) / 4.f;
206 #if !defined(WINOGRAD_FILTER_TRANSFORM_HORIZONTAL) && !defined(WINOGRAD_FILTER_TRANSFORM_VERTICAL)
210 out1.s0 = (-w0.s0 - w1.s0 - w2.s0) / 24.f;
211 out1.s1 = (w0.s0 + w1.s0 + w2.s0 + w0.s1 + w1.s1 + w2.s1 + w0.s2 + w1.s2 + w2.s2) / 36.f;
212 out1.s2 = (w0.s0 + w1.s0 + w2.s0 - w0.s1 - w1.s1 - w2.s1 + w0.s2 + w1.s2 + w2.s2) / 36.f;
213 out1.s3 = (-w0.s0 - w1.s0 - w2.s0 + 2.f * (-w0.s1 - w1.s1 - w2.s1) + 4.f * (-w0.s2 - w1.s2 - w2.s2)) / 144.f;
214 out1.s4 = (-w0.s0 - w1.s0 - w2.s0 + 2.f * (w0.s1 + w1.s1 + w2.s1) + 4.f * (-w0.s2 - w1.s2 - w2.s2)) / 144.f;
215 out1.s5 = (-w0.s2 - w1.s2 - w2.s2) / 6.f;
220 out2.s0 = (-w0.s0 + w1.s0 - w2.s0) / 24.f;
221 out2.s1 = (w0.s0 - w1.s0 + w2.s0 + w0.s1 - w1.s1 + w2.s1 + w0.s2 - w1.s2 + w2.s2) / 36.f;
222 out2.s2 = (w0.s0 - w1.s0 + w2.s0 - w0.s1 + w1.s1 - w2.s1 + w0.s2 - w1.s2 + w2.s2) / 36.f;
223 out2.s3 = (-w0.s0 + w1.s0 - w2.s0 + 2.f * (-w0.s1 + w1.s1 - w2.s1) + 4.f * (-w0.s2 + w1.s2 - w2.s2)) / 144.f;
224 out2.s4 = (-w0.s0 + w1.s0 - w2.s0 + 2.f * (w0.s1 - w1.s1 + w2.s1) + 4.f * (-w0.s2 + w1.s2 - w2.s2)) / 144.f;
225 out2.s5 = (-w0.s2 + w1.s2 - w2.s2) / 6.f;
230 out3.s0 = (w0.s0 + 2.f * w1.s0 + 4.f * w2.s0) / 96.f;
231 out3.s1 = (-w0.s0 - 2.f * w1.s0 - 4.f * w2.s0 - w0.s1 - 2.f * w1.s1 - 4.f * w2.s1 - w0.s2 - 2.f * w1.s2 - 4.f * w2.s2) / 144.f;
232 out3.s2 = (-w0.s0 - 2.f * w1.s0 - 4.f * w2.s0 + w0.s1 + 2.f * w1.s1 + 4.f * w2.s1 - w0.s2 - 2.f * w1.s2 - 4.f * w2.s2) / 144.f;
233 out3.s3 = ((w0.s0 + 2.f * w1.s0 + 4.f * w2.s0) + 2.f * (w0.s1 + 2.f * w1.s1 + 4.f * w2.s1) + 4.f * (w0.s2 + 2.f * w1.s2 + 4.f * w2.s2)) / 576.f;
234 out3.s4 = ((w0.s0 + 2.f * w1.s0 + 4.f * w2.s0) + 2.f * (-w0.s1 - 2.f * w1.s1 - 4.f * w2.s1) + 4.f * (w0.s2 + 2.f * w1.s2 + 4.f * w2.s2)) / 576.f;
235 out3.s5 = (w0.s2 + 2.f * w1.s2 + 4.f * w2.s2) / 24.f;
240 out4.s0 = (w0.s0 - 2.f * w1.s0 + 4.f * w2.s0) / 96.f;
241 out4.s1 = (-w0.s0 + 2.f * w1.s0 - 4.f * w2.s0 - w0.s1 + 2.f * w1.s1 - 4.f * w2.s1 - w0.s2 + 2.f * w1.s2 - 4.f * w2.s2) / 144.f;
242 out4.s2 = (-w0.s0 + 2.f * w1.s0 - 4.f * w2.s0 + w0.s1 - 2.f * w1.s1 + 4.f * w2.s1 - w0.s2 + 2.f * w1.s2 - 4.f * w2.s2) / 144.f;
243 out4.s3 = ((w0.s0 - 2.f * w1.s0 + 4.f * w2.s0) + 2.f * (w0.s1 - 2.f * w1.s1 + 4.f * w2.s1) + 4.f * (w0.s2 - 2.f * w1.s2 + 4.f * w2.s2)) / 576.f;
244 out4.s4 = ((w0.s0 - 2.f * w1.s0 + 4.f * w2.s0) + 2.f * (-w0.s1 + 2.f * w1.s1 - 4.f * w2.s1) + 4.f * (w0.s2 - 2.f * w1.s2 + 4.f * w2.s2)) / 576.f;
245 out4.s5 = (w0.s2 - 2.f * w1.s2 + 4.f * w2.s2) / 24.f;
250 out5.s0 = (w2.s0) / 4.f;
251 out5.s1 = (-w2.s0 - w2.s1 - w2.s2) / 6.f;
252 out5.s2 = (-w2.s0 + w2.s1 - w2.s2) / 6.f;
253 out5.s3 = (w2.s0 + 2.f * w2.s1 + 4.f * w2.s2) / 24.f;
254 out5.s4 = (w2.s0 - 2.f * w2.s1 + 4.f * w2.s2) / 24.f;
256 #endif // !defined(WINOGRAD_FILTER_TRANSFORM_HORIZONTAL) && !defined(WINOGRAD_FILTER_TRANSFORM_VERTICAL)
258 int z = get_global_id(2);
259 int x0 = z / SRC_DIM_Z;
260 int y0 = z % SRC_DIM_Z;
263 __global uchar *dst_addr = dst_ptr + dst_offset_first_element_in_bytes + x0 * dst_stride_x + y0 * dst_stride_y;
268 *(__global DATA_TYPE *)(dst_addr + 0 * dst_stride_z) = out0.s0;
269 *(__global DATA_TYPE *)(dst_addr + 1 * dst_stride_z) = out0.s1;
270 *(__global DATA_TYPE *)(dst_addr + 2 * dst_stride_z) = out0.s2;
271 *(__global DATA_TYPE *)(dst_addr + 3 * dst_stride_z) = out0.s3;
272 *(__global DATA_TYPE *)(dst_addr + 4 * dst_stride_z) = out0.s4;
273 *(__global DATA_TYPE *)(dst_addr + 5 * dst_stride_z) = out0.s5;
275 #if !defined(WINOGRAD_FILTER_TRANSFORM_HORIZONTAL) && !defined(WINOGRAD_FILTER_TRANSFORM_VERTICAL)
276 *(__global DATA_TYPE *)(dst_addr + 6 * dst_stride_z) = out1.s0;
277 *(__global DATA_TYPE *)(dst_addr + 7 * dst_stride_z) = out1.s1;
278 *(__global DATA_TYPE *)(dst_addr + 8 * dst_stride_z) = out1.s2;
279 *(__global DATA_TYPE *)(dst_addr + 9 * dst_stride_z) = out1.s3;
280 *(__global DATA_TYPE *)(dst_addr + 10 * dst_stride_z) = out1.s4;
281 *(__global DATA_TYPE *)(dst_addr + 11 * dst_stride_z) = out1.s5;
282 *(__global DATA_TYPE *)(dst_addr + 12 * dst_stride_z) = out2.s0;
283 *(__global DATA_TYPE *)(dst_addr + 13 * dst_stride_z) = out2.s1;
284 *(__global DATA_TYPE *)(dst_addr + 14 * dst_stride_z) = out2.s2;
285 *(__global DATA_TYPE *)(dst_addr + 15 * dst_stride_z) = out2.s3;
286 *(__global DATA_TYPE *)(dst_addr + 16 * dst_stride_z) = out2.s4;
287 *(__global DATA_TYPE *)(dst_addr + 17 * dst_stride_z) = out2.s5;
288 *(__global DATA_TYPE *)(dst_addr + 18 * dst_stride_z) = out3.s0;
289 *(__global DATA_TYPE *)(dst_addr + 19 * dst_stride_z) = out3.s1;
290 *(__global DATA_TYPE *)(dst_addr + 20 * dst_stride_z) = out3.s2;
291 *(__global DATA_TYPE *)(dst_addr + 21 * dst_stride_z) = out3.s3;
292 *(__global DATA_TYPE *)(dst_addr + 22 * dst_stride_z) = out3.s4;
293 *(__global DATA_TYPE *)(dst_addr + 23 * dst_stride_z) = out3.s5;
294 *(__global DATA_TYPE *)(dst_addr + 24 * dst_stride_z) = out4.s0;
295 *(__global DATA_TYPE *)(dst_addr + 25 * dst_stride_z) = out4.s1;
296 *(__global DATA_TYPE *)(dst_addr + 26 * dst_stride_z) = out4.s2;
297 *(__global DATA_TYPE *)(dst_addr + 27 * dst_stride_z) = out4.s3;
298 *(__global DATA_TYPE *)(dst_addr + 28 * dst_stride_z) = out4.s4;
299 *(__global DATA_TYPE *)(dst_addr + 29 * dst_stride_z) = out4.s5;
300 *(__global DATA_TYPE *)(dst_addr + 30 * dst_stride_z) = out5.s0;
301 *(__global DATA_TYPE *)(dst_addr + 31 * dst_stride_z) = out5.s1;
302 *(__global DATA_TYPE *)(dst_addr + 32 * dst_stride_z) = out5.s2;
303 *(__global DATA_TYPE *)(dst_addr + 33 * dst_stride_z) = out5.s3;
304 *(__global DATA_TYPE *)(dst_addr + 34 * dst_stride_z) = out5.s4;
305 *(__global DATA_TYPE *)(dst_addr + 35 * dst_stride_z) = out5.s5;
306 #endif // !defined(WINOGRAD_FILTER_TRANSFORM_HORIZONTAL) && !defined(WINOGRAD_FILTER_TRANSFORM_VERTICAL)
336 __kernel
void winograd_filter_transform_4x4_5x5_nchw(
345 #if defined(WINOGRAD_FILTER_TRANSFORM_HORIZONTAL)
347 w00 = vload4(0, (__global DATA_TYPE *)(src_addr + 0 * src_stride_y));
348 DATA_TYPE w01 = *((__global DATA_TYPE *)(src_addr + 0 * src_stride_y) + 4);
349 #elif defined(WINOGRAD_FILTER_TRANSFORM_VERTICAL)
351 w00 = (
VEC_DATA_TYPE(DATA_TYPE, 4))(*((__global DATA_TYPE *)(src_addr + 0 * src_stride_y)),
352 *((__global DATA_TYPE *)(src_addr + 1 * src_stride_y)),
353 *((__global DATA_TYPE *)(src_addr + 2 * src_stride_y)),
354 *((__global DATA_TYPE *)(src_addr + 3 * src_stride_y)));
355 DATA_TYPE w01 = *((__global DATA_TYPE *)(src_addr + 4 * src_stride_y));
356 #else // defined(WINOGRAD_FILTER_TRANSFORM_VERTICAL)
358 w00 = vload4(0, (__global DATA_TYPE *)(src_addr + 0 * src_stride_y));
359 DATA_TYPE w01 = *((__global DATA_TYPE *)(src_addr + 0 * src_stride_y) + 4);
361 w10 = vload4(0, (__global DATA_TYPE *)(src_addr + 1 * src_stride_y));
362 DATA_TYPE w11 = *((__global DATA_TYPE *)(src_addr + 1 * src_stride_y) + 4);
364 w20 = vload4(0, (__global DATA_TYPE *)(src_addr + 2 * src_stride_y));
365 DATA_TYPE w21 = *((__global DATA_TYPE *)(src_addr + 2 * src_stride_y) + 4);
367 w30 = vload4(0, (__global DATA_TYPE *)(src_addr + 3 * src_stride_y));
368 DATA_TYPE w31 = *((__global DATA_TYPE *)(src_addr + 3 * src_stride_y) + 4);
370 w40 = vload4(0, (__global DATA_TYPE *)(src_addr + 4 * src_stride_y));
371 DATA_TYPE w41 = *((__global DATA_TYPE *)(src_addr + 4 * src_stride_y) + 4);
372 #endif // defined(WINOGRAD_FILTER_TRANSFORM_HORIZONTAL)
380 out0.s1 = -2.f * (w00.s0 + w00.s1 + w00.s2 + w00.s3 + w01) / 9.f;
381 out0.s2 = -2.f * (w00.s0 - w00.s1 + w00.s2 - w00.s3 + w01) / 9.f;
382 out0.s3 = (w00.s0 + 2.f * w00.s1 + 4.f * w00.s2 + 8.f * w00.s3 + 16.f * w01) / 90.f;
383 out0.s4 = (w00.s0 - 2.f * w00.s1 + 4.f * w00.s2 - 8.f * w00.s3 + 16.f * w01) / 90.f;
384 out0.s5 = (16.f * w00.s0 + 8.f * w00.s1 + 4.f * w00.s2 + 2.f * w00.s3 + w01) / 180.f;
385 out0.s6 = (16.f * w00.s0 - 8.f * w00.s1 + 4.f * w00.s2 - 2.f * w00.s3 + w01) / 180.f;
388 #if !defined(WINOGRAD_FILTER_TRANSFORM_HORIZONTAL) && !defined(WINOGRAD_FILTER_TRANSFORM_VERTICAL)
392 out1.s0 = -2.f * (w00.s0 + w10.s0 + w20.s0 + w30.s0 + w40.s0) / 9.f;
393 out1.s1 = 4.f * ((w00.s0 + w10.s0 + w20.s0 + w30.s0 + w40.s0) + (w00.s1 + w10.s1 + w20.s1 + w30.s1 + w40.s1) + (w00.s2 + w10.s2 + w20.s2 + w30.s2 + w40.s2) +
394 (w00.s3 + w10.s3 + w20.s3 + w30.s3 + w40.s3) + (w01 + w11 + w21 + w31 + w41)) / 81.f;
395 out1.s2 = 4.f * ((w00.s0 + w10.s0 + w20.s0 + w30.s0 + w40.s0) - (w00.s1 + w10.s1 + w20.s1 + w30.s1 + w40.s1) + (w00.s2 + w10.s2 + w20.s2 + w30.s2 + w40.s2) -
396 (w00.s3 + w10.s3 + w20.s3 + w30.s3 + w40.s3) + (w01 + w11 + w21 + w31 + w41)) / 81.f;
397 out1.s3 = -((w00.s0 + w10.s0 + w20.s0 + w30.s0 + w40.s0) + 2.f * (w00.s1 + w10.s1 + w20.s1 + w30.s1 + w40.s1) + 4.f * (w00.s2 + w10.s2 + w20.s2 + w30.s2 + w40.s2) + 8.f *
398 (w00.s3 + w10.s3 + w20.s3 + w30.s3 + w40.s3) + 16.f * (w01 + w11 + w21 + w31 + w41)) / 405.f;
399 out1.s4 = -((w00.s0 + w10.s0 + w20.s0 + w30.s0 + w40.s0) - 2.f * (w00.s1 + w10.s1 + w20.s1 + w30.s1 + w40.s1) + 4.f * (w00.s2 + w10.s2 + w20.s2 + w30.s2 + w40.s2) - 8.f *
400 (w00.s3 + w10.s3 + w20.s3 + w30.s3 + w40.s3) + 16.f * (w01 + w11 + w21 + w31 + w41)) / 405.f;
401 out1.s5 = -(16.f * (w00.s0 + w10.s0 + w20.s0 + w30.s0 + w40.s0) + 8.f * (w00.s1 + w10.s1 + w20.s1 + w30.s1 + w40.s1) + 4.f * (w00.s2 + w10.s2 + w20.s2 + w30.s2 + w40.s2) + 2.f *
402 (w00.s3 + w10.s3 + w20.s3 + w30.s3 + w40.s3) + (w01 + w11 + w21 + w31 + w41)) / 810.f;
403 out1.s6 = -(16.f * (w00.s0 + w10.s0 + w20.s0 + w30.s0 + w40.s0) - 8.f * (w00.s1 + w10.s1 + w20.s1 + w30.s1 + w40.s1) + 4.f * (w00.s2 + w10.s2 + w20.s2 + w30.s2 + w40.s2) - 2.f *
404 (w00.s3 + w10.s3 + w20.s3 + w30.s3 + w40.s3) + (w01 + w11 + w21 + w31 + w41)) / 810.f;
405 out1.s7 = -2.f * (w01 + w11 + w21 + w31 + w41) / 9.f;
410 out2.s0 = -2.f * (w00.s0 - w10.s0 + w20.s0 - w30.s0 + w40.s0) / 9.f;
411 out2.s1 = 4.f * ((w00.s0 - w10.s0 + w20.s0 - w30.s0 + w40.s0) + (w00.s1 - w10.s1 + w20.s1 - w30.s1 + w40.s1) + (w00.s2 - w10.s2 + w20.s2 - w30.s2 + w40.s2) +
412 (w00.s3 - w10.s3 + w20.s3 - w30.s3 + w40.s3) + (w01 - w11 + w21 - w31 + w41)) / 81.f;
413 out2.s2 = 4.f * ((w00.s0 - w10.s0 + w20.s0 - w30.s0 + w40.s0) - (w00.s1 - w10.s1 + w20.s1 - w30.s1 + w40.s1) + (w00.s2 - w10.s2 + w20.s2 - w30.s2 + w40.s2) -
414 (w00.s3 - w10.s3 + w20.s3 - w30.s3 + w40.s3) + (w01 - w11 + w21 - w31 + w41)) / 81.f;
415 out2.s3 = -((w00.s0 - w10.s0 + w20.s0 - w30.s0 + w40.s0) + 2.f * (w00.s1 - w10.s1 + w20.s1 - w30.s1 + w40.s1) + 4.f * (w00.s2 - w10.s2 + w20.s2 - w30.s2 + w40.s2) + 8.f *
416 (w00.s3 - w10.s3 + w20.s3 - w30.s3 + w40.s3) + 16.f * (w01 - w11 + w21 - w31 + w41)) / 405.f;
417 out2.s4 = -((w00.s0 - w10.s0 + w20.s0 - w30.s0 + w40.s0) - 2.f * (w00.s1 - w10.s1 + w20.s1 - w30.s1 + w40.s1) + 4.f * (w00.s2 - w10.s2 + w20.s2 - w30.s2 + w40.s2) - 8.f *
418 (w00.s3 - w10.s3 + w20.s3 - w30.s3 + w40.s3) + 16.f * (w01 - w11 + w21 - w31 + w41)) / 405.f;
419 out2.s5 = -(16.f * (w00.s0 - w10.s0 + w20.s0 - w30.s0 + w40.s0) + 8.f * (w00.s1 - w10.s1 + w20.s1 - w30.s1 + w40.s1) + 4.f * (w00.s2 - w10.s2 + w20.s2 - w30.s2 + w40.s2) + 2.f *
420 (w00.s3 - w10.s3 + w20.s3 - w30.s3 + w40.s3) + (w01 - w11 + w21 - w31 + w41)) / 810.f;
421 out2.s6 = -(16.f * (w00.s0 - w10.s0 + w20.s0 - w30.s0 + w40.s0) - 8.f * (w00.s1 - w10.s1 + w20.s1 - w30.s1 + w40.s1) + 4.f * (w00.s2 - w10.s2 + w20.s2 - w30.s2 + w40.s2) - 2.f *
422 (w00.s3 - w10.s3 + w20.s3 - w30.s3 + w40.s3) + (w01 - w11 + w21 - w31 + w41)) / 810.f;
423 out2.s7 = -2.f * (w01 - w11 + w21 - w31 + w41) / 9.f;
428 out3.s0 = (w00.s0 + 2.f * w10.s0 + 4.f * w20.s0 + 8.f * w30.s0 + 16.f * w40.s0) / 90.f;
429 out3.s1 = -((w00.s0 + 2.f * w10.s0 + 4.f * w20.s0 + 8.f * w30.s0 + 16.f * w40.s0) + (w00.s1 + 2.f * w10.s1 + 4.f * w20.s1 + 8.f * w30.s1 + 16.f * w40.s1) +
430 (w00.s2 + 2.f * w10.s2 + 4.f * w20.s2 + 8.f * w30.s2 + 16.f * w40.s2) + (w00.s3 + 2.f * w10.s3 + 4.f * w20.s3 + 8.f * w30.s3 + 16.f * w40.s3) +
431 (w01 + 2.f * w11 + 4.f * w21 + 8.f * w31 + 16.f * w41)) / 405.f;
432 out3.s2 = -((w00.s0 + 2.f * w10.s0 + 4.f * w20.s0 + 8.f * w30.s0 + 16.f * w40.s0) - (w00.s1 + 2.f * w10.s1 + 4.f * w20.s1 + 8.f * w30.s1 + 16.f * w40.s1) +
433 (w00.s2 + 2.f * w10.s2 + 4.f * w20.s2 + 8.f * w30.s2 + 16.f * w40.s2) - (w00.s3 + 2.f * w10.s3 + 4.f * w20.s3 + 8.f * w30.s3 + 16.f * w40.s3) +
434 (w01 + 2.f * w11 + 4.f * w21 + 8.f * w31 + 16.f * w41)) / 405.f;
435 out3.s3 = ((w00.s0 + 2.f * w10.s0 + 4.f * w20.s0 + 8.f * w30.s0 + 16.f * w40.s0) + 2.f * (w00.s1 + 2.f * w10.s1 + 4.f * w20.s1 + 8.f * w30.s1 + 16.f * w40.s1) + 4.f *
436 (w00.s2 + 2.f * w10.s2 + 4.f * w20.s2 + 8.f * w30.s2 + 16.f * w40.s2) + 8.f * (w00.s3 + 2.f * w10.s3 + 4.f * w20.s3 + 8.f * w30.s3 + 16.f * w40.s3) + 16.f *
437 (w01 + 2.f * w11 + 4.f * w21 + 8.f * w31 + 16.f * w41)) / 8100.f;
438 out3.s4 = ((w00.s0 + 2.f * w10.s0 + 4.f * w20.s0 + 8.f * w30.s0 + 16.f * w40.s0) - 2.f * (w00.s1 + 2.f * w10.s1 + 4.f * w20.s1 + 8.f * w30.s1 + 16.f * w40.s1) + 4.f *
439 (w00.s2 + 2.f * w10.s2 + 4.f * w20.s2 + 8.f * w30.s2 + 16.f * w40.s2) - 8.f * (w00.s3 + 2.f * w10.s3 + 4.f * w20.s3 + 8.f * w30.s3 + 16.f * w40.s3) + 16.f *
440 (w01 + 2.f * w11 + 4.f * w21 + 8.f * w31 + 16.f * w41)) / 8100.f;
441 out3.s5 = (16.f * (w00.s0 + 2.f * w10.s0 + 4.f * w20.s0 + 8.f * w30.s0 + 16.f * w40.s0) + 8.f * (w00.s1 + 2.f * w10.s1 + 4.f * w20.s1 + 8.f * w30.s1 + 16.f * w40.s1) + 4.f *
442 (w00.s2 + 2.f * w10.s2 + 4.f * w20.s2 + 8.f * w30.s2 + 16.f * w40.s2) + 2.f * (w00.s3 + 2.f * w10.s3 + 4.f * w20.s3 + 8.f * w30.s3 + 16.f * w40.s3) +
443 (w01 + 2.f * w11 + 4.f * w21 + 8.f * w31 + 16.f * w41)) / 16200.f;
444 out3.s6 = (16.f * (w00.s0 + 2.f * w10.s0 + 4.f * w20.s0 + 8.f * w30.s0 + 16.f * w40.s0) - 8.f * (w00.s1 + 2.f * w10.s1 + 4.f * w20.s1 + 8.f * w30.s1 + 16.f * w40.s1) + 4.f *
445 (w00.s2 + 2.f * w10.s2 + 4.f * w20.s2 + 8.f * w30.s2 + 16.f * w40.s2) - 2.f * (w00.s3 + 2.f * w10.s3 + 4.f * w20.s3 + 8.f * w30.s3 + 16.f * w40.s3) +
446 (w01 + 2.f * w11 + 4.f * w21 + 8.f * w31 + 16.f * w41)) / 16200.f;
447 out3.s7 = (w01 + 2.f * w11 + 4.f * w21 + 8.f * w31 + 16.f * w41) / 90.f;
452 out4.s0 = (w00.s0 - 2.f * w10.s0 + 4.f * w20.s0 - 8.f * w30.s0 + 16.f * w40.s0) / 90.f;
453 out4.s1 = -((w00.s0 - 2.f * w10.s0 + 4.f * w20.s0 - 8.f * w30.s0 + 16.f * w40.s0) + (w00.s1 - 2.f * w10.s1 + 4.f * w20.s1 - 8.f * w30.s1 + 16.f * w40.s1) +
454 (w00.s2 - 2.f * w10.s2 + 4.f * w20.s2 - 8.f * w30.s2 + 16.f * w40.s2) + (w00.s3 - 2.f * w10.s3 + 4.f * w20.s3 - 8.f * w30.s3 + 16.f * w40.s3) +
455 (w01 - 2.f * w11 + 4.f * w21 - 8.f * w31 + 16.f * w41)) / 405.f;
456 out4.s2 = -((w00.s0 - 2.f * w10.s0 + 4.f * w20.s0 - 8.f * w30.s0 + 16.f * w40.s0) - (w00.s1 - 2.f * w10.s1 + 4.f * w20.s1 - 8.f * w30.s1 + 16.f * w40.s1) +
457 (w00.s2 - 2.f * w10.s2 + 4.f * w20.s2 - 8.f * w30.s2 + 16.f * w40.s2) - (w00.s3 - 2.f * w10.s3 + 4.f * w20.s3 - 8.f * w30.s3 + 16.f * w40.s3) +
458 (w01 - 2.f * w11 + 4.f * w21 - 8.f * w31 + 16.f * w41)) / 405.f;
459 out4.s3 = ((w00.s0 - 2.f * w10.s0 + 4.f * w20.s0 - 8.f * w30.s0 + 16.f * w40.s0) + 2.f * (w00.s1 - 2.f * w10.s1 + 4.f * w20.s1 - 8.f * w30.s1 + 16.f * w40.s1) + 4.f *
460 (w00.s2 - 2.f * w10.s2 + 4.f * w20.s2 - 8.f * w30.s2 + 16.f * w40.s2) + 8.f * (w00.s3 - 2.f * w10.s3 + 4.f * w20.s3 - 8.f * w30.s3 + 16.f * w40.s3) + 16.f *
461 (w01 - 2.f * w11 + 4.f * w21 - 8.f * w31 + 16.f * w41)) / 8100.f;
462 out4.s4 = ((w00.s0 - 2.f * w10.s0 + 4.f * w20.s0 - 8.f * w30.s0 + 16.f * w40.s0) - 2.f * (w00.s1 - 2.f * w10.s1 + 4.f * w20.s1 - 8.f * w30.s1 + 16.f * w40.s1) + 4.f *
463 (w00.s2 - 2.f * w10.s2 + 4.f * w20.s2 - 8.f * w30.s2 + 16.f * w40.s2) - 8.f * (w00.s3 - 2.f * w10.s3 + 4.f * w20.s3 - 8.f * w30.s3 + 16.f * w40.s3) + 16.f *
464 (w01 - 2.f * w11 + 4.f * w21 - 8.f * w31 + 16.f * w41)) / 8100.f;
465 out4.s5 = (16.f * (w00.s0 - 2.f * w10.s0 + 4.f * w20.s0 - 8.f * w30.s0 + 16.f * w40.s0) + 8.f * (w00.s1 - 2.f * w10.s1 + 4.f * w20.s1 - 8.f * w30.s1 + 16.f * w40.s1) + 4.f *
466 (w00.s2 - 2.f * w10.s2 + 4.f * w20.s2 - 8.f * w30.s2 + 16.f * w40.s2) + 2.f * (w00.s3 - 2.f * w10.s3 + 4.f * w20.s3 - 8.f * w30.s3 + 16.f * w40.s3) +
467 (w01 - 2.f * w11 + 4.f * w21 - 8.f * w31 + 16.f * w41)) / 16200.f;
468 out4.s6 = (16.f * (w00.s0 - 2.f * w10.s0 + 4.f * w20.s0 - 8.f * w30.s0 + 16.f * w40.s0) - 8.f * (w00.s1 - 2.f * w10.s1 + 4.f * w20.s1 - 8.f * w30.s1 + 16.f * w40.s1) + 4.f *
469 (w00.s2 - 2.f * w10.s2 + 4.f * w20.s2 - 8.f * w30.s2 + 16.f * w40.s2) - 2.f * (w00.s3 - 2.f * w10.s3 + 4.f * w20.s3 - 8.f * w30.s3 + 16.f * w40.s3) +
470 (w01 - 2.f * w11 + 4.f * w21 - 8.f * w31 + 16.f * w41)) / 16200.f;
471 out4.s7 = (w01 - 2.f * w11 + 4.f * w21 - 8.f * w31 + 16.f * w41) / 90.f;
476 out5.s0 = (16.f * w00.s0 + 8.f * w10.s0 + 4.f * w20.s0 + 2.f * w30.s0 + w40.s0) / 180.f;
477 out5.s1 = -((16.f * w00.s0 + 8.f * w10.s0 + 4.f * w20.s0 + 2.f * w30.s0 + w40.s0) + (16.f * w00.s1 + 8.f * w10.s1 + 4.f * w20.s1 + 2.f * w30.s1 + w40.s1) +
478 (16.f * w00.s2 + 8.f * w10.s2 + 4.f * w20.s2 + 2.f * w30.s2 + w40.s2) + (16.f * w00.s3 + 8.f * w10.s3 + 4.f * w20.s3 + 2.f * w30.s3 + w40.s3) +
479 (16.f * w01 + 8.f * w11 + 4.f * w21 + 2.f * w31 + w41)) / 810.f;
480 out5.s2 = -((16.f * w00.s0 + 8.f * w10.s0 + 4.f * w20.s0 + 2.f * w30.s0 + w40.s0) - (16.f * w00.s1 + 8.f * w10.s1 + 4.f * w20.s1 + 2.f * w30.s1 + w40.s1) +
481 (16.f * w00.s2 + 8.f * w10.s2 + 4.f * w20.s2 + 2.f * w30.s2 + w40.s2) - (16.f * w00.s3 + 8.f * w10.s3 + 4.f * w20.s3 + 2.f * w30.s3 + w40.s3) +
482 (16.f * w01 + 8.f * w11 + 4.f * w21 + 2.f * w31 + w41)) / 810.f;
483 out5.s3 = ((16.f * w00.s0 + 8.f * w10.s0 + 4.f * w20.s0 + 2.f * w30.s0 + w40.s0) + 2.f * (16.f * w00.s1 + 8.f * w10.s1 + 4.f * w20.s1 + 2.f * w30.s1 + w40.s1) + 4.f *
484 (16.f * w00.s2 + 8.f * w10.s2 + 4.f * w20.s2 + 2.f * w30.s2 + w40.s2) + 8.f * (16.f * w00.s3 + 8.f * w10.s3 + 4.f * w20.s3 + 2.f * w30.s3 + w40.s3) + 16.f *
485 (16.f * w01 + 8.f * w11 + 4.f * w21 + 2.f * w31 + w41)) / 16200.f;
486 out5.s4 = ((16.f * w00.s0 + 8.f * w10.s0 + 4.f * w20.s0 + 2.f * w30.s0 + w40.s0) - 2.f * (16.f * w00.s1 + 8.f * w10.s1 + 4.f * w20.s1 + 2.f * w30.s1 + w40.s1) + 4.f *
487 (16.f * w00.s2 + 8.f * w10.s2 + 4.f * w20.s2 + 2.f * w30.s2 + w40.s2) - 8.f * (16.f * w00.s3 + 8.f * w10.s3 + 4.f * w20.s3 + 2.f * w30.s3 + w40.s3) + 16.f *
488 (16.f * w01 + 8.f * w11 + 4.f * w21 + 2.f * w31 + w41)) / 16200.f;
489 out5.s5 = (16.f * (16.f * w00.s0 + 8.f * w10.s0 + 4.f * w20.s0 + 2.f * w30.s0 + w40.s0) + 8.f * (16.f * w00.s1 + 8.f * w10.s1 + 4.f * w20.s1 + 2.f * w30.s1 + w40.s1) + 4.f *
490 (16.f * w00.s2 + 8.f * w10.s2 + 4.f * w20.s2 + 2.f * w30.s2 + w40.s2) + 2.f * (16.f * w00.s3 + 8.f * w10.s3 + 4.f * w20.s3 + 2.f * w30.s3 + w40.s3) +
491 (16.f * w01 + 8.f * w11 + 4.f * w21 + 2.f * w31 + w41)) / 32400.f;
492 out5.s6 = (16.f * (16.f * w00.s0 + 8.f * w10.s0 + 4.f * w20.s0 + 2.f * w30.s0 + w40.s0) - 8.f * (16.f * w00.s1 + 8.f * w10.s1 + 4.f * w20.s1 + 2.f * w30.s1 + w40.s1) + 4.f *
493 (16.f * w00.s2 + 8.f * w10.s2 + 4.f * w20.s2 + 2.f * w30.s2 + w40.s2) - 2.f * (16.f * w00.s3 + 8.f * w10.s3 + 4.f * w20.s3 + 2.f * w30.s3 + w40.s3) +
494 (16.f * w01 + 8.f * w11 + 4.f * w21 + 2.f * w31 + w41)) / 32400.f;
495 out5.s7 = (16.f * w01 + 8.f * w11 + 4.f * w21 + 2.f * w31 + w41) / 180.f;
500 out6.s0 = (16.f * w00.s0 - 8.f * w10.s0 + 4.f * w20.s0 - 2.f * w30.s0 + w40.s0) / 180.f;
501 out6.s1 = -((16.f * w00.s0 - 8.f * w10.s0 + 4.f * w20.s0 - 2.f * w30.s0 + w40.s0) + (16.f * w00.s1 - 8.f * w10.s1 + 4.f * w20.s1 - 2.f * w30.s1 + w40.s1) +
502 (16.f * w00.s2 - 8.f * w10.s2 + 4.f * w20.s2 - 2.f * w30.s2 + w40.s2) + (16.f * w00.s3 - 8.f * w10.s3 + 4.f * w20.s3 - 2.f * w30.s3 + w40.s3) +
503 (16.f * w01 - 8.f * w11 + 4.f * w21 - 2.f * w31 + w41)) / 810.f;
504 out6.s2 = -((16.f * w00.s0 - 8.f * w10.s0 + 4.f * w20.s0 - 2.f * w30.s0 + w40.s0) - (16.f * w00.s1 - 8.f * w10.s1 + 4.f * w20.s1 - 2.f * w30.s1 + w40.s1) +
505 (16.f * w00.s2 - 8.f * w10.s2 + 4.f * w20.s2 - 2.f * w30.s2 + w40.s2) - (16.f * w00.s3 - 8.f * w10.s3 + 4.f * w20.s3 - 2.f * w30.s3 + w40.s3) +
506 (16.f * w01 - 8.f * w11 + 4.f * w21 - 2.f * w31 + w41)) / 810.f;
507 out6.s3 = ((16.f * w00.s0 - 8.f * w10.s0 + 4.f * w20.s0 - 2.f * w30.s0 + w40.s0) + 2.f * (16.f * w00.s1 - 8.f * w10.s1 + 4.f * w20.s1 - 2.f * w30.s1 + w40.s1) + 4.f *
508 (16.f * w00.s2 - 8.f * w10.s2 + 4.f * w20.s2 - 2.f * w30.s2 + w40.s2) + 8.f * (16.f * w00.s3 - 8.f * w10.s3 + 4.f * w20.s3 - 2.f * w30.s3 + w40.s3) + 16.f *
509 (16.f * w01 - 8.f * w11 + 4.f * w21 - 2.f * w31 + w41)) / 16200.f;
510 out6.s4 = ((16.f * w00.s0 - 8.f * w10.s0 + 4.f * w20.s0 - 2.f * w30.s0 + w40.s0) - 2.f * (16.f * w00.s1 - 8.f * w10.s1 + 4.f * w20.s1 - 2.f * w30.s1 + w40.s1) + 4.f *
511 (16.f * w00.s2 - 8.f * w10.s2 + 4.f * w20.s2 - 2.f * w30.s2 + w40.s2) - 8.f * (16.f * w00.s3 - 8.f * w10.s3 + 4.f * w20.s3 - 2.f * w30.s3 + w40.s3) + 16.f *
512 (16.f * w01 - 8.f * w11 + 4.f * w21 - 2.f * w31 + w41)) / 16200.f;
513 out6.s5 = (16.f * (16.f * w00.s0 - 8.f * w10.s0 + 4.f * w20.s0 - 2.f * w30.s0 + w40.s0) + 8.f * (16.f * w00.s1 - 8.f * w10.s1 + 4.f * w20.s1 - 2.f * w30.s1 + w40.s1) + 4.f *
514 (16.f * w00.s2 - 8.f * w10.s2 + 4.f * w20.s2 - 2.f * w30.s2 + w40.s2) + 2.f * (16.f * w00.s3 - 8.f * w10.s3 + 4.f * w20.s3 - 2.f * w30.s3 + w40.s3) +
515 (16.f * w01 - 8.f * w11 + 4.f * w21 - 2.f * w31 + w41)) / 32400.f;
516 out6.s6 = (16.f * (16.f * w00.s0 - 8.f * w10.s0 + 4.f * w20.s0 - 2.f * w30.s0 + w40.s0) - 8.f * (16.f * w00.s1 - 8.f * w10.s1 + 4.f * w20.s1 - 2.f * w30.s1 + w40.s1) + 4.f *
517 (16.f * w00.s2 - 8.f * w10.s2 + 4.f * w20.s2 - 2.f * w30.s2 + w40.s2) - 2.f * (16.f * w00.s3 - 8.f * w10.s3 + 4.f * w20.s3 - 2.f * w30.s3 + w40.s3) +
518 (16.f * w01 - 8.f * w11 + 4.f * w21 - 2.f * w31 + w41)) / 32400.f;
519 out6.s7 = (16.f * w01 - 8.f * w11 + 4.f * w21 - 2.f * w31 + w41) / 180.f;
525 out7.s1 = -2.f * (w40.s0 + w40.s1 + w40.s2 + w40.s3 + w41) / 9.f;
526 out7.s2 = -2.f * (w40.s0 - w40.s1 + w40.s2 - w40.s3 + w41) / 9.f;
527 out7.s3 = (w40.s0 + 2.f * w40.s1 + 4.f * w40.s2 + 8.f * w40.s3 + 16.f * w41) / 90.f;
528 out7.s4 = (w40.s0 - 2.f * w40.s1 + 4.f * w40.s2 - 8.f * w40.s3 + 16.f * w41) / 90.f;
529 out7.s5 = (16.f * w40.s0 + 8.f * w40.s1 + 4.f * w40.s2 + 2.f * w40.s3 + w41) / 180.f;
530 out7.s6 = (16.f * w40.s0 - 8.f * w40.s1 + 4.f * w40.s2 - 2.f * w40.s3 + w41) / 180.f;
532 #endif // !defined(WINOGRAD_FILTER_TRANSFORM_HORIZONTAL) && !defined(WINOGRAD_FILTER_TRANSFORM_VERTICAL)
534 int z = get_global_id(2);
535 int x0 = z / SRC_DIM_Z;
536 int y0 = z % SRC_DIM_Z;
539 __global uchar *dst_addr = dst_ptr + dst_offset_first_element_in_bytes + x0 *
sizeof(DATA_TYPE) + y0 * dst_stride_y;
542 *(__global DATA_TYPE *)(dst_addr + 0 * dst_stride_z) = out0.s0;
543 *(__global DATA_TYPE *)(dst_addr + 1 * dst_stride_z) = out0.s1;
544 *(__global DATA_TYPE *)(dst_addr + 2 * dst_stride_z) = out0.s2;
545 *(__global DATA_TYPE *)(dst_addr + 3 * dst_stride_z) = out0.s3;
546 *(__global DATA_TYPE *)(dst_addr + 4 * dst_stride_z) = out0.s4;
547 *(__global DATA_TYPE *)(dst_addr + 5 * dst_stride_z) = out0.s5;
548 *(__global DATA_TYPE *)(dst_addr + 6 * dst_stride_z) = out0.s6;
549 *(__global DATA_TYPE *)(dst_addr + 7 * dst_stride_z) = out0.s7;
551 #if !defined(WINOGRAD_FILTER_TRANSFORM_HORIZONTAL) && !defined(WINOGRAD_FILTER_TRANSFORM_VERTICAL)
552 *(__global DATA_TYPE *)(dst_addr + 8 * dst_stride_z) = out1.s0;
553 *(__global DATA_TYPE *)(dst_addr + 9 * dst_stride_z) = out1.s1;
554 *(__global DATA_TYPE *)(dst_addr + 10 * dst_stride_z) = out1.s2;
555 *(__global DATA_TYPE *)(dst_addr + 11 * dst_stride_z) = out1.s3;
556 *(__global DATA_TYPE *)(dst_addr + 12 * dst_stride_z) = out1.s4;
557 *(__global DATA_TYPE *)(dst_addr + 13 * dst_stride_z) = out1.s5;
558 *(__global DATA_TYPE *)(dst_addr + 14 * dst_stride_z) = out1.s6;
559 *(__global DATA_TYPE *)(dst_addr + 15 * dst_stride_z) = out1.s7;
560 *(__global DATA_TYPE *)(dst_addr + 16 * dst_stride_z) = out2.s0;
561 *(__global DATA_TYPE *)(dst_addr + 17 * dst_stride_z) = out2.s1;
562 *(__global DATA_TYPE *)(dst_addr + 18 * dst_stride_z) = out2.s2;
563 *(__global DATA_TYPE *)(dst_addr + 19 * dst_stride_z) = out2.s3;
564 *(__global DATA_TYPE *)(dst_addr + 20 * dst_stride_z) = out2.s4;
565 *(__global DATA_TYPE *)(dst_addr + 21 * dst_stride_z) = out2.s5;
566 *(__global DATA_TYPE *)(dst_addr + 22 * dst_stride_z) = out2.s6;
567 *(__global DATA_TYPE *)(dst_addr + 23 * dst_stride_z) = out2.s7;
568 *(__global DATA_TYPE *)(dst_addr + 24 * dst_stride_z) = out3.s0;
569 *(__global DATA_TYPE *)(dst_addr + 25 * dst_stride_z) = out3.s1;
570 *(__global DATA_TYPE *)(dst_addr + 26 * dst_stride_z) = out3.s2;
571 *(__global DATA_TYPE *)(dst_addr + 27 * dst_stride_z) = out3.s3;
572 *(__global DATA_TYPE *)(dst_addr + 28 * dst_stride_z) = out3.s4;
573 *(__global DATA_TYPE *)(dst_addr + 29 * dst_stride_z) = out3.s5;
574 *(__global DATA_TYPE *)(dst_addr + 30 * dst_stride_z) = out3.s6;
575 *(__global DATA_TYPE *)(dst_addr + 31 * dst_stride_z) = out3.s7;
576 *(__global DATA_TYPE *)(dst_addr + 32 * dst_stride_z) = out4.s0;
577 *(__global DATA_TYPE *)(dst_addr + 33 * dst_stride_z) = out4.s1;
578 *(__global DATA_TYPE *)(dst_addr + 34 * dst_stride_z) = out4.s2;
579 *(__global DATA_TYPE *)(dst_addr + 35 * dst_stride_z) = out4.s3;
580 *(__global DATA_TYPE *)(dst_addr + 36 * dst_stride_z) = out4.s4;
581 *(__global DATA_TYPE *)(dst_addr + 37 * dst_stride_z) = out4.s5;
582 *(__global DATA_TYPE *)(dst_addr + 38 * dst_stride_z) = out4.s6;
583 *(__global DATA_TYPE *)(dst_addr + 39 * dst_stride_z) = out4.s7;
584 *(__global DATA_TYPE *)(dst_addr + 40 * dst_stride_z) = out5.s0;
585 *(__global DATA_TYPE *)(dst_addr + 41 * dst_stride_z) = out5.s1;
586 *(__global DATA_TYPE *)(dst_addr + 42 * dst_stride_z) = out5.s2;
587 *(__global DATA_TYPE *)(dst_addr + 43 * dst_stride_z) = out5.s3;
588 *(__global DATA_TYPE *)(dst_addr + 44 * dst_stride_z) = out5.s4;
589 *(__global DATA_TYPE *)(dst_addr + 45 * dst_stride_z) = out5.s5;
590 *(__global DATA_TYPE *)(dst_addr + 46 * dst_stride_z) = out5.s6;
591 *(__global DATA_TYPE *)(dst_addr + 47 * dst_stride_z) = out5.s7;
592 *(__global DATA_TYPE *)(dst_addr + 48 * dst_stride_z) = out6.s0;
593 *(__global DATA_TYPE *)(dst_addr + 49 * dst_stride_z) = out6.s1;
594 *(__global DATA_TYPE *)(dst_addr + 50 * dst_stride_z) = out6.s2;
595 *(__global DATA_TYPE *)(dst_addr + 51 * dst_stride_z) = out6.s3;
596 *(__global DATA_TYPE *)(dst_addr + 52 * dst_stride_z) = out6.s4;
597 *(__global DATA_TYPE *)(dst_addr + 53 * dst_stride_z) = out6.s5;
598 *(__global DATA_TYPE *)(dst_addr + 54 * dst_stride_z) = out6.s6;
599 *(__global DATA_TYPE *)(dst_addr + 55 * dst_stride_z) = out6.s7;
600 *(__global DATA_TYPE *)(dst_addr + 56 * dst_stride_z) = out7.s0;
601 *(__global DATA_TYPE *)(dst_addr + 57 * dst_stride_z) = out7.s1;
602 *(__global DATA_TYPE *)(dst_addr + 58 * dst_stride_z) = out7.s2;
603 *(__global DATA_TYPE *)(dst_addr + 59 * dst_stride_z) = out7.s3;
604 *(__global DATA_TYPE *)(dst_addr + 60 * dst_stride_z) = out7.s4;
605 *(__global DATA_TYPE *)(dst_addr + 61 * dst_stride_z) = out7.s5;
606 *(__global DATA_TYPE *)(dst_addr + 62 * dst_stride_z) = out7.s6;
607 *(__global DATA_TYPE *)(dst_addr + 63 * dst_stride_z) = out7.s7;
608 #endif // !defined(WINOGRAD_FILTER_TRANSFORM_HORIZONTAL) && !defined(WINOGRAD_FILTER_TRANSFORM_VERTICAL)
611 #endif // defined(SRC_DIM_Z)
613 #if defined(WINOGRAD_FILTER_TRANSFORM_HORIZONTAL)
639 __kernel
void winograd_filter_transform_2x1_3x1_nchw(
643 winograd_filter_transform_2x2_3x3_nchw(src_ptr,
652 src_offset_first_element_in_bytes,
660 dst_offset_first_element_in_bytes);
688 __kernel
void winograd_filter_transform_4x1_3x1_nchw(
692 winograd_filter_transform_4x4_3x3_nchw(src_ptr,
701 src_offset_first_element_in_bytes,
709 dst_offset_first_element_in_bytes);
737 __kernel
void winograd_filter_transform_4x1_5x1_nchw(
741 winograd_filter_transform_4x4_5x5_nchw(src_ptr,
750 src_offset_first_element_in_bytes,
758 dst_offset_first_element_in_bytes);
761 #endif // defined(WINOGRAD_FILTER_TRANSFORM_HORIZONTAL)
763 #if defined(WINOGRAD_FILTER_TRANSFORM_VERTICAL)
789 __kernel
void winograd_filter_transform_1x2_1x3_nchw(
793 winograd_filter_transform_2x2_3x3_nchw(src_ptr,
802 src_offset_first_element_in_bytes,
810 dst_offset_first_element_in_bytes);
838 __kernel
void winograd_filter_transform_1x4_1x3_nchw(
842 winograd_filter_transform_4x4_3x3_nchw(src_ptr,
851 src_offset_first_element_in_bytes,
859 dst_offset_first_element_in_bytes);
887 __kernel
void winograd_filter_transform_1x4_1x5_nchw(
891 winograd_filter_transform_4x4_5x5_nchw(src_ptr,
900 src_offset_first_element_in_bytes,
908 dst_offset_first_element_in_bytes);
911 #endif // defined(WINOGRAD_FILTER_TRANSFORM_VERTICAL)