diff options
| author | Giorgio Arena <giorgio.arena@arm.com> | 2021-01-11 15:31:15 +0000 |
|---|---|---|
| committer | Giorgio Arena <giorgio.arena@arm.com> | 2021-01-13 14:01:40 +0000 |
| commit | 3d3a01c2122278febd6e77ec255d2557d0005546 (patch) | |
| tree | 837f7eb76afed5b56af6d64db00ef69e122ee2d2 /src/core/CL/cl_kernels/range.cl | |
| parent | d23a251df7b248067e06d5559e985ae1c523be27 (diff) | |
| download | ComputeLibrary-3d3a01c2122278febd6e77ec255d2557d0005546.tar.gz | |
Remove padding for CLArgMinMaxLayerKernel and fix CLRange mismatches
- Cast the destination pointer to (__global DATA_TYPE*) when VEC_SIZE == 1 in range.cl
Resolves: COMPMID-3906, COMPMID-4093
Signed-off-by: Giorgio Arena <giorgio.arena@arm.com>
Change-Id: Ic0a334d98785ea434ed81f89dbe34e7674991f82
Reviewed-on: https://review.mlplatform.org/c/ml/ComputeLibrary/+/4792
Comments-Addressed: Arm Jenkins <bsgcomp@arm.com>
Tested-by: Arm Jenkins <bsgcomp@arm.com>
Reviewed-by: Michele Di Giorgio <michele.digiorgio@arm.com>
Diffstat (limited to 'src/core/CL/cl_kernels/range.cl')
| -rw-r--r-- | src/core/CL/cl_kernels/range.cl | 14 |
1 files changed, 7 insertions, 7 deletions
diff --git a/src/core/CL/cl_kernels/range.cl b/src/core/CL/cl_kernels/range.cl index 467b962b0d..d25d10e207 100644 --- a/src/core/CL/cl_kernels/range.cl +++ b/src/core/CL/cl_kernels/range.cl @@ -61,7 +61,7 @@ __kernel void range( DATA_TYPE seq; seq = (DATA_TYPE)START + (DATA_TYPE)id * (DATA_TYPE)STEP; - *dst_ptr = seq; + *(__global DATA_TYPE *)dst_ptr = seq; #else // VECTOR_SIZE == 1 VEC_DATA_TYPE(DATA_TYPE, VECTOR_SIZE) seq0 = ((DATA_TYPE)START + (DATA_TYPE)id * (DATA_TYPE)STEP); @@ -108,18 +108,18 @@ __kernel void range_quantized( __global uchar *dst_ptr = out_ptr + out_offset_first_element_in_bytes + id * sizeof(DATA_TYPE); #if VECTOR_SIZE == 1 float seq; - seq = (float)START + (float)id * (float)STEP; - seq = (DATA_TYPE)(int)(seq / ((float)SCALE_OUT) + (float)OFFSET_OUT); - seq = max(0.0f, min(seq, 255.0f)); - *dst_ptr = CONVERT_SAT(CONVERT_DOWN(seq, int), uchar); + seq = (float)START + (float)id * (float)STEP; + seq = (DATA_TYPE)(int)(seq / ((float)SCALE_OUT) + (float)OFFSET_OUT); + seq = max(0.0f, min(seq, 255.0f)); + *(__global DATA_TYPE *)dst_ptr = CONVERT_SAT(CONVERT_DOWN(seq, int), DATA_TYPE); #else // VECTOR_SIZE == 1 VEC_DATA_TYPE(float, VECTOR_SIZE) seq = (float)START + id * (float)STEP; seq = seq + STEP_VEC; seq = seq / ((VEC_DATA_TYPE(float, VECTOR_SIZE))((float)SCALE_OUT)) + ((VEC_DATA_TYPE(float, VECTOR_SIZE))((float)OFFSET_OUT)); seq = max((VEC_DATA_TYPE(float, VECTOR_SIZE))(0.0f), min(seq, (VEC_DATA_TYPE(float, VECTOR_SIZE))(255.0f))); - VEC_DATA_TYPE(uchar, VECTOR_SIZE) - res0 = CONVERT_SAT(CONVERT_DOWN(seq, VEC_DATA_TYPE(int, VECTOR_SIZE)), VEC_DATA_TYPE(uchar, VECTOR_SIZE)); + VEC_DATA_TYPE(DATA_TYPE, VECTOR_SIZE) + res0 = CONVERT_SAT(CONVERT_DOWN(seq, VEC_DATA_TYPE(int, VECTOR_SIZE)), VEC_DATA_TYPE(DATA_TYPE, VECTOR_SIZE)); STORE_VECTOR_SELECT(res, DATA_TYPE, dst_ptr, VECTOR_SIZE, VEC_SIZE_LEFTOVER, VEC_SIZE_LEFTOVER != 0 && get_global_id(0) == 0) #endif // VECTOR_SIZE == 1 } |