Gitiles
Code Review Sign In
review.mlplatform.org / ml / ComputeLibrary / 014333d73883c3872e458cedda5ccef586a7ccd4 / . / src / core / CL / cl_kernels / pooling_layer.cl
blob: 2c7ddfdf234ebde0ccee2608303bc6739600595a [file] [log] [blame]
Anthony Barbier6ff3b192017-09-04 18:44:23 +0100[diff] [blame]1/*
Isabella Gottardia527e8c2018-01-31 17:49:25 +0000[diff] [blame]2 * Copyright (c) 2017-2018 ARM Limited.
Anthony Barbier6ff3b192017-09-04 18:44:23 +0100[diff] [blame]3 *
4 * SPDX-License-Identifier: MIT
5 *
6 * Permission is hereby granted, free of charge, to any person obtaining a copy
7 * of this software and associated documentation files (the "Software"), to
8 * deal in the Software without restriction, including without limitation the
9 * rights to use, copy, modify, merge, publish, distribute, sublicense, and/or
10 * sell copies of the Software, and to permit persons to whom the Software is
11 * furnished to do so, subject to the following conditions:
12 *
13 * The above copyright notice and this permission notice shall be included in all
14 * copies or substantial portions of the Software.
15 *
16 * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
17 * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
18 * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
19 * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
20 * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
21 * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE
22 * SOFTWARE.
23 */
24#include "helpers.h"
25
steniu010c7614f2017-06-23 17:00:26 +0100[diff] [blame]26#ifdef FIXED_POINT_POSITION
27
28#include "fixed_point.h"
29
30#if defined(POOL_AVG)
31#define POOL_OP(x, y) add_sat(x, y)
32#else /* POOL_AVG */
33#define POOL_OP(x, y) (max((x), (y)))
34#endif /* POOL_AVG */
35
Georgios Pinitascdf51452017-08-31 14:21:36 +0100[diff] [blame]36#define DIV_OP1(x, y) DIV_SAT_OP_EXPAND((x), (y), DATA_TYPE, FIXED_POINT_POSITION)
steniu010c7614f2017-06-23 17:00:26 +0100[diff] [blame]37#define DIV_OP(x, y) DIV_OP1(x, y << FIXED_POINT_POSITION)
Georgios Pinitascdf51452017-08-31 14:21:36 +0100[diff] [blame]38#define SQRT_OP(x) DIV_OP1((1 << FIXED_POINT_POSITION), (INVSQRT_OP_EXPAND((x), DATA_TYPE, 1, FIXED_POINT_POSITION)))
39
40#if defined(POOL_L2)
41#define POW2_OP(x, vec_size) MUL_SAT_OP_EXPAND((x), (x), DATA_TYPE, vec_size, FIXED_POINT_POSITION)
42#else /* defined(POOL_L2) */
43#define POW2_OP(x, vec_size) (x)
44#endif /* defined(POOL_L2) */
steniu010c7614f2017-06-23 17:00:26 +0100[diff] [blame]45
46#else /* FIXED_POINT_POSITION */
47
Georgios Pinitascdf51452017-08-31 14:21:36 +0100[diff] [blame]48#if defined(POOL_AVG) || defined(POOL_L2)
Anthony Barbier6ff3b192017-09-04 18:44:23 +0100[diff] [blame]49#define POOL_OP(x, y) ((x) + (y))
Georgios Pinitascdf51452017-08-31 14:21:36 +0100[diff] [blame]50#else /* defined(POOL_AVG) || defined(POOL_L2) */
Anthony Barbier6ff3b192017-09-04 18:44:23 +0100[diff] [blame]51#define POOL_OP(x, y) (fmax((x), (y)))
Georgios Pinitascdf51452017-08-31 14:21:36 +0100[diff] [blame]52#endif /* defined(POOL_AVG) || defined(POOL_L2) */
53
54#if defined(POOL_L2)
55#define POW2_OP(x, vec_size) ((x) * (x))
56#else /* defined(POOL_L2) */
57#define POW2_OP(x, vec_size) (x)
58#endif /* defined(POOL_L2) */
Anthony Barbier6ff3b192017-09-04 18:44:23 +0100[diff] [blame]59
steniu010c7614f2017-06-23 17:00:26 +0100[diff] [blame]60#define DIV_OP(x, y) (x * (1.f / y))
Georgios Pinitascdf51452017-08-31 14:21:36 +0100[diff] [blame]61#define SQRT_OP(x) sqrt((x))
steniu010c7614f2017-06-23 17:00:26 +0100[diff] [blame]62
63#endif /* FIXED_POINT_POSITION */
64
Michalis Spyroue74b2012018-04-18 09:49:16 +0100[diff] [blame]65#define DIV_OP_NHWC(x, y) (x * (VEC_DATA_TYPE(DATA_TYPE, 8))(1.f / y))
66
Gian Marco Iodicecb292832017-08-02 13:19:48 +0100[diff] [blame]67#if STRIDE_X == 1
68#define POOLING3x3(res, input, output) POOLING3x3_STRIDE1(res, input, output)
69#elif STRIDE_X == 2 /* STRIDE_X == 1 */
70#define POOLING3x3(res, input, output) POOLING3x3_STRIDE2(res, input, output)
71#elif STRIDE_X == 3 /* STRIDE_X not equals 1 or 2 */
72#define POOLING3x3(res, input, output) POOLING3x3_STRIDE3(res, input, output)
73#endif /* STRIDE_X == 3 */
74
Gian Marco Iodicecb292832017-08-02 13:19:48 +0100[diff] [blame]75#define POOLING3x3_STRIDE1(res, input, output) \
76 ({ \
77 VEC_DATA_TYPE(DATA_TYPE, 4) \
78 data00 = vload4(0, (__global DATA_TYPE *)tensor3D_offset(&input, 0, 0, 0)); \
79 VEC_DATA_TYPE(DATA_TYPE, 2) \
80 data01 = vload2(0, (__global DATA_TYPE *)tensor3D_offset(&input, 0, 0, 0) + 4); \
81 VEC_DATA_TYPE(DATA_TYPE, 4) \
82 data10 = vload4(0, (__global DATA_TYPE *)tensor3D_offset(&input, 0, 1, 0)); \
83 VEC_DATA_TYPE(DATA_TYPE, 2) \
84 data11 = vload2(0, (__global DATA_TYPE *)tensor3D_offset(&input, 0, 1, 0) + 4); \
85 VEC_DATA_TYPE(DATA_TYPE, 4) \
86 data20 = vload4(0, (__global DATA_TYPE *)tensor3D_offset(&input, 0, 2, 0)); \
87 VEC_DATA_TYPE(DATA_TYPE, 2) \
88 data21 = vload2(0, (__global DATA_TYPE *)tensor3D_offset(&input, 0, 2, 0) + 4); \
Georgios Pinitascdf51452017-08-31 14:21:36 +0100[diff] [blame]89 data00 = POW2_OP(data00, 4); \
90 data01 = POW2_OP(data01, 2); \
91 data10 = POW2_OP(data10, 4); \
92 data11 = POW2_OP(data11, 2); \
93 data20 = POW2_OP(data20, 4); \
94 data21 = POW2_OP(data21, 2); \
Gian Marco Iodicecb292832017-08-02 13:19:48 +0100[diff] [blame]95 \
96 VEC_DATA_TYPE(DATA_TYPE, 8) \
97 values00 = (VEC_DATA_TYPE(DATA_TYPE, 8))(data00.s01212323); \
98 VEC_DATA_TYPE(DATA_TYPE, 4) \
99 values01 = (VEC_DATA_TYPE(DATA_TYPE, 4))(data01.s0, data00.s3, data01.s01); \
100 VEC_DATA_TYPE(DATA_TYPE, 8) \
101 values10 = (VEC_DATA_TYPE(DATA_TYPE, 8))(data10.s01212323); \
102 VEC_DATA_TYPE(DATA_TYPE, 4) \
103 values11 = (VEC_DATA_TYPE(DATA_TYPE, 4))(data11.s0, data10.s3, data11.s01); \
104 VEC_DATA_TYPE(DATA_TYPE, 8) \
105 values20 = (VEC_DATA_TYPE(DATA_TYPE, 8))(data20.s01212323); \
106 VEC_DATA_TYPE(DATA_TYPE, 4) \
107 values21 = (VEC_DATA_TYPE(DATA_TYPE, 4))(data21.s0, data20.s3, data21.s01); \
108 \
109 values00 = POOL_OP(values00, values10); \
110 values01 = POOL_OP(values01, values11); \
111 values00 = POOL_OP(values00, values20); \
112 values01 = POOL_OP(values01, values21); \
113 \
114 res = POOL_OP((VEC_DATA_TYPE(DATA_TYPE, 4))(values00.s036, values01.s1), (VEC_DATA_TYPE(DATA_TYPE, 4))(values00.s147, values01.s2)); \
115 res = POOL_OP(res, (VEC_DATA_TYPE(DATA_TYPE, 4))(values00.s25, values01.s03)); \
116 })
117
118#define POOLING3x3_STRIDE2(res, input, output) \
119 ({ \
120 VEC_DATA_TYPE(DATA_TYPE, 8) \
121 data00 = vload8(0, (__global DATA_TYPE *)tensor3D_offset(&input, 0, 0, 0)); \
122 DATA_TYPE data01 = *((__global DATA_TYPE *)tensor3D_offset(&input, 0, 0, 0) + 8); \
123 VEC_DATA_TYPE(DATA_TYPE, 8) \
124 data10 = vload8(0, (__global DATA_TYPE *)tensor3D_offset(&input, 0, 1, 0)); \
125 DATA_TYPE data11 = *((__global DATA_TYPE *)tensor3D_offset(&input, 0, 1, 0) + 8); \
126 VEC_DATA_TYPE(DATA_TYPE, 8) \
127 data20 = vload8(0, (__global DATA_TYPE *)tensor3D_offset(&input, 0, 2, 0)); \
128 DATA_TYPE data21 = *((__global DATA_TYPE *)tensor3D_offset(&input, 0, 2, 0) + 8); \
Georgios Pinitascdf51452017-08-31 14:21:36 +0100[diff] [blame]129 data00 = POW2_OP(data00, 8); \
130 data01 = POW2_OP(data01, 1); \
131 data10 = POW2_OP(data10, 8); \
132 data11 = POW2_OP(data11, 1); \
133 data20 = POW2_OP(data20, 8); \
134 data21 = POW2_OP(data21, 1); \
Gian Marco Iodicecb292832017-08-02 13:19:48 +0100[diff] [blame]135 \
136 VEC_DATA_TYPE(DATA_TYPE, 8) \
137 values00 = (VEC_DATA_TYPE(DATA_TYPE, 8))(data00.s01223445); \
138 VEC_DATA_TYPE(DATA_TYPE, 4) \
139 values01 = (VEC_DATA_TYPE(DATA_TYPE, 4))(data00.s667, data01); \
140 VEC_DATA_TYPE(DATA_TYPE, 8) \
141 values10 = (VEC_DATA_TYPE(DATA_TYPE, 8))(data10.s01223445); \
142 VEC_DATA_TYPE(DATA_TYPE, 4) \
143 values11 = (VEC_DATA_TYPE(DATA_TYPE, 4))(data10.s667, data11); \
144 VEC_DATA_TYPE(DATA_TYPE, 8) \
145 values20 = (VEC_DATA_TYPE(DATA_TYPE, 8))(data20.s01223445); \
146 VEC_DATA_TYPE(DATA_TYPE, 4) \
147 values21 = (VEC_DATA_TYPE(DATA_TYPE, 4))(data20.s667, data21); \
148 \
149 values00 = POOL_OP(values00, values10); \
150 values01 = POOL_OP(values01, values11); \
151 values00 = POOL_OP(values00, values20); \
152 values01 = POOL_OP(values01, values21); \
153 \
154 res = POOL_OP((VEC_DATA_TYPE(DATA_TYPE, 4))(values00.s036, values01.s1), (VEC_DATA_TYPE(DATA_TYPE, 4))(values00.s147, values01.s2)); \
155 res = POOL_OP(res, (VEC_DATA_TYPE(DATA_TYPE, 4))(values00.s25, values01.s03)); \
156 })
157
158#define POOLING3x3_STRIDE3(res, input, output) \
159 ({ \
160 VEC_DATA_TYPE(DATA_TYPE, 8) \
161 data00 = vload8(0, (__global DATA_TYPE *)tensor3D_offset(&input, 0, 0, 0)); \
162 VEC_DATA_TYPE(DATA_TYPE, 4) \
163 data01 = vload4(0, (__global DATA_TYPE *)tensor3D_offset(&input, 0, 0, 0) + 8); \
164 VEC_DATA_TYPE(DATA_TYPE, 8) \
165 data10 = vload8(0, (__global DATA_TYPE *)tensor3D_offset(&input, 0, 1, 0)); \
166 VEC_DATA_TYPE(DATA_TYPE, 4) \
167 data11 = vload4(0, (__global DATA_TYPE *)tensor3D_offset(&input, 0, 1, 0) + 8); \
168 VEC_DATA_TYPE(DATA_TYPE, 8) \
169 data20 = vload8(0, (__global DATA_TYPE *)tensor3D_offset(&input, 0, 2, 0)); \
170 VEC_DATA_TYPE(DATA_TYPE, 4) \
171 data21 = vload4(0, (__global DATA_TYPE *)tensor3D_offset(&input, 0, 2, 0) + 8); \
Georgios Pinitascdf51452017-08-31 14:21:36 +0100[diff] [blame]172 data00 = POW2_OP(data00, 8); \
173 data01 = POW2_OP(data01, 4); \
174 data10 = POW2_OP(data10, 8); \
175 data11 = POW2_OP(data11, 4); \
176 data20 = POW2_OP(data20, 8); \
177 data21 = POW2_OP(data21, 4); \
Gian Marco Iodicecb292832017-08-02 13:19:48 +0100[diff] [blame]178 \
179 data00 = POOL_OP(data00, data10); \
180 data01 = POOL_OP(data01, data11); \
181 data00 = POOL_OP(data00, data20); \
182 data01 = POOL_OP(data01, data21); \
183 \
184 res = POOL_OP((VEC_DATA_TYPE(DATA_TYPE, 4))(data00.s036, data01.s1), (VEC_DATA_TYPE(DATA_TYPE, 4))(data00.s147, data01.s2)); \
185 res = POOL_OP(res, (VEC_DATA_TYPE(DATA_TYPE, 4))(data00.s25, data01.s03)); \
186 })
187
Isabella Gottardia527e8c2018-01-31 17:49:25 +0000[diff] [blame]188DATA_TYPE calculate_avg_scale(const int pool_size_x, const int pool_size_y, const int upper_bound_w, const int upper_bound_h,
Gian Marco Iodicecb292832017-08-02 13:19:48 +0100[diff] [blame]189 const int pad_x, const int pad_y, const int stride_x, const int stride_y)
Anthony Barbier6ff3b192017-09-04 18:44:23 +0100[diff] [blame]190{
Georgios Pinitasadaae7e2017-10-30 15:56:32 +0000[diff] [blame]191 int start_x = get_global_id(0) * stride_x - pad_x;
192 int start_y = get_global_id(1) * stride_y - pad_y;
Isabella Gottardia527e8c2018-01-31 17:49:25 +0000[diff] [blame]193 const int end_x = min(start_x + pool_size_x, upper_bound_w);
194 const int end_y = min(start_y + pool_size_y, upper_bound_h);
Georgios Pinitasadaae7e2017-10-30 15:56:32 +0000[diff] [blame]195#if defined(EXCLUDE_PADDING)
196 start_x = max(0, start_x);
197 start_y = max(0, start_y);
198#endif /* defined(EXCLUDE_PADDING) */
steniu010c7614f2017-06-23 17:00:26 +0100[diff] [blame]199 return ((end_y - start_y) * (end_x - start_x));
Gian Marco Iodicecb292832017-08-02 13:19:48 +0100[diff] [blame]200}
201
Anthony Barbier6ff3b192017-09-04 18:44:23 +0100[diff] [blame]202/** Performs a pooling function of pool size equal to 2.
203 *
steniu010c7614f2017-06-23 17:00:26 +0100[diff] [blame]204 * @note Datatype must be passed using -DDATA_TYPE e.g. -DDATA_TYPE=float. Supported data types are QS8/QS16/F16/F32;
Gian Marco Iodicecb292832017-08-02 13:19:48 +0100[diff] [blame]205 * @note In case of average pooling the following information must be passed at compile time:
Georgios Pinitascdf51452017-08-31 14:21:36 +0100[diff] [blame]206 * -DPOOL_AVG or -DPOOL_L2 must be provided otherwise max pooling will be performed.
Gian Marco Iodicecb292832017-08-02 13:19:48 +0100[diff] [blame]207 * -DMAX_WIDTH and -DMAX_HEIGHT which are the maximum accessible indeces in x and y dimensions (width + pad)
208 * -DSTRIDE_X and -DSTRIDE_Y which are the steps of the window along the x and y directions
209 * -DPAD_X and -DPAD_Y which are the pooling paddings in x and y dimension
Anthony Barbier6ff3b192017-09-04 18:44:23 +0100[diff] [blame]210 *
steniu010c7614f2017-06-23 17:00:26 +0100[diff] [blame]211 * @param[in] input_ptr Pointer to the source image. Supported data types: QS8/QS16/F16/F32
Anthony Barbier6ff3b192017-09-04 18:44:23 +0100[diff] [blame]212 * @param[in] input_stride_x Stride of the source image in X dimension (in bytes)
213 * @param[in] input_step_x input_stride_x * number of elements along X processed per workitem(in bytes)
214 * @param[in] input_stride_y Stride of the source image in Y dimension (in bytes)
215 * @param[in] input_step_y input_stride_y * number of elements along Y processed per workitem(in bytes)
216 * @param[in] input_stride_z Stride of the source tensor in Z dimension (in bytes)
217 * @param[in] input_step_z input_stride_z * number of elements along Z processed per workitem(in bytes)
218 * @param[in] input_offset_first_element_in_bytes The offset of the first element in the source image
Gian Marco Iodicecb292832017-08-02 13:19:48 +0100[diff] [blame]219 * @param[out] output_ptr Pointer to the destination image. Supported data types: same as @p input_ptr
Anthony Barbier6ff3b192017-09-04 18:44:23 +0100[diff] [blame]220 * @param[in] output_stride_x Stride of the destination image in X dimension (in bytes)
221 * @param[in] output_step_x output_stride_x * number of elements along X processed per workitem(in bytes)
222 * @param[in] output_stride_y Stride of the destination image in Y dimension (in bytes)
223 * @param[in] output_step_y output_stride_y * number of elements along Y processed per workitem(in bytes)
224 * @param[in] output_stride_z Stride of the source tensor in Z dimension (in bytes)
225 * @param[in] output_step_z output_stride_z * number of elements along Z processed per workitem(in bytes)
226 * @param[in] output_offset_first_element_in_bytes The offset of the first element in the destination image
Anthony Barbier6ff3b192017-09-04 18:44:23 +0100[diff] [blame]227 */
228__kernel void pooling_layer_2(
229 TENSOR3D_DECLARATION(input),
Gian Marco Iodicecb292832017-08-02 13:19:48 +0100[diff] [blame]230 TENSOR3D_DECLARATION(output))
Anthony Barbier6ff3b192017-09-04 18:44:23 +0100[diff] [blame]231{
232 // Get pixels pointer
233 Tensor3D input = CONVERT_TO_TENSOR3D_STRUCT(input);
234 Tensor3D output = CONVERT_TO_TENSOR3D_STRUCT(output);
235
236 // Load data
237 VEC_DATA_TYPE(DATA_TYPE, 2)
238 data0 = vload2(0, (__global DATA_TYPE *)tensor3D_offset(&input, 0, 0, 0));
239 VEC_DATA_TYPE(DATA_TYPE, 2)
240 data1 = vload2(0, (__global DATA_TYPE *)tensor3D_offset(&input, 0, 1, 0));
241
Georgios Pinitascdf51452017-08-31 14:21:36 +0100[diff] [blame]242#if defined(POOL_L2)
243 // Raise to power of 2 for L2 Pooling
244 data0 = POW2_OP(data0, 2);
245 data1 = POW2_OP(data1, 2);
246#endif /* defined(POOL_L2) */
247
Anthony Barbier6ff3b192017-09-04 18:44:23 +0100[diff] [blame]248 // Perform calculations
249 data0 = POOL_OP(data0, data1);
250 DATA_TYPE res = POOL_OP(data0.s0, data0.s1);
251
Georgios Pinitascdf51452017-08-31 14:21:36 +0100[diff] [blame]252#if defined(POOL_AVG) || defined(POOL_L2)
253 // Divide by pool region in case of average or l2 pooling
Isabella Gottardia527e8c2018-01-31 17:49:25 +0000[diff] [blame]254 res = DIV_OP(res, calculate_avg_scale(2, 2, MAX_WIDTH, MAX_HEIGHT, PAD_X, PAD_Y, STRIDE_X, STRIDE_Y));
Georgios Pinitascdf51452017-08-31 14:21:36 +0100[diff] [blame]255#endif /* defined(POOL_AVG) || defined(POOL_L2) */
256
257#if defined(POOL_L2)
258 // Take square root of the result in L2 pooling
259 res = SQRT_OP(res);
260#endif /* defined(POOL_L2) */
Anthony Barbier6ff3b192017-09-04 18:44:23 +0100[diff] [blame]261
262 // Store result
263 *(__global DATA_TYPE *)output.ptr = res;
264}
265
Gian Marco Iodicecb292832017-08-02 13:19:48 +0100[diff] [blame]266/** Performs a pooling function of pool size equal to 3
Anthony Barbier6ff3b192017-09-04 18:44:23 +0100[diff] [blame]267 *
steniu010c7614f2017-06-23 17:00:26 +0100[diff] [blame]268 * @note Datatype must be passed using -DDATA_TYPE e.g. -DDATA_TYPE=float. Supported data types are QS8/QS16/F16/F32;
Gian Marco Iodicecb292832017-08-02 13:19:48 +0100[diff] [blame]269 * @note In case of average pooling the following information must be passed at compile time:
Georgios Pinitascdf51452017-08-31 14:21:36 +0100[diff] [blame]270 * -DPOOL_AVG or -DPOOL_L2 must be provided otherwise max pooling will be performed.
Gian Marco Iodicecb292832017-08-02 13:19:48 +0100[diff] [blame]271 * -DMAX_WIDTH and -DMAX_HEIGHT which are the maximum accessible indeces in x and y dimensions (width + pad)
272 * -DSTRIDE_X and -DSTRIDE_Y which are the steps of the window along the x and y directions
273 * -DPAD_X and -DPAD_Y which are the pooling paddings in x and y dimension
Anthony Barbier6ff3b192017-09-04 18:44:23 +0100[diff] [blame]274 *
steniu010c7614f2017-06-23 17:00:26 +0100[diff] [blame]275 * @param[in] input_ptr Pointer to the source image. Supported data types: QS8/QS16/F16/F32
Anthony Barbier6ff3b192017-09-04 18:44:23 +0100[diff] [blame]276 * @param[in] input_stride_x Stride of the source image in X dimension (in bytes)
277 * @param[in] input_step_x input_stride_x * number of elements along X processed per workitem(in bytes)
278 * @param[in] input_stride_y Stride of the source image in Y dimension (in bytes)
279 * @param[in] input_step_y input_stride_y * number of elements along Y processed per workitem(in bytes)
280 * @param[in] input_stride_z Stride of the source tensor in Z dimension (in bytes)
281 * @param[in] input_step_z input_stride_z * number of elements along Z processed per workitem(in bytes)
282 * @param[in] input_offset_first_element_in_bytes The offset of the first element in the source image
Gian Marco Iodicecb292832017-08-02 13:19:48 +0100[diff] [blame]283 * @param[out] output_ptr Pointer to the destination image. Supported data types: same as @p input_ptr
Anthony Barbier6ff3b192017-09-04 18:44:23 +0100[diff] [blame]284 * @param[in] output_stride_x Stride of the destination image in X dimension (in bytes)
285 * @param[in] output_step_x output_stride_x * number of elements along X processed per workitem(in bytes)
286 * @param[in] output_stride_y Stride of the destination image in Y dimension (in bytes)
287 * @param[in] output_step_y output_stride_y * number of elements along Y processed per workitem(in bytes)
288 * @param[in] output_stride_z Stride of the source tensor in Z dimension (in bytes)
289 * @param[in] output_step_z output_stride_z * number of elements along Z processed per workitem(in bytes)
290 * @param[in] output_offset_first_element_in_bytes The offset of the first element in the destination image
Anthony Barbier6ff3b192017-09-04 18:44:23 +0100[diff] [blame]291 */
292__kernel void pooling_layer_3(
293 TENSOR3D_DECLARATION(input),
Gian Marco Iodicecb292832017-08-02 13:19:48 +0100[diff] [blame]294 TENSOR3D_DECLARATION(output))
Anthony Barbier6ff3b192017-09-04 18:44:23 +0100[diff] [blame]295{
296 // Get pixels pointer
297 Tensor3D input = CONVERT_TO_TENSOR3D_STRUCT(input);
298 Tensor3D output = CONVERT_TO_TENSOR3D_STRUCT(output);
299
300 // Load data
301 VEC_DATA_TYPE(DATA_TYPE, 3)
302 data0 = vload3(0, (__global DATA_TYPE *)tensor3D_offset(&input, 0, 0, 0));
303 VEC_DATA_TYPE(DATA_TYPE, 3)
304 data1 = vload3(0, (__global DATA_TYPE *)tensor3D_offset(&input, 0, 1, 0));
305 VEC_DATA_TYPE(DATA_TYPE, 3)
306 data2 = vload3(0, (__global DATA_TYPE *)tensor3D_offset(&input, 0, 2, 0));
307
Georgios Pinitascdf51452017-08-31 14:21:36 +0100[diff] [blame]308#if defined(POOL_L2)
309 // Raise to power of 2 for L2 Pooling
310 data0 = POW2_OP(data0, 3);
311 data1 = POW2_OP(data1, 3);
312 data2 = POW2_OP(data2, 3);
313#endif /* defined(POOL_L2) */
314
Anthony Barbier6ff3b192017-09-04 18:44:23 +0100[diff] [blame]315 // Perform calculations
316 data0 = POOL_OP(data0, data1);
317 data0 = POOL_OP(data0, data2);
318 DATA_TYPE res = POOL_OP(POOL_OP(data0.s0, data0.s1), data0.s2);
319
Georgios Pinitascdf51452017-08-31 14:21:36 +0100[diff] [blame]320#if defined(POOL_AVG) || defined(POOL_L2)
Georgios Pinitasce093142017-06-19 16:11:53 +0100[diff] [blame]321 // Divide by pool region in case of average pooling
Isabella Gottardia527e8c2018-01-31 17:49:25 +0000[diff] [blame]322 res = DIV_OP(res, calculate_avg_scale(3, 3, MAX_WIDTH, MAX_HEIGHT, PAD_X, PAD_Y, STRIDE_X, STRIDE_Y));
Georgios Pinitascdf51452017-08-31 14:21:36 +0100[diff] [blame]323#endif /* defined(POOL_AVG) || defined(POOL_L2) */
324
325#if defined(POOL_L2)
326 // Take square root of the result in L2 pooling
327 res = SQRT_OP(res);
328#endif /* defined(POOL_L2) */
Anthony Barbier6ff3b192017-09-04 18:44:23 +0100[diff] [blame]329
330 // Store result
331 *(__global DATA_TYPE *)output.ptr = res;
332}
Georgios Pinitasce093142017-06-19 16:11:53 +0100[diff] [blame]333
steniu010c7614f2017-06-23 17:00:26 +0100[diff] [blame]334#if defined(POOLING3x3) && !defined(FIXED_POINT_POSITION)
335
336#define CONVERT_OP(data_type) convert_##data_type##4
337#define CONVERT_VECTOR4(data_type) CONVERT_OP(data_type)
338
339VEC_DATA_TYPE(DATA_TYPE, 4)
340calculate_avg_scale4(const int pool_size, const int upper_bound_w, const int upper_bound_h,
341 const int pad_x, const int pad_y, const int stride_x, const int stride_y)
342{
Georgios Pinitasadaae7e2017-10-30 15:56:32 +0000[diff] [blame]343 int4 start_x = ((int4)get_global_id(0) * 4 + (int4)(0, 1, 2, 3)) * (int4)stride_x - (int4)pad_x;
344 int start_y = get_global_id(1) * stride_y - pad_y;
steniu010c7614f2017-06-23 17:00:26 +0100[diff] [blame]345 const int4 end_x = min(start_x + (int4)pool_size, (int4)upper_bound_w);
346 const int end_y = min(start_y + pool_size, upper_bound_h);
Georgios Pinitasadaae7e2017-10-30 15:56:32 +0000[diff] [blame]347#if defined(EXCLUDE_PADDING)
348 start_x = max((int4)0, start_x);
349 start_y = max(0, start_y);
350#endif /* defined(EXCLUDE_PADDING) */
steniu010c7614f2017-06-23 17:00:26 +0100[diff] [blame]351 return (VEC_DATA_TYPE(DATA_TYPE, 4))(1.f) / CONVERT_VECTOR4(DATA_TYPE)(((int4)(end_y - start_y)) * (end_x - start_x));
352}
353
Gian Marco Iodicecb292832017-08-02 13:19:48 +0100[diff] [blame]354/** Performs an optimized pooling function of pool size equal to 3 when the stride_x is less equal than 3
Georgios Pinitasce093142017-06-19 16:11:53 +0100[diff] [blame]355 *
steniu010c7614f2017-06-23 17:00:26 +0100[diff] [blame]356 * @note Datatype must be passed using -DDATA_TYPE e.g. -DDATA_TYPE=float. Supported data types are QS8/QS16/F16/F32;
Gian Marco Iodicecb292832017-08-02 13:19:48 +0100[diff] [blame]357 * @note In case of average pooling the following information must be passed at compile time:
Georgios Pinitascdf51452017-08-31 14:21:36 +0100[diff] [blame]358 * -DPOOL_AVG or -DPOOL_L2 must be provided otherwise max pooling will be performed.
Gian Marco Iodicecb292832017-08-02 13:19:48 +0100[diff] [blame]359 * -DMAX_WIDTH and -DMAX_HEIGHT which are the maximum accessible indeces in x and y dimensions (width + pad)
360 * -DSTRIDE_X and -DSTRIDE_Y which are the steps of the window along the x and y directions
361 * -DPAD_X and -DPAD_Y which are the pooling paddings in x and y dimension
Georgios Pinitasce093142017-06-19 16:11:53 +0100[diff] [blame]362 *
Gian Marco Iodicecb292832017-08-02 13:19:48 +0100[diff] [blame]363 * @param[in] input_ptr Pointer to the source image. Supported data types: F16/F32
Georgios Pinitasce093142017-06-19 16:11:53 +0100[diff] [blame]364 * @param[in] input_stride_x Stride of the source image in X dimension (in bytes)
365 * @param[in] input_step_x input_stride_x * number of elements along X processed per workitem(in bytes)
366 * @param[in] input_stride_y Stride of the source image in Y dimension (in bytes)
367 * @param[in] input_step_y input_stride_y * number of elements along Y processed per workitem(in bytes)
368 * @param[in] input_stride_z Stride of the source tensor in Z dimension (in bytes)
369 * @param[in] input_step_z input_stride_z * number of elements along Z processed per workitem(in bytes)
370 * @param[in] input_offset_first_element_in_bytes The offset of the first element in the source image
Gian Marco Iodicecb292832017-08-02 13:19:48 +0100[diff] [blame]371 * @param[out] output_ptr Pointer to the destination image. Supported data types: same as @p input_ptr
Georgios Pinitasce093142017-06-19 16:11:53 +0100[diff] [blame]372 * @param[in] output_stride_x Stride of the destination image in X dimension (in bytes)
373 * @param[in] output_step_x output_stride_x * number of elements along X processed per workitem(in bytes)
374 * @param[in] output_stride_y Stride of the destination image in Y dimension (in bytes)
375 * @param[in] output_step_y output_stride_y * number of elements along Y processed per workitem(in bytes)
376 * @param[in] output_stride_z Stride of the source tensor in Z dimension (in bytes)
377 * @param[in] output_step_z output_stride_z * number of elements along Z processed per workitem(in bytes)
378 * @param[in] output_offset_first_element_in_bytes The offset of the first element in the destination image
Gian Marco Iodicecb292832017-08-02 13:19:48 +0100[diff] [blame]379 */
Anton Lokhmotovaf6204c2017-11-08 09:34:19 +0000[diff] [blame]380__kernel void pooling_layer_optimized_3(
Gian Marco Iodicecb292832017-08-02 13:19:48 +0100[diff] [blame]381 TENSOR3D_DECLARATION(input),
382 TENSOR3D_DECLARATION(output))
383{
384 // Get pixels pointer
385 Tensor3D input = CONVERT_TO_TENSOR3D_STRUCT(input);
386 Tensor3D output = CONVERT_TO_TENSOR3D_STRUCT(output);
387
388 VEC_DATA_TYPE(DATA_TYPE, 4)
389 res;
390
391 // Perform pooling 3x3 for 4 output elements
392 POOLING3x3(res, input, output);
393
Georgios Pinitascdf51452017-08-31 14:21:36 +0100[diff] [blame]394#if defined(POOL_AVG) || defined(POOL_L2)
Gian Marco Iodicecb292832017-08-02 13:19:48 +0100[diff] [blame]395 // Divide by pool region in case of average pooling
Gian Marco Iodicecb292832017-08-02 13:19:48 +0100[diff] [blame]396 res *= calculate_avg_scale4(3, MAX_WIDTH, MAX_HEIGHT, PAD_X, PAD_Y, STRIDE_X, STRIDE_Y);
Georgios Pinitascdf51452017-08-31 14:21:36 +0100[diff] [blame]397#endif /* defined(POOL_AVG) || defined(POOL_L2) */
398
399#if defined(POOL_L2)
400 // Take square root of the result in L2 pooling
401 res = SQRT_OP(res);
402#endif /* defined(POOL_L2) */
Gian Marco Iodicecb292832017-08-02 13:19:48 +0100[diff] [blame]403
404 vstore4(res, 0, (__global DATA_TYPE *)output.ptr);
405}
steniu010c7614f2017-06-23 17:00:26 +0100[diff] [blame]406#endif // defined(POOLING3x3) && !defined(FIXED_POINT_POSITION)
Gian Marco Iodicecb292832017-08-02 13:19:48 +0100[diff] [blame]407
Isabella Gottardia527e8c2018-01-31 17:49:25 +0000[diff] [blame]408#if defined(POOL_SIZE_X) && defined(POOL_SIZE_Y)
Gian Marco Iodicebf179552017-09-05 13:51:21 +0100[diff] [blame]409
410// Set the initial value for the pooling operation accordingly with the data type
Georgios Pinitascdf51452017-08-31 14:21:36 +0100[diff] [blame]411#if defined(POOL_AVG) || defined(POOL_L2)
Gian Marco Iodicebf179552017-09-05 13:51:21 +0100[diff] [blame]412#define INITIAL_VALUE 0
Georgios Pinitascdf51452017-08-31 14:21:36 +0100[diff] [blame]413#else /* defined(POOL_AVG) || defined(POOL_L2) */
Gian Marco Iodicebf179552017-09-05 13:51:21 +0100[diff] [blame]414#ifdef FIXED_POINT_POSITION
415#define MIN_VAL_EXPAND(type) type##_MIN
416#define MIN_VAL(type) MIN_VAL_EXPAND(type)
417#define INITIAL_VALUE MIN_VAL(DATA_TYPE)
Gian Marco Iodicebf179552017-09-05 13:51:21 +0100[diff] [blame]418#else // FIXED_POINT_POSITION
419#if FP16
420#define INITIAL_VALUE -HALF_MAX
421#else // FP16
422#define INITIAL_VALUE -FLT_MAX
423#endif // FP16
424#endif // FIXED_POINT_POSITION
425
426#endif // POOL_AVG
427
Michalis Spyroue74b2012018-04-18 09:49:16 +0100[diff] [blame]428/** Performs a pooling function of pool size equal to N (NCHW)
Gian Marco Iodicebf179552017-09-05 13:51:21 +0100[diff] [blame]429 *
Georgios Pinitas13fc22c2017-10-19 18:35:59 +0100[diff] [blame]430 * @note Datatype must be passed using -DDATA_TYPE e.g. -DDATA_TYPE=float. Supported data types are QS8/QS16/F16/F32;
Gian Marco Iodicebf179552017-09-05 13:51:21 +0100[diff] [blame]431 * @note -DFP16 must be passed at compile time if half float data type is used
Isabella Gottardia527e8c2018-01-31 17:49:25 +0000[diff] [blame]432 * @note Pool sizes must be passed using -DPOOL_SIZE_X and -DPOOL_SIZE_Y e.g. -DPOOL_SIZE_X=13;
Gian Marco Iodicebf179552017-09-05 13:51:21 +0100[diff] [blame]433 * @note In case of average pooling the following information must be passed at compile time:
434 * -DPOOL_AVG must be provided otherwise max pooling will be performed.
435 * -DMAX_WIDTH and -DMAX_HEIGHT which are the maximum accessible indeces in x and y dimensions (width + pad)
436 * -DSTRIDE_X and -DSTRIDE_Y which are the steps of the window along the x and y directions
437 * -DPAD_X and -DPAD_Y which are the pooling paddings in x and y dimension
438 *
Georgios Pinitas13fc22c2017-10-19 18:35:59 +0100[diff] [blame]439 * @param[in] input_ptr Pointer to the source image. Supported data types: QS8/QS16/F16/F32
Gian Marco Iodicebf179552017-09-05 13:51:21 +0100[diff] [blame]440 * @param[in] input_stride_x Stride of the source image in X dimension (in bytes)
441 * @param[in] input_step_x input_stride_x * number of elements along X processed per workitem(in bytes)
442 * @param[in] input_stride_y Stride of the source image in Y dimension (in bytes)
443 * @param[in] input_step_y input_stride_y * number of elements along Y processed per workitem(in bytes)
444 * @param[in] input_stride_z Stride of the source tensor in Z dimension (in bytes)
445 * @param[in] input_step_z input_stride_z * number of elements along Z processed per workitem(in bytes)
446 * @param[in] input_offset_first_element_in_bytes The offset of the first element in the source image
447 * @param[out] output_ptr Pointer to the destination image. Supported data types: same as @p input_ptr
448 * @param[in] output_stride_x Stride of the destination image in X dimension (in bytes)
449 * @param[in] output_step_x output_stride_x * number of elements along X processed per workitem(in bytes)
450 * @param[in] output_stride_y Stride of the destination image in Y dimension (in bytes)
451 * @param[in] output_step_y output_stride_y * number of elements along Y processed per workitem(in bytes)
452 * @param[in] output_stride_z Stride of the source tensor in Z dimension (in bytes)
453 * @param[in] output_step_z output_stride_z * number of elements along Z processed per workitem(in bytes)
454 * @param[in] output_offset_first_element_in_bytes The offset of the first element in the destination image
455 */
Michalis Spyroue74b2012018-04-18 09:49:16 +0100[diff] [blame]456__kernel void pooling_layer_MxN_nchw(
Gian Marco Iodicebf179552017-09-05 13:51:21 +0100[diff] [blame]457 TENSOR3D_DECLARATION(input),
458 TENSOR3D_DECLARATION(output))
459{
460 // Get pixels pointer
461 Tensor3D input = CONVERT_TO_TENSOR3D_STRUCT(input);
462 Tensor3D output = CONVERT_TO_TENSOR3D_STRUCT(output);
463
464 VEC_DATA_TYPE(DATA_TYPE, 8)
465 vdata = INITIAL_VALUE;
466 DATA_TYPE sdata = INITIAL_VALUE;
467
468 // Load data
Isabella Gottardia527e8c2018-01-31 17:49:25 +0000[diff] [blame]469 for(int y = 0; y < POOL_SIZE_Y; y++)
Gian Marco Iodicebf179552017-09-05 13:51:21 +0100[diff] [blame]470 {
471 int x = 0;
Isabella Gottardia527e8c2018-01-31 17:49:25 +0000[diff] [blame]472 for(; x <= ((int)POOL_SIZE_X - 8); x += 8)
Gian Marco Iodicebf179552017-09-05 13:51:21 +0100[diff] [blame]473 {
474 VEC_DATA_TYPE(DATA_TYPE, 8)
475 data0 = vload8(0, (__global DATA_TYPE *)tensor3D_offset(&input, x, y, 0));
Georgios Pinitascdf51452017-08-31 14:21:36 +0100[diff] [blame]476#if defined(POOL_L2)
477 // Raise to power of 2 for L2 Pooling
478 data0 *= data0;
479#endif /* defined(POOL_L2) */
Gian Marco Iodicebf179552017-09-05 13:51:21 +0100[diff] [blame]480 vdata = POOL_OP(vdata, data0);
481 }
482
483 // Leftover
Isabella Gottardia527e8c2018-01-31 17:49:25 +0000[diff] [blame]484 for(; x < (int)POOL_SIZE_X; ++x)
Gian Marco Iodicebf179552017-09-05 13:51:21 +0100[diff] [blame]485 {
486 DATA_TYPE data0 = *((__global DATA_TYPE *)tensor3D_offset(&input, x, y, 0));
Georgios Pinitascdf51452017-08-31 14:21:36 +0100[diff] [blame]487#if defined(POOL_L2)
488 // Raise to power of 2 for L2 Pooling
489 data0 *= data0;
490#endif /* defined(POOL_L2) */
491 sdata = POOL_OP(sdata, data0);
Gian Marco Iodicebf179552017-09-05 13:51:21 +0100[diff] [blame]492 }
493 }
494
495 // Reduce result
496 VEC_DATA_TYPE(DATA_TYPE, 4)
497 reduce4 = POOL_OP(vdata.s0123, vdata.s4567);
498 VEC_DATA_TYPE(DATA_TYPE, 2)
499 reduce2 = POOL_OP(reduce4.s01, reduce4.s23);
500 DATA_TYPE res = POOL_OP(reduce2.s0, reduce2.s1);
501 res = POOL_OP(res, sdata);
502
Georgios Pinitascdf51452017-08-31 14:21:36 +0100[diff] [blame]503#if defined(POOL_AVG) || defined(POOL_L2)
Gian Marco Iodicebf179552017-09-05 13:51:21 +0100[diff] [blame]504 // Divide by pool region in case of average pooling
Isabella Gottardia527e8c2018-01-31 17:49:25 +0000[diff] [blame]505 res = DIV_OP(res, calculate_avg_scale(POOL_SIZE_X, POOL_SIZE_Y, MAX_WIDTH, MAX_HEIGHT, PAD_X, PAD_Y, STRIDE_X, STRIDE_Y));
Georgios Pinitascdf51452017-08-31 14:21:36 +0100[diff] [blame]506#endif /* defined(POOL_AVG) || defined(POOL_L2) */
507
508#if defined(POOL_L2)
509 // Take square root of the result in L2 pooling
510 res = SQRT_OP(res);
511#endif /* defined(POOL_L2) */
Gian Marco Iodicebf179552017-09-05 13:51:21 +0100[diff] [blame]512
513 // Store result
514 *(__global DATA_TYPE *)output.ptr = res;
515}
Isabella Gottardia527e8c2018-01-31 17:49:25 +0000[diff] [blame]516#endif // defined(POOL_SIZE_X) && defined(POOL_SIZE_Y)
Michalis Spyroue74b2012018-04-18 09:49:16 +0100[diff] [blame]517
518DATA_TYPE calculate_avg_scale_nhwc(const int pool_size_x, const int pool_size_y, int upper_bound_w, int upper_bound_h,
519 const int pad_x, const int pad_y, const int stride_x, const int stride_y)
520{
521 int start_x = get_global_id(1) * stride_x - pad_x;
522 int start_y = get_global_id(2) * stride_y - pad_y;
523
524#if !defined(EXCLUDE_PADDING)
525 upper_bound_w += pad_x;
526 upper_bound_h += pad_y;
527#endif /* defined(EXCLUDE_PADDING) */
528 const int end_x = min(start_x + pool_size_x, upper_bound_w);
529 const int end_y = min(start_y + pool_size_y, upper_bound_h);
530#if defined(EXCLUDE_PADDING)
531 start_x = max(0, start_x);
532 start_y = max(0, start_y);
533#endif /* defined(EXCLUDE_PADDING) */
534 return ((end_y - start_y) * (end_x - start_x));
535}
536
537/** Performs a pooling function of pool size equal to N (NHWC)
538 *
539 * @note Datatype must be passed using -DDATA_TYPE e.g. -DDATA_TYPE=float. Supported data types are F16/F32
540 * @note -DFP16 must be passed at compile time if half float data type is used
541 * @note Pool sizes must be passed using -DPOOL_SIZE_X and -DPOOL_SIZE_Y e.g. -DPOOL_SIZE_X=13;
542 * @note Tensors width and height must be passed at compile time using -DMAX_WIDTH and -DMAX_HEIGHT
543 * @note Strides must be passed at compile time using -DSTRIDE_X and -DSTRIDE_Y which are the steps of the window along the x and y directions
544 * @note Pad values must be passed at compile time using -DPAD_X and -DPAD_Y which are the pooling paddings in x and y dimension
545 * @note In case of average pooling the following information must be passed at compile time:
546 * -DPOOL_AVG must be provided otherwise max pooling will be performed.
547 *
548 * @param[in] input_ptr Pointer to the source image. Supported data types: F16/F32
549 * @param[in] input_stride_x Stride of the source image in X dimension (in bytes)
550 * @param[in] input_step_x input_stride_x * number of elements along X processed per workitem(in bytes)
551 * @param[in] input_stride_y Stride of the source image in Y dimension (in bytes)
552 * @param[in] input_step_y input_stride_y * number of elements along Y processed per workitem(in bytes)
553 * @param[in] input_stride_z Stride of the source tensor in Z dimension (in bytes)
554 * @param[in] input_step_z input_stride_z * number of elements along Z processed per workitem(in bytes)
555 * @param[in] input_offset_first_element_in_bytes The offset of the first element in the source image
556 * @param[out] output_ptr Pointer to the destination image. Supported data types: same as @p input_ptr
557 * @param[in] output_stride_x Stride of the destination image in X dimension (in bytes)
558 * @param[in] output_step_x output_stride_x * number of elements along X processed per workitem(in bytes)
559 * @param[in] output_stride_y Stride of the destination image in Y dimension (in bytes)
560 * @param[in] output_step_y output_stride_y * number of elements along Y processed per workitem(in bytes)
561 * @param[in] output_stride_z Stride of the source tensor in Z dimension (in bytes)
562 * @param[in] output_step_z output_stride_z * number of elements along Z processed per workitem(in bytes)
563 * @param[in] output_offset_first_element_in_bytes The offset of the first element in the destination image
564 */
565__kernel void pooling_layer_MxN_nhwc(
566 TENSOR3D_DECLARATION(input),
567 TENSOR3D_DECLARATION(output))
568{
569 // Get pixels pointer
570 Tensor3D input = CONVERT_TO_TENSOR3D_STRUCT(input);
571 Tensor3D output = CONVERT_TO_TENSOR3D_STRUCT(output);
572
573 VEC_DATA_TYPE(DATA_TYPE, 8)
574 vdata = INITIAL_VALUE;
575 DATA_TYPE sdata = INITIAL_VALUE;
576
577 const int idx_width = get_global_id(1) * STRIDE_X;
578 const int idx_height = get_global_id(2) * STRIDE_Y;
579
580 for(int y = 0; y < POOL_SIZE_Y; ++y)
581 {
582 int y1 = select(y, PAD_Y - idx_height, y + idx_height < PAD_Y || y + idx_height > MAX_HEIGHT);
583 for(int x = 0; x < POOL_SIZE_X; ++x)
584 {
585 int x1 = select(x, PAD_X - idx_width - 1, x + idx_width < PAD_X || x + idx_width > MAX_WIDTH);
586 x1 = select(x1, PAD_X - idx_width - 1, y != y1);
587
588 VEC_DATA_TYPE(DATA_TYPE, 8)
589 data0 = vload8(0, (__global DATA_TYPE *)tensor3D_offset(&input, 0, x1 - PAD_X, y1 - PAD_Y));
590#if defined(POOL_L2)
591 // Raise to power of 2 for L2 Pooling
592 data0 *= data0;
593#endif /* defined(POOL_L2) */
594 vdata = POOL_OP(vdata, data0);
595 }
596 }
597
598#if defined(POOL_AVG) || defined(POOL_L2)
599 // Divide by pool region in case of average pooling
600 vdata = DIV_OP_NHWC(vdata, calculate_avg_scale_nhwc(POOL_SIZE_X, POOL_SIZE_Y, MAX_WIDTH, MAX_HEIGHT, PAD_X, PAD_Y, STRIDE_X, STRIDE_Y));
601#endif /* defined(POOL_AVG) || defined(POOL_L2) */
602
603#if defined(POOL_L2)
604 // Take square root of the result in L2 pooling
605 vdata = SQRT_OP(vdata);
606#endif /* defined(POOL_L2) */
607
608 // Store result
609 vstore8(vdata, 0, (__global DATA_TYPE *)output.ptr);
610}