diff options
| author | Gian Marco Iodice <gianmarco.iodice@arm.com> | 2018-08-09 15:50:07 +0100 |
|---|---|---|
| committer | Anthony Barbier <anthony.barbier@arm.com> | 2018-11-02 16:54:54 +0000 |
| commit | 2318fcfd0dc8360126bfec71fff88a2015cbc56d (patch) | |
| tree | 3c152bea201f6c9bb289cdec070d9e95bc4817b9 /src/core/CL/cl_kernels/im2col.cl | |
| parent | d8b03dd029261091e34dd8831d546299c60ce094 (diff) | |
| download | ComputeLibrary-2318fcfd0dc8360126bfec71fff88a2015cbc56d.tar.gz | |
COMPMID-1486 - CLGEMMDilatedConvolutionLayer FP16 / FP32 failing in nightlies
Wrong boundary condition in the im2col3x3_nhwc kernel
Change-Id: I83e9dd9b425fd0e3227decb1da3d08a3f5e2536d
Reviewed-on: https://eu-gerrit-1.euhpc.arm.com/143489
Reviewed-by: Anthony Barbier <anthony.barbier@arm.com>
Tested-by: Jenkins <bsgcomp@arm.com>
Diffstat (limited to 'src/core/CL/cl_kernels/im2col.cl')
| -rw-r--r-- | src/core/CL/cl_kernels/im2col.cl | 4 |
1 files changed, 2 insertions, 2 deletions
diff --git a/src/core/CL/cl_kernels/im2col.cl b/src/core/CL/cl_kernels/im2col.cl index 274ec20046..5db1d6ce33 100644 --- a/src/core/CL/cl_kernels/im2col.cl +++ b/src/core/CL/cl_kernels/im2col.cl @@ -850,7 +850,7 @@ __kernel void im2col3x3_nhwc( #if PAD_TOP != 0 || PAD_LEFT != 0 || PAD_BOTTOM != 0 || PAD_RIGHT != 0 // Replace invalid values with zeros - y_cond = (int)((uint)(yi - (int)PAD_TOP + 1) >= (uint)(SRC_HEIGHT)); + y_cond = (int)((uint)(yi - (int)PAD_TOP + 1 * DILATION_Y) >= (uint)(SRC_HEIGHT)); values3 = select(values3, (VECTOR_N)PAD_VALUE, (VEC_DATA_TYPE(COND_DATA_TYPE, VECTOR_SIZE))y_cond || (VEC_DATA_TYPE(COND_DATA_TYPE, VECTOR_SIZE))(x_cond.s0)); values4 = select(values4, (VECTOR_N)PAD_VALUE, (VEC_DATA_TYPE(COND_DATA_TYPE, VECTOR_SIZE))y_cond || (VEC_DATA_TYPE(COND_DATA_TYPE, VECTOR_SIZE))(x_cond.s1)); values5 = select(values5, (VECTOR_N)PAD_VALUE, (VEC_DATA_TYPE(COND_DATA_TYPE, VECTOR_SIZE))y_cond || (VEC_DATA_TYPE(COND_DATA_TYPE, VECTOR_SIZE))(x_cond.s2)); @@ -875,7 +875,7 @@ __kernel void im2col3x3_nhwc( #if PAD_TOP != 0 || PAD_LEFT != 0 || PAD_BOTTOM != 0 || PAD_RIGHT != 0 // Replace invalid values with PAD_VALUE - y_cond = (int)((uint)(yi - (int)PAD_TOP + 2) >= (uint)(SRC_HEIGHT)); + y_cond = (int)((uint)(yi - (int)PAD_TOP + 2 * DILATION_Y) >= (uint)(SRC_HEIGHT)); values6 = select(values6, (VECTOR_N)PAD_VALUE, (VEC_DATA_TYPE(COND_DATA_TYPE, VECTOR_SIZE))y_cond || (VEC_DATA_TYPE(COND_DATA_TYPE, VECTOR_SIZE))(x_cond.s0)); values7 = select(values7, (VECTOR_N)PAD_VALUE, (VEC_DATA_TYPE(COND_DATA_TYPE, VECTOR_SIZE))y_cond || (VEC_DATA_TYPE(COND_DATA_TYPE, VECTOR_SIZE))(x_cond.s1)); values8 = select(values8, (VECTOR_N)PAD_VALUE, (VEC_DATA_TYPE(COND_DATA_TYPE, VECTOR_SIZE))y_cond || (VEC_DATA_TYPE(COND_DATA_TYPE, VECTOR_SIZE))(x_cond.s2)); |