diff options
| author | SiCong Li <sicong.li@arm.com> | 2017-09-12 19:06:28 +0100 |
|---|---|---|
| committer | Anthony Barbier <anthony.barbier@arm.com> | 2018-11-02 16:35:24 +0000 |
| commit | cfb6553e6114fd84e1972b1a5c5cc42784996333 (patch) | |
| tree | 8932e23781591c9999613f16ec3b34a2d3890517 /src/core/CL/cl_kernels/roi_pooling_layer.cl | |
| parent | 5b61fd3fbaf41031232296abde56258d12ba3340 (diff) | |
| download | ComputeLibrary-cfb6553e6114fd84e1972b1a5c5cc42784996333.tar.gz | |
COMPMID-417 Fix ROIPooling
* Fix ROIPooling in NEON, CL and Reference.
Change-Id: Id5066625e5073e0bfebe69391f7941e993003296
Reviewed-on: http://mpd-gerrit.cambridge.arm.com/87435
Reviewed-by: Moritz Pflanzer <moritz.pflanzer@arm.com>
Tested-by: Kaizen <jeremy.johnson+kaizengerrit@arm.com>
Reviewed-by: Georgios Pinitas <georgios.pinitas@arm.com>
Reviewed-by: Anthony Barbier <anthony.barbier@arm.com>
Diffstat (limited to 'src/core/CL/cl_kernels/roi_pooling_layer.cl')
| -rw-r--r-- | src/core/CL/cl_kernels/roi_pooling_layer.cl | 17 |
1 files changed, 8 insertions, 9 deletions
diff --git a/src/core/CL/cl_kernels/roi_pooling_layer.cl b/src/core/CL/cl_kernels/roi_pooling_layer.cl index 35a9c0a21f..042b102a15 100644 --- a/src/core/CL/cl_kernels/roi_pooling_layer.cl +++ b/src/core/CL/cl_kernels/roi_pooling_layer.cl @@ -138,24 +138,23 @@ __kernel void roi_pooling_layer( // Load roi parameters // roi is laid out as follows: // { x, y, width, height, batch_index } - const ushort8 roi = vload8(0, (__global ushort *)vector_offset(&rois, pw)); - const int2 roi_anchor = convert_int2_sat(round(convert_float2(roi.s01) * (float)SPATIAL_SCALE)); - const int2 roi_dims = convert_int2_sat(fmax(round(convert_float2(roi.s23) * (float)SPATIAL_SCALE), 1.f)); - - // Determine pooled region in input image to pooled region in output image ratio - const float2 pool_region_ratio = convert_float2(roi_dims) / (float2)(POOLED_DIM_X, POOLED_DIM_Y); + const ushort4 roi = vload4(0, (__global ushort *)vector_offset(&rois, pw)); + const ushort roi_batch = *((__global ushort *)vector_offset(&rois, pw) + 4); + const int2 roi_anchor = convert_int2_sat(round(convert_float2(roi.s01) * (float)SPATIAL_SCALE)); + const int2 roi_dims = convert_int2_sat(fmax(round(convert_float2(roi.s23) * (float)SPATIAL_SCALE), 1.f)); // Calculate pooled region start and end const float2 spatial_indx = (float2)(px, py); + const float2 pooled_dims = (float2)(POOLED_DIM_X, POOLED_DIM_Y); const int2 max_spatial_dims = (int2)(MAX_DIM_X, MAX_DIM_Y); - int2 region_start = convert_int2_sat(floor(spatial_indx * pool_region_ratio)) + roi_anchor; - int2 region_end = convert_int2_sat(ceil((spatial_indx + 1) * pool_region_ratio)) + roi_anchor; + int2 region_start = convert_int2_sat(floor(spatial_indx / pooled_dims * convert_float2(roi_dims))) + roi_anchor; + int2 region_end = convert_int2_sat(floor((spatial_indx + 1) / pooled_dims * convert_float2(roi_dims))) + roi_anchor; region_start = clamp(region_start, 0, max_spatial_dims); region_end = clamp(region_end, 0, max_spatial_dims); // Move input and output pointer across the fourth dimension - input.ptr += roi.s4 * input_stride_w; + input.ptr += roi_batch * input_stride_w; output.ptr += pw * output_stride_w; for(int pz = 0; pz < MAX_DIM_Z; ++pz) |