Gitiles
Code Review Sign In
review.mlplatform.org / ml / ComputeLibrary / a527e8c0ac7a82de4618dfe6aa312d4f6ca2e485 / . / src / core / CL / cl_kernels / pooling_layer.cl
blob: dae0b999083500e1b3d1a29cc5fedd3e54619efc [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
Gian Marco Iodicecb292832017-08-02 13:19:48 +0100[diff] [blame]65#if STRIDE_X == 1
66#define POOLING3x3(res, input, output) POOLING3x3_STRIDE1(res, input, output)
67#elif STRIDE_X == 2 /* STRIDE_X == 1 */
68#define POOLING3x3(res, input, output) POOLING3x3_STRIDE2(res, input, output)
69#elif STRIDE_X == 3 /* STRIDE_X not equals 1 or 2 */
70#define POOLING3x3(res, input, output) POOLING3x3_STRIDE3(res, input, output)
71#endif /* STRIDE_X == 3 */
72
Gian Marco Iodicecb292832017-08-02 13:19:48 +0100[diff] [blame]73#define POOLING3x3_STRIDE1(res, input, output) \
74 ({ \
75 VEC_DATA_TYPE(DATA_TYPE, 4) \
76 data00 = vload4(0, (__global DATA_TYPE *)tensor3D_offset(&input, 0, 0, 0)); \
77 VEC_DATA_TYPE(DATA_TYPE, 2) \
78 data01 = vload2(0, (__global DATA_TYPE *)tensor3D_offset(&input, 0, 0, 0) + 4); \
79 VEC_DATA_TYPE(DATA_TYPE, 4) \
80 data10 = vload4(0, (__global DATA_TYPE *)tensor3D_offset(&input, 0, 1, 0)); \
81 VEC_DATA_TYPE(DATA_TYPE, 2) \
82 data11 = vload2(0, (__global DATA_TYPE *)tensor3D_offset(&input, 0, 1, 0) + 4); \
83 VEC_DATA_TYPE(DATA_TYPE, 4) \
84 data20 = vload4(0, (__global DATA_TYPE *)tensor3D_offset(&input, 0, 2, 0)); \
85 VEC_DATA_TYPE(DATA_TYPE, 2) \
86 data21 = vload2(0, (__global DATA_TYPE *)tensor3D_offset(&input, 0, 2, 0) + 4); \
Georgios Pinitascdf51452017-08-31 14:21:36 +0100[diff] [blame]87 data00 = POW2_OP(data00, 4); \
88 data01 = POW2_OP(data01, 2); \
89 data10 = POW2_OP(data10, 4); \
90 data11 = POW2_OP(data11, 2); \
91 data20 = POW2_OP(data20, 4); \
92 data21 = POW2_OP(data21, 2); \
Gian Marco Iodicecb292832017-08-02 13:19:48 +0100[diff] [blame]93 \
94 VEC_DATA_TYPE(DATA_TYPE, 8) \
95 values00 = (VEC_DATA_TYPE(DATA_TYPE, 8))(data00.s01212323); \
96 VEC_DATA_TYPE(DATA_TYPE, 4) \
97 values01 = (VEC_DATA_TYPE(DATA_TYPE, 4))(data01.s0, data00.s3, data01.s01); \
98 VEC_DATA_TYPE(DATA_TYPE, 8) \
99 values10 = (VEC_DATA_TYPE(DATA_TYPE, 8))(data10.s01212323); \
100 VEC_DATA_TYPE(DATA_TYPE, 4) \
101 values11 = (VEC_DATA_TYPE(DATA_TYPE, 4))(data11.s0, data10.s3, data11.s01); \
102 VEC_DATA_TYPE(DATA_TYPE, 8) \
103 values20 = (VEC_DATA_TYPE(DATA_TYPE, 8))(data20.s01212323); \
104 VEC_DATA_TYPE(DATA_TYPE, 4) \
105 values21 = (VEC_DATA_TYPE(DATA_TYPE, 4))(data21.s0, data20.s3, data21.s01); \
106 \
107 values00 = POOL_OP(values00, values10); \
108 values01 = POOL_OP(values01, values11); \
109 values00 = POOL_OP(values00, values20); \
110 values01 = POOL_OP(values01, values21); \
111 \
112 res = POOL_OP((VEC_DATA_TYPE(DATA_TYPE, 4))(values00.s036, values01.s1), (VEC_DATA_TYPE(DATA_TYPE, 4))(values00.s147, values01.s2)); \
113 res = POOL_OP(res, (VEC_DATA_TYPE(DATA_TYPE, 4))(values00.s25, values01.s03)); \
114 })
115
116#define POOLING3x3_STRIDE2(res, input, output) \
117 ({ \
118 VEC_DATA_TYPE(DATA_TYPE, 8) \
119 data00 = vload8(0, (__global DATA_TYPE *)tensor3D_offset(&input, 0, 0, 0)); \
120 DATA_TYPE data01 = *((__global DATA_TYPE *)tensor3D_offset(&input, 0, 0, 0) + 8); \
121 VEC_DATA_TYPE(DATA_TYPE, 8) \
122 data10 = vload8(0, (__global DATA_TYPE *)tensor3D_offset(&input, 0, 1, 0)); \
123 DATA_TYPE data11 = *((__global DATA_TYPE *)tensor3D_offset(&input, 0, 1, 0) + 8); \
124 VEC_DATA_TYPE(DATA_TYPE, 8) \
125 data20 = vload8(0, (__global DATA_TYPE *)tensor3D_offset(&input, 0, 2, 0)); \
126 DATA_TYPE data21 = *((__global DATA_TYPE *)tensor3D_offset(&input, 0, 2, 0) + 8); \
Georgios Pinitascdf51452017-08-31 14:21:36 +0100[diff] [blame]127 data00 = POW2_OP(data00, 8); \
128 data01 = POW2_OP(data01, 1); \
129 data10 = POW2_OP(data10, 8); \
130 data11 = POW2_OP(data11, 1); \
131 data20 = POW2_OP(data20, 8); \
132 data21 = POW2_OP(data21, 1); \
Gian Marco Iodicecb292832017-08-02 13:19:48 +0100[diff] [blame]133 \
134 VEC_DATA_TYPE(DATA_TYPE, 8) \
135 values00 = (VEC_DATA_TYPE(DATA_TYPE, 8))(data00.s01223445); \
136 VEC_DATA_TYPE(DATA_TYPE, 4) \
137 values01 = (VEC_DATA_TYPE(DATA_TYPE, 4))(data00.s667, data01); \
138 VEC_DATA_TYPE(DATA_TYPE, 8) \
139 values10 = (VEC_DATA_TYPE(DATA_TYPE, 8))(data10.s01223445); \
140 VEC_DATA_TYPE(DATA_TYPE, 4) \
141 values11 = (VEC_DATA_TYPE(DATA_TYPE, 4))(data10.s667, data11); \
142 VEC_DATA_TYPE(DATA_TYPE, 8) \
143 values20 = (VEC_DATA_TYPE(DATA_TYPE, 8))(data20.s01223445); \
144 VEC_DATA_TYPE(DATA_TYPE, 4) \
145 values21 = (VEC_DATA_TYPE(DATA_TYPE, 4))(data20.s667, data21); \
146 \
147 values00 = POOL_OP(values00, values10); \
148 values01 = POOL_OP(values01, values11); \
149 values00 = POOL_OP(values00, values20); \
150 values01 = POOL_OP(values01, values21); \
151 \
152 res = POOL_OP((VEC_DATA_TYPE(DATA_TYPE, 4))(values00.s036, values01.s1), (VEC_DATA_TYPE(DATA_TYPE, 4))(values00.s147, values01.s2)); \
153 res = POOL_OP(res, (VEC_DATA_TYPE(DATA_TYPE, 4))(values00.s25, values01.s03)); \
154 })
155
156#define POOLING3x3_STRIDE3(res, input, output) \
157 ({ \
158 VEC_DATA_TYPE(DATA_TYPE, 8) \
159 data00 = vload8(0, (__global DATA_TYPE *)tensor3D_offset(&input, 0, 0, 0)); \
160 VEC_DATA_TYPE(DATA_TYPE, 4) \
161 data01 = vload4(0, (__global DATA_TYPE *)tensor3D_offset(&input, 0, 0, 0) + 8); \
162 VEC_DATA_TYPE(DATA_TYPE, 8) \
163 data10 = vload8(0, (__global DATA_TYPE *)tensor3D_offset(&input, 0, 1, 0)); \
164 VEC_DATA_TYPE(DATA_TYPE, 4) \
165 data11 = vload4(0, (__global DATA_TYPE *)tensor3D_offset(&input, 0, 1, 0) + 8); \
166 VEC_DATA_TYPE(DATA_TYPE, 8) \
167 data20 = vload8(0, (__global DATA_TYPE *)tensor3D_offset(&input, 0, 2, 0)); \
168 VEC_DATA_TYPE(DATA_TYPE, 4) \
169 data21 = vload4(0, (__global DATA_TYPE *)tensor3D_offset(&input, 0, 2, 0) + 8); \
Georgios Pinitascdf51452017-08-31 14:21:36 +0100[diff] [blame]170 data00 = POW2_OP(data00, 8); \
171 data01 = POW2_OP(data01, 4); \
172 data10 = POW2_OP(data10, 8); \
173 data11 = POW2_OP(data11, 4); \
174 data20 = POW2_OP(data20, 8); \
175 data21 = POW2_OP(data21, 4); \
Gian Marco Iodicecb292832017-08-02 13:19:48 +0100[diff] [blame]176 \
177 data00 = POOL_OP(data00, data10); \
178 data01 = POOL_OP(data01, data11); \
179 data00 = POOL_OP(data00, data20); \
180 data01 = POOL_OP(data01, data21); \
181 \
182 res = POOL_OP((VEC_DATA_TYPE(DATA_TYPE, 4))(data00.s036, data01.s1), (VEC_DATA_TYPE(DATA_TYPE, 4))(data00.s147, data01.s2)); \
183 res = POOL_OP(res, (VEC_DATA_TYPE(DATA_TYPE, 4))(data00.s25, data01.s03)); \
184 })
185
Isabella Gottardia527e8c2018-01-31 17:49:25 +0000[diff] [blame^]186DATA_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]187 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]188{
Georgios Pinitasadaae7e2017-10-30 15:56:32 +0000[diff] [blame]189 int start_x = get_global_id(0) * stride_x - pad_x;
190 int start_y = get_global_id(1) * stride_y - pad_y;
Isabella Gottardia527e8c2018-01-31 17:49:25 +0000[diff] [blame^]191 const int end_x = min(start_x + pool_size_x, upper_bound_w);
192 const int end_y = min(start_y + pool_size_y, upper_bound_h);
Georgios Pinitasadaae7e2017-10-30 15:56:32 +0000[diff] [blame]193#if defined(EXCLUDE_PADDING)
194 start_x = max(0, start_x);
195 start_y = max(0, start_y);
196#endif /* defined(EXCLUDE_PADDING) */
steniu010c7614f2017-06-23 17:00:26 +0100[diff] [blame]197 return ((end_y - start_y) * (end_x - start_x));
Gian Marco Iodicecb292832017-08-02 13:19:48 +0100[diff] [blame]198}
199
Anthony Barbier6ff3b192017-09-04 18:44:23 +0100[diff] [blame]200/** Performs a pooling function of pool size equal to 2.
201 *
steniu010c7614f2017-06-23 17:00:26 +0100[diff] [blame]202 * @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]203 * @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]204 * -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]205 * -DMAX_WIDTH and -DMAX_HEIGHT which are the maximum accessible indeces in x and y dimensions (width + pad)
206 * -DSTRIDE_X and -DSTRIDE_Y which are the steps of the window along the x and y directions
207 * -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]208 *
steniu010c7614f2017-06-23 17:00:26 +0100[diff] [blame]209 * @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]210 * @param[in] input_stride_x Stride of the source image in X dimension (in bytes)
211 * @param[in] input_step_x input_stride_x * number of elements along X processed per workitem(in bytes)
212 * @param[in] input_stride_y Stride of the source image in Y dimension (in bytes)
213 * @param[in] input_step_y input_stride_y * number of elements along Y processed per workitem(in bytes)
214 * @param[in] input_stride_z Stride of the source tensor in Z dimension (in bytes)
215 * @param[in] input_step_z input_stride_z * number of elements along Z processed per workitem(in bytes)
216 * @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]217 * @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]218 * @param[in] output_stride_x Stride of the destination image in X dimension (in bytes)
219 * @param[in] output_step_x output_stride_x * number of elements along X processed per workitem(in bytes)
220 * @param[in] output_stride_y Stride of the destination image in Y dimension (in bytes)
221 * @param[in] output_step_y output_stride_y * number of elements along Y processed per workitem(in bytes)
222 * @param[in] output_stride_z Stride of the source tensor in Z dimension (in bytes)
223 * @param[in] output_step_z output_stride_z * number of elements along Z processed per workitem(in bytes)
224 * @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]225 */
226__kernel void pooling_layer_2(
227 TENSOR3D_DECLARATION(input),
Gian Marco Iodicecb292832017-08-02 13:19:48 +0100[diff] [blame]228 TENSOR3D_DECLARATION(output))
Anthony Barbier6ff3b192017-09-04 18:44:23 +0100[diff] [blame]229{
230 // Get pixels pointer
231 Tensor3D input = CONVERT_TO_TENSOR3D_STRUCT(input);
232 Tensor3D output = CONVERT_TO_TENSOR3D_STRUCT(output);
233
234 // Load data
235 VEC_DATA_TYPE(DATA_TYPE, 2)
236 data0 = vload2(0, (__global DATA_TYPE *)tensor3D_offset(&input, 0, 0, 0));
237 VEC_DATA_TYPE(DATA_TYPE, 2)
238 data1 = vload2(0, (__global DATA_TYPE *)tensor3D_offset(&input, 0, 1, 0));
239
Georgios Pinitascdf51452017-08-31 14:21:36 +0100[diff] [blame]240#if defined(POOL_L2)
241 // Raise to power of 2 for L2 Pooling
242 data0 = POW2_OP(data0, 2);
243 data1 = POW2_OP(data1, 2);
244#endif /* defined(POOL_L2) */
245
Anthony Barbier6ff3b192017-09-04 18:44:23 +0100[diff] [blame]246 // Perform calculations
247 data0 = POOL_OP(data0, data1);
248 DATA_TYPE res = POOL_OP(data0.s0, data0.s1);
249
Georgios Pinitascdf51452017-08-31 14:21:36 +0100[diff] [blame]250#if defined(POOL_AVG) || defined(POOL_L2)
251 // Divide by pool region in case of average or l2 pooling
Isabella Gottardia527e8c2018-01-31 17:49:25 +0000[diff] [blame^]252 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]253#endif /* defined(POOL_AVG) || defined(POOL_L2) */
254
255#if defined(POOL_L2)
256 // Take square root of the result in L2 pooling
257 res = SQRT_OP(res);
258#endif /* defined(POOL_L2) */
Anthony Barbier6ff3b192017-09-04 18:44:23 +0100[diff] [blame]259
260 // Store result
261 *(__global DATA_TYPE *)output.ptr = res;
262}
263
Gian Marco Iodicecb292832017-08-02 13:19:48 +0100[diff] [blame]264/** Performs a pooling function of pool size equal to 3
Anthony Barbier6ff3b192017-09-04 18:44:23 +0100[diff] [blame]265 *
steniu010c7614f2017-06-23 17:00:26 +0100[diff] [blame]266 * @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]267 * @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]268 * -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]269 * -DMAX_WIDTH and -DMAX_HEIGHT which are the maximum accessible indeces in x and y dimensions (width + pad)
270 * -DSTRIDE_X and -DSTRIDE_Y which are the steps of the window along the x and y directions
271 * -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]272 *
steniu010c7614f2017-06-23 17:00:26 +0100[diff] [blame]273 * @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]274 * @param[in] input_stride_x Stride of the source image in X dimension (in bytes)
275 * @param[in] input_step_x input_stride_x * number of elements along X processed per workitem(in bytes)
276 * @param[in] input_stride_y Stride of the source image in Y dimension (in bytes)
277 * @param[in] input_step_y input_stride_y * number of elements along Y processed per workitem(in bytes)
278 * @param[in] input_stride_z Stride of the source tensor in Z dimension (in bytes)
279 * @param[in] input_step_z input_stride_z * number of elements along Z processed per workitem(in bytes)
280 * @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]281 * @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]282 * @param[in] output_stride_x Stride of the destination image in X dimension (in bytes)
283 * @param[in] output_step_x output_stride_x * number of elements along X processed per workitem(in bytes)
284 * @param[in] output_stride_y Stride of the destination image in Y dimension (in bytes)
285 * @param[in] output_step_y output_stride_y * number of elements along Y processed per workitem(in bytes)
286 * @param[in] output_stride_z Stride of the source tensor in Z dimension (in bytes)
287 * @param[in] output_step_z output_stride_z * number of elements along Z processed per workitem(in bytes)
288 * @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]289 */
290__kernel void pooling_layer_3(
291 TENSOR3D_DECLARATION(input),
Gian Marco Iodicecb292832017-08-02 13:19:48 +0100[diff] [blame]292 TENSOR3D_DECLARATION(output))
Anthony Barbier6ff3b192017-09-04 18:44:23 +0100[diff] [blame]293{
294 // Get pixels pointer
295 Tensor3D input = CONVERT_TO_TENSOR3D_STRUCT(input);
296 Tensor3D output = CONVERT_TO_TENSOR3D_STRUCT(output);
297
298 // Load data
299 VEC_DATA_TYPE(DATA_TYPE, 3)
300 data0 = vload3(0, (__global DATA_TYPE *)tensor3D_offset(&input, 0, 0, 0));
301 VEC_DATA_TYPE(DATA_TYPE, 3)
302 data1 = vload3(0, (__global DATA_TYPE *)tensor3D_offset(&input, 0, 1, 0));
303 VEC_DATA_TYPE(DATA_TYPE, 3)
304 data2 = vload3(0, (__global DATA_TYPE *)tensor3D_offset(&input, 0, 2, 0));
305
Georgios Pinitascdf51452017-08-31 14:21:36 +0100[diff] [blame]306#if defined(POOL_L2)
307 // Raise to power of 2 for L2 Pooling
308 data0 = POW2_OP(data0, 3);
309 data1 = POW2_OP(data1, 3);
310 data2 = POW2_OP(data2, 3);
311#endif /* defined(POOL_L2) */
312
Anthony Barbier6ff3b192017-09-04 18:44:23 +0100[diff] [blame]313 // Perform calculations
314 data0 = POOL_OP(data0, data1);
315 data0 = POOL_OP(data0, data2);
316 DATA_TYPE res = POOL_OP(POOL_OP(data0.s0, data0.s1), data0.s2);
317
Georgios Pinitascdf51452017-08-31 14:21:36 +0100[diff] [blame]318#if defined(POOL_AVG) || defined(POOL_L2)
Georgios Pinitasce093142017-06-19 16:11:53 +0100[diff] [blame]319 // Divide by pool region in case of average pooling
Isabella Gottardia527e8c2018-01-31 17:49:25 +0000[diff] [blame^]320 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]321#endif /* defined(POOL_AVG) || defined(POOL_L2) */
322
323#if defined(POOL_L2)
324 // Take square root of the result in L2 pooling
325 res = SQRT_OP(res);
326#endif /* defined(POOL_L2) */
Anthony Barbier6ff3b192017-09-04 18:44:23 +0100[diff] [blame]327
328 // Store result
329 *(__global DATA_TYPE *)output.ptr = res;
330}
Georgios Pinitasce093142017-06-19 16:11:53 +0100[diff] [blame]331
steniu010c7614f2017-06-23 17:00:26 +0100[diff] [blame]332#if defined(POOLING3x3) && !defined(FIXED_POINT_POSITION)
333
334#define CONVERT_OP(data_type) convert_##data_type##4
335#define CONVERT_VECTOR4(data_type) CONVERT_OP(data_type)
336
337VEC_DATA_TYPE(DATA_TYPE, 4)
338calculate_avg_scale4(const int pool_size, const int upper_bound_w, const int upper_bound_h,
339 const int pad_x, const int pad_y, const int stride_x, const int stride_y)
340{
Georgios Pinitasadaae7e2017-10-30 15:56:32 +0000[diff] [blame]341 int4 start_x = ((int4)get_global_id(0) * 4 + (int4)(0, 1, 2, 3)) * (int4)stride_x - (int4)pad_x;
342 int start_y = get_global_id(1) * stride_y - pad_y;
steniu010c7614f2017-06-23 17:00:26 +0100[diff] [blame]343 const int4 end_x = min(start_x + (int4)pool_size, (int4)upper_bound_w);
344 const int end_y = min(start_y + pool_size, upper_bound_h);
Georgios Pinitasadaae7e2017-10-30 15:56:32 +0000[diff] [blame]345#if defined(EXCLUDE_PADDING)
346 start_x = max((int4)0, start_x);
347 start_y = max(0, start_y);
348#endif /* defined(EXCLUDE_PADDING) */
steniu010c7614f2017-06-23 17:00:26 +0100[diff] [blame]349 return (VEC_DATA_TYPE(DATA_TYPE, 4))(1.f) / CONVERT_VECTOR4(DATA_TYPE)(((int4)(end_y - start_y)) * (end_x - start_x));
350}
351
Gian Marco Iodicecb292832017-08-02 13:19:48 +0100[diff] [blame]352/** 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]353 *
steniu010c7614f2017-06-23 17:00:26 +0100[diff] [blame]354 * @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]355 * @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]356 * -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]357 * -DMAX_WIDTH and -DMAX_HEIGHT which are the maximum accessible indeces in x and y dimensions (width + pad)
358 * -DSTRIDE_X and -DSTRIDE_Y which are the steps of the window along the x and y directions
359 * -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]360 *
Gian Marco Iodicecb292832017-08-02 13:19:48 +0100[diff] [blame]361 * @param[in] input_ptr Pointer to the source image. Supported data types: F16/F32
Georgios Pinitasce093142017-06-19 16:11:53 +0100[diff] [blame]362 * @param[in] input_stride_x Stride of the source image in X dimension (in bytes)
363 * @param[in] input_step_x input_stride_x * number of elements along X processed per workitem(in bytes)
364 * @param[in] input_stride_y Stride of the source image in Y dimension (in bytes)
365 * @param[in] input_step_y input_stride_y * number of elements along Y processed per workitem(in bytes)
366 * @param[in] input_stride_z Stride of the source tensor in Z dimension (in bytes)
367 * @param[in] input_step_z input_stride_z * number of elements along Z processed per workitem(in bytes)
368 * @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]369 * @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]370 * @param[in] output_stride_x Stride of the destination image in X dimension (in bytes)
371 * @param[in] output_step_x output_stride_x * number of elements along X processed per workitem(in bytes)
372 * @param[in] output_stride_y Stride of the destination image in Y dimension (in bytes)
373 * @param[in] output_step_y output_stride_y * number of elements along Y processed per workitem(in bytes)
374 * @param[in] output_stride_z Stride of the source tensor in Z dimension (in bytes)
375 * @param[in] output_step_z output_stride_z * number of elements along Z processed per workitem(in bytes)
376 * @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]377 */
Anton Lokhmotovaf6204c2017-11-08 09:34:19 +0000[diff] [blame]378__kernel void pooling_layer_optimized_3(
Gian Marco Iodicecb292832017-08-02 13:19:48 +0100[diff] [blame]379 TENSOR3D_DECLARATION(input),
380 TENSOR3D_DECLARATION(output))
381{
382 // Get pixels pointer
383 Tensor3D input = CONVERT_TO_TENSOR3D_STRUCT(input);
384 Tensor3D output = CONVERT_TO_TENSOR3D_STRUCT(output);
385
386 VEC_DATA_TYPE(DATA_TYPE, 4)
387 res;
388
389 // Perform pooling 3x3 for 4 output elements
390 POOLING3x3(res, input, output);
391
Georgios Pinitascdf51452017-08-31 14:21:36 +0100[diff] [blame]392#if defined(POOL_AVG) || defined(POOL_L2)
Gian Marco Iodicecb292832017-08-02 13:19:48 +0100[diff] [blame]393 // Divide by pool region in case of average pooling
Gian Marco Iodicecb292832017-08-02 13:19:48 +0100[diff] [blame]394 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]395#endif /* defined(POOL_AVG) || defined(POOL_L2) */
396
397#if defined(POOL_L2)
398 // Take square root of the result in L2 pooling
399 res = SQRT_OP(res);
400#endif /* defined(POOL_L2) */
Gian Marco Iodicecb292832017-08-02 13:19:48 +0100[diff] [blame]401
402 vstore4(res, 0, (__global DATA_TYPE *)output.ptr);
403}
steniu010c7614f2017-06-23 17:00:26 +0100[diff] [blame]404#endif // defined(POOLING3x3) && !defined(FIXED_POINT_POSITION)
Gian Marco Iodicecb292832017-08-02 13:19:48 +0100[diff] [blame]405
Isabella Gottardia527e8c2018-01-31 17:49:25 +0000[diff] [blame^]406#if defined(POOL_SIZE_X) && defined(POOL_SIZE_Y)
Gian Marco Iodicebf179552017-09-05 13:51:21 +0100[diff] [blame]407
408// Set the initial value for the pooling operation accordingly with the data type
Georgios Pinitascdf51452017-08-31 14:21:36 +0100[diff] [blame]409#if defined(POOL_AVG) || defined(POOL_L2)
Gian Marco Iodicebf179552017-09-05 13:51:21 +0100[diff] [blame]410#define INITIAL_VALUE 0
Georgios Pinitascdf51452017-08-31 14:21:36 +0100[diff] [blame]411#else /* defined(POOL_AVG) || defined(POOL_L2) */
Gian Marco Iodicebf179552017-09-05 13:51:21 +0100[diff] [blame]412#ifdef FIXED_POINT_POSITION
413#define MIN_VAL_EXPAND(type) type##_MIN
414#define MIN_VAL(type) MIN_VAL_EXPAND(type)
415#define INITIAL_VALUE MIN_VAL(DATA_TYPE)
Gian Marco Iodicebf179552017-09-05 13:51:21 +0100[diff] [blame]416#else // FIXED_POINT_POSITION
417#if FP16
418#define INITIAL_VALUE -HALF_MAX
419#else // FP16
420#define INITIAL_VALUE -FLT_MAX
421#endif // FP16
422#endif // FIXED_POINT_POSITION
423
424#endif // POOL_AVG
425
426/** Performs a pooling function of pool size equal to N
427 *
Georgios Pinitas13fc22c2017-10-19 18:35:59 +0100[diff] [blame]428 * @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]429 * @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^]430 * @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]431 * @note In case of average pooling the following information must be passed at compile time:
432 * -DPOOL_AVG must be provided otherwise max pooling will be performed.
433 * -DMAX_WIDTH and -DMAX_HEIGHT which are the maximum accessible indeces in x and y dimensions (width + pad)
434 * -DSTRIDE_X and -DSTRIDE_Y which are the steps of the window along the x and y directions
435 * -DPAD_X and -DPAD_Y which are the pooling paddings in x and y dimension
436 *
Georgios Pinitas13fc22c2017-10-19 18:35:59 +0100[diff] [blame]437 * @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]438 * @param[in] input_stride_x Stride of the source image in X dimension (in bytes)
439 * @param[in] input_step_x input_stride_x * number of elements along X processed per workitem(in bytes)
440 * @param[in] input_stride_y Stride of the source image in Y dimension (in bytes)
441 * @param[in] input_step_y input_stride_y * number of elements along Y processed per workitem(in bytes)
442 * @param[in] input_stride_z Stride of the source tensor in Z dimension (in bytes)
443 * @param[in] input_step_z input_stride_z * number of elements along Z processed per workitem(in bytes)
444 * @param[in] input_offset_first_element_in_bytes The offset of the first element in the source image
445 * @param[out] output_ptr Pointer to the destination image. Supported data types: same as @p input_ptr
446 * @param[in] output_stride_x Stride of the destination image in X dimension (in bytes)
447 * @param[in] output_step_x output_stride_x * number of elements along X processed per workitem(in bytes)
448 * @param[in] output_stride_y Stride of the destination image in Y dimension (in bytes)
449 * @param[in] output_step_y output_stride_y * number of elements along Y processed per workitem(in bytes)
450 * @param[in] output_stride_z Stride of the source tensor in Z dimension (in bytes)
451 * @param[in] output_step_z output_stride_z * number of elements along Z processed per workitem(in bytes)
452 * @param[in] output_offset_first_element_in_bytes The offset of the first element in the destination image
453 */
Isabella Gottardia527e8c2018-01-31 17:49:25 +0000[diff] [blame^]454__kernel void pooling_layer_MxN(
Gian Marco Iodicebf179552017-09-05 13:51:21 +0100[diff] [blame]455 TENSOR3D_DECLARATION(input),
456 TENSOR3D_DECLARATION(output))
457{
458 // Get pixels pointer
459 Tensor3D input = CONVERT_TO_TENSOR3D_STRUCT(input);
460 Tensor3D output = CONVERT_TO_TENSOR3D_STRUCT(output);
461
462 VEC_DATA_TYPE(DATA_TYPE, 8)
463 vdata = INITIAL_VALUE;
464 DATA_TYPE sdata = INITIAL_VALUE;
465
466 // Load data
Isabella Gottardia527e8c2018-01-31 17:49:25 +0000[diff] [blame^]467 for(int y = 0; y < POOL_SIZE_Y; y++)
Gian Marco Iodicebf179552017-09-05 13:51:21 +0100[diff] [blame]468 {
469 int x = 0;
Isabella Gottardia527e8c2018-01-31 17:49:25 +0000[diff] [blame^]470 for(; x <= ((int)POOL_SIZE_X - 8); x += 8)
Gian Marco Iodicebf179552017-09-05 13:51:21 +0100[diff] [blame]471 {
472 VEC_DATA_TYPE(DATA_TYPE, 8)
473 data0 = vload8(0, (__global DATA_TYPE *)tensor3D_offset(&input, x, y, 0));
Georgios Pinitascdf51452017-08-31 14:21:36 +0100[diff] [blame]474#if defined(POOL_L2)
475 // Raise to power of 2 for L2 Pooling
476 data0 *= data0;
477#endif /* defined(POOL_L2) */
Gian Marco Iodicebf179552017-09-05 13:51:21 +0100[diff] [blame]478 vdata = POOL_OP(vdata, data0);
479 }
480
481 // Leftover
Isabella Gottardia527e8c2018-01-31 17:49:25 +0000[diff] [blame^]482 for(; x < (int)POOL_SIZE_X; ++x)
Gian Marco Iodicebf179552017-09-05 13:51:21 +0100[diff] [blame]483 {
484 DATA_TYPE data0 = *((__global DATA_TYPE *)tensor3D_offset(&input, x, y, 0));
Georgios Pinitascdf51452017-08-31 14:21:36 +0100[diff] [blame]485#if defined(POOL_L2)
486 // Raise to power of 2 for L2 Pooling
487 data0 *= data0;
488#endif /* defined(POOL_L2) */
489 sdata = POOL_OP(sdata, data0);
Gian Marco Iodicebf179552017-09-05 13:51:21 +0100[diff] [blame]490 }
491 }
492
493 // Reduce result
494 VEC_DATA_TYPE(DATA_TYPE, 4)
495 reduce4 = POOL_OP(vdata.s0123, vdata.s4567);
496 VEC_DATA_TYPE(DATA_TYPE, 2)
497 reduce2 = POOL_OP(reduce4.s01, reduce4.s23);
498 DATA_TYPE res = POOL_OP(reduce2.s0, reduce2.s1);
499 res = POOL_OP(res, sdata);
500
Georgios Pinitascdf51452017-08-31 14:21:36 +0100[diff] [blame]501#if defined(POOL_AVG) || defined(POOL_L2)
Gian Marco Iodicebf179552017-09-05 13:51:21 +0100[diff] [blame]502 // Divide by pool region in case of average pooling
Isabella Gottardia527e8c2018-01-31 17:49:25 +0000[diff] [blame^]503 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]504#endif /* defined(POOL_AVG) || defined(POOL_L2) */
505
506#if defined(POOL_L2)
507 // Take square root of the result in L2 pooling
508 res = SQRT_OP(res);
509#endif /* defined(POOL_L2) */
Gian Marco Iodicebf179552017-09-05 13:51:21 +0100[diff] [blame]510
511 // Store result
512 *(__global DATA_TYPE *)output.ptr = res;
513}
Isabella Gottardia527e8c2018-01-31 17:49:25 +0000[diff] [blame^]514#endif // defined(POOL_SIZE_X) && defined(POOL_SIZE_Y)