arm_compute v17.09
[platform/upstream/armcl.git] / src / core / CL / cl_kernels / roi_pooling_layer.cl
1 /*
2  * Copyright (c) 2017 ARM Limited.
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
26 #if DATA_SIZE == 32
27 #define VEC_SIZE 4
28 #define VEC_MAX vec4_max
29 #elif DATA_SIZE == 16
30 #define VEC_SIZE 8
31 #define VEC_MAX vec8_max
32 #else /* DATA_SIZE not equals 32 or 16 */
33 #error "Unsupported data size"
34 #endif /* DATA_SIZE == 32 */
35
36 inline DATA_TYPE vec4_max(VEC_DATA_TYPE(DATA_TYPE, 4) vec)
37 {
38     VEC_DATA_TYPE(DATA_TYPE, 2)
39     temp = fmax(vec.lo, vec.hi);
40     return fmax(temp.x, temp.y);
41 }
42
43 inline DATA_TYPE vec8_max(VEC_DATA_TYPE(DATA_TYPE, 8) vec)
44 {
45     VEC_DATA_TYPE(DATA_TYPE, 4)
46     temp = fmax(vec.lo, vec.hi);
47     return vec4_max(temp);
48 }
49
50 /** Performs a roi pooling on a single output pixel.
51  *
52  * @param[in] input          Pointer to input Tensor3D struct.
53  * @param[in] region_start_x Start x index projected onto the input tensor.
54  * @param[in] region_end_x   End x index projected onto the input tensor.
55  * @param[in] region_start_y Start y index projected onto the input tensor.
56  * @param[in] region_end_y   End y index projected onto the input tensor.
57  * @param[in] pz             z index of the input tensor.
58  *
59  * @return A max pooled value from the region specified in the input tensor.
60  */
61 inline DATA_TYPE roi_pool_1x1(const Tensor3D *input, int region_start_x, int region_end_x, int region_start_y, int region_end_y, int pz)
62 {
63     // Iterate through the pooling region
64     if((region_end_x <= region_start_x) || (region_end_y <= region_start_y))
65     {
66         return (DATA_TYPE)0;
67     }
68     else
69     {
70         int num_iter = (int)((region_end_x - region_start_x) / VEC_SIZE);
71         VEC_DATA_TYPE(DATA_TYPE, VEC_SIZE)
72         curr_max = (VEC_DATA_TYPE(DATA_TYPE, VEC_SIZE))(-FLT_MAX);
73         for(int j = region_start_y; j < region_end_y; ++j)
74         {
75             int i = region_start_x;
76             for(; i < region_start_x + num_iter * VEC_SIZE; i += VEC_SIZE)
77             {
78                 VEC_DATA_TYPE(DATA_TYPE, VEC_SIZE)
79                 val      = VLOAD(VEC_SIZE)(0, (__global DATA_TYPE *)tensor3D_offset(input, i, j, pz));
80                 curr_max = fmax(val, curr_max);
81             }
82             for(; i < region_end_x; ++i)
83             {
84                 DATA_TYPE val = *(__global DATA_TYPE *)tensor3D_offset(input, i, j, pz);
85                 curr_max      = fmax(curr_max, val);
86             }
87         }
88         return (DATA_TYPE)VEC_MAX(curr_max);
89     }
90 }
91
92 /** Performs a roi pooling function.
93  *
94  * @note Datatype must be passed using -DDATA_TYPE e.g. -DDATA_TYPE=float. Supported data types are F16, F32;
95  * @note Datasize must be passed using -DDATA_SIZE e.g. -DDATA_SIZE=32;
96  * @note Input dimensions must be passed using -DMAX_DIM_X, -DMAX_DIM_Y and -DMAX_DIM_Z;
97  * @note Pooled region dimensions must be passed using -DPOOLED_DIM_X and -DPOOLED_DIM_Y;
98  * @note Spatial scale must be passed using -DSPATIAL_SCALE;
99  *
100  * @param[in]  input_ptr                            Pointer to the source image. Supported data types: F16, F32
101  * @param[in]  input_stride_x                       Stride of the source image in X dimension (in bytes)
102  * @param[in]  input_step_x                         input_stride_x * number of elements along X processed per workitem(in bytes)
103  * @param[in]  input_stride_y                       Stride of the source image in Y dimension (in bytes)
104  * @param[in]  input_step_y                         input_stride_y * number of elements along Y processed per workitem(in bytes)
105  * @param[in]  input_stride_z                       Stride of the source tensor in Z dimension (in bytes)
106  * @param[in]  input_step_z                         input_stride_z * number of elements along Z processed per workitem(in bytes)
107  * @param[in]  input_offset_first_element_in_bytes  The offset of the first element in the pooled region of the source image as specifed by ROI
108  * @param[in]  rois_ptr                             Pointer to the rois array. Layout: {x, y, width, height, batch_indx}
109  * @param[in]  rois_stride_x                        Stride of the rois array in X dimension (in bytes)
110  * @param[in]  rois_step_x                          rois_stride_x * number of elements along X processed per workitem(in bytes)
111  * @param[in]  rois_offset_first_element_in_bytes   The offset of the first element in the rois array
112  * @param[out] output_ptr                           Pointer to the destination image. Supported data types: F16, F32
113  * @param[in]  output_stride_x                      Stride of the destination image in X dimension (in bytes)
114  * @param[in]  output_step_x                        output_stride_x * number of elements along X processed per workitem(in bytes)
115  * @param[in]  output_stride_y                      Stride of the destination image in Y dimension (in bytes)
116  * @param[in]  output_step_y                        output_stride_y * number of elements along Y processed per workitem(in bytes)
117  * @param[in]  output_stride_z                      Stride of the destination tensor in Z dimension (in bytes)
118  * @param[in]  output_step_z                        output_stride_z * number of elements along Z processed per workitem(in bytes)
119  * @param[in]  output_offset_first_element_in_bytes The offset of the first element in the destination image
120  * @param[in]  input_stride_w                       Stride of the source image in W dimension (in bytes)
121  * @param[in]  output_stride_w                      Stride of the destination image in W dimension (in bytes)
122  */
123 __kernel void roi_pooling_layer(
124     TENSOR3D_DECLARATION(input),
125     VECTOR_DECLARATION(rois),
126     TENSOR3D_DECLARATION(output),
127     unsigned int input_stride_w, unsigned int output_stride_w)
128 {
129     // Get pixels pointer
130     Tensor3D input  = CONVERT_TO_TENSOR3D_STRUCT_NO_STEP(input);
131     Vector   rois   = CONVERT_TO_VECTOR_STRUCT_NO_STEP(rois);
132     Tensor3D output = CONVERT_TO_TENSOR3D_STRUCT_NO_STEP(output);
133
134     const int px = get_global_id(0);
135     const int py = get_global_id(1);
136     const int pw = get_global_id(2);
137
138     // Load roi parameters
139     // roi is laid out as follows:
140     // { x, y, width, height, batch_index }
141     const ushort4 roi      = vload4(0, (__global ushort *)vector_offset(&rois, pw));
142     const ushort roi_batch = *((__global ushort *)vector_offset(&rois, pw) + 4);
143     const int2 roi_anchor  = convert_int2_sat(round(convert_float2(roi.s01) * (float)SPATIAL_SCALE));
144     const int2 roi_dims    = convert_int2_sat(fmax(round(convert_float2(roi.s23) * (float)SPATIAL_SCALE), 1.f));
145
146     // Calculate pooled region start and end
147     const float2 spatial_indx     = (float2)(px, py);
148     const float2 pooled_dims      = (float2)(POOLED_DIM_X, POOLED_DIM_Y);
149     const int2   max_spatial_dims = (int2)(MAX_DIM_X, MAX_DIM_Y);
150     int2         region_start     = convert_int2_sat(floor(spatial_indx / pooled_dims * convert_float2(roi_dims))) + roi_anchor;
151     int2         region_end       = convert_int2_sat(floor((spatial_indx + 1) / pooled_dims * convert_float2(roi_dims))) + roi_anchor;
152
153     region_start = clamp(region_start, 0, max_spatial_dims);
154     region_end   = clamp(region_end, 0, max_spatial_dims);
155
156     // Move input and output pointer across the fourth dimension
157     input.ptr += roi_batch * input_stride_w;
158     output.ptr += pw * output_stride_w;
159
160     for(int pz = 0; pz < MAX_DIM_Z; ++pz)
161     {
162         *(__global DATA_TYPE *)tensor3D_offset(&output, px, py, pz) = (__global DATA_TYPE)roi_pool_1x1(&input,
163                                                                                                        region_start.x,
164                                                                                                        region_end.x,
165                                                                                                        region_start.y,
166                                                                                                        region_end.y, pz);
167     }
168 }