diff options
author | Michele Di Giorgio <michele.digiorgio@arm.com> | 2019-08-28 16:27:26 +0100 |
---|---|---|
committer | Pablo Marquez <pablo.tello@arm.com> | 2019-09-04 14:32:18 +0000 |
commit | 4aff98fcfd3c736115f3983dc448c3280e570841 (patch) | |
tree | fd9ce801272b0a1aec9f14fbe89923760a5c2182 /src/core/CL/cl_kernels/roi_align_layer_quantized.cl | |
parent | 49be2e32e697f3b7a124018bc3cee91adb5f9478 (diff) | |
download | ComputeLibrary-4aff98fcfd3c736115f3983dc448c3280e570841.tar.gz |
COMPMID-2247: Extend support of CLBoundingBoxTransform for QUANT16_ASYMM
Change-Id: I8af7a382c0bccf55cf7f4a64f46ce9e6cd965afe
Signed-off-by: Michele Di Giorgio <michele.digiorgio@arm.com>
Reviewed-on: https://review.mlplatform.org/c/1833
Comments-Addressed: Arm Jenkins <bsgcomp@arm.com>
Reviewed-by: Pablo Marquez <pablo.tello@arm.com>
Tested-by: Arm Jenkins <bsgcomp@arm.com>
Diffstat (limited to 'src/core/CL/cl_kernels/roi_align_layer_quantized.cl')
-rw-r--r-- | src/core/CL/cl_kernels/roi_align_layer_quantized.cl | 24 |
1 files changed, 2 insertions, 22 deletions
diff --git a/src/core/CL/cl_kernels/roi_align_layer_quantized.cl b/src/core/CL/cl_kernels/roi_align_layer_quantized.cl index f9360e98f1..030731b7d3 100644 --- a/src/core/CL/cl_kernels/roi_align_layer_quantized.cl +++ b/src/core/CL/cl_kernels/roi_align_layer_quantized.cl @@ -21,7 +21,7 @@ * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE * SOFTWARE. */ -#include "helpers.h" +#include "helpers_asymm.h" // This specifies the value to shift the result of roi_dims / pooled_dims before ceiling. // It is close to the epsilon machine (for a floating point system, x and x+EPS are the same number). @@ -29,26 +29,6 @@ #if defined(DATA_TYPE) && defined(POOLED_DIM_X) && defined(POOLED_DIM_Y) && defined(MAX_DIM_X) && defined(MAX_DIM_Y) && defined(MAX_DIM_Z) && defined(SPATIAL_SCALE) && defined(OFFSET_IN) && defined(OFFSET_OUT) && defined(SCALE_IN) && defined(SCALE_OUT) && defined(OFFSET_ROIS) && defined(SCALE_ROIS) // Check for compile time constants -#define CONVERT_RTE(x, type) (convert_##type##_rte((x))) -#define CONVERT_DOWN(x, type) CONVERT_RTE(x, type) -inline float dequantize_qasymm8(uchar input, float offset, float scale) -{ - return ((float)input - offset) * scale; -} - -inline uchar quantize_qasymm8(float input, float offset, float scale) -{ - float out_f32 = input / scale + offset; - uchar res_u8 = CONVERT_SAT(CONVERT_DOWN(out_f32, int), uchar); - return res_u8; -} - -inline float4 dequantize_qasymm16(ushort4 input, float offset, float scale) -{ - float4 in_f32 = (CONVERT(input, float4) - (float4)(offset)) * (float4)(scale); - return in_f32; -} - /** Performs a roi align on a single output pixel. * * @param[in] input Pointer to input Tensor3D struct. @@ -178,7 +158,7 @@ __kernel void roi_align_layer_quantized( // Load roi parameters // roi is laid out as follows { batch_index, x1, y1, x2, y2 } const ushort roi_batch = *((__global ushort *)offset(&rois, 0, pw)); - float4 roi = dequantize_qasymm16(vload4(0, (__global ushort *)offset(&rois, 1, pw)), OFFSET_ROIS, SCALE_ROIS); + float4 roi = DEQUANTIZE(vload4(0, (__global ushort *)offset(&rois, 1, pw)), OFFSET_ROIS, SCALE_ROIS, ushort, 4); float2 roi_anchor = roi.s01 * convert_float(SPATIAL_SCALE); float2 roi_dims = fmax((roi.s23 - roi.s01) * convert_float(SPATIAL_SCALE), 1.f); |