diff options
| author | Giorgio Arena <giorgio.arena@arm.com> | 2020-10-23 13:14:26 +0100 |
|---|---|---|
| committer | Giorgio Arena <giorgio.arena@arm.com> | 2020-10-23 14:58:54 +0000 |
| commit | 69153b3744542691cc205bafda3bbc9c84d394ce (patch) | |
| tree | 55635d6eb71f775cfbefd8fb5e47bd1a101ca135 | |
| parent | b333758b5e2b00ba7b36cc02c169605f55fb0251 (diff) | |
| download | ComputeLibrary-69153b3744542691cc205bafda3bbc9c84d394ce.tar.gz | |
COMPMID-3880 Nightly failure 947: hikey960 mismatches CL PoolingLayer FP32
Signed-off-by: Giorgio Arena <giorgio.arena@arm.com>
Change-Id: I448f2b52e2089da1b859301eade7b988d816a816
Reviewed-on: https://review.mlplatform.org/c/ml/ComputeLibrary/+/4245
Tested-by: Arm Jenkins <bsgcomp@arm.com>
Reviewed-by: Gian Marco Iodice <gianmarco.iodice@arm.com>
Comments-Addressed: Arm Jenkins <bsgcomp@arm.com>
| -rw-r--r-- | src/core/CL/cl_kernels/pooling_layer.cl | 14 |
1 files changed, 11 insertions, 3 deletions
diff --git a/src/core/CL/cl_kernels/pooling_layer.cl b/src/core/CL/cl_kernels/pooling_layer.cl index e69c3c35e9..680e947149 100644 --- a/src/core/CL/cl_kernels/pooling_layer.cl +++ b/src/core/CL/cl_kernels/pooling_layer.cl @@ -786,6 +786,8 @@ __kernel void pooling_layer_MxN_nhwc( } #endif // defined(POOL_SIZE_X) && defined(POOL_SIZE_Y) +#define SELECT_TYPE SELECT_DATA_TYPE(ACC_DATA_TYPE, VEC_SIZE) + /** Performs pooling layer of size equal to 2. This OpenCL kernel can perform the following pooling types: * -# max, -DPOOL_MAX must be passed at compile time * -# max extracting the max index, -DPOOL_MAX and -DEXTRACT_MAX_INDEX must be passed at compile time @@ -899,10 +901,16 @@ __kernel void pooling_layer_2x2_nhwc( #if !defined(POOL_MAX) if(filter_size != 4) { + SELECT_TYPE cond_w_s = (SELECT_TYPE)idx_in_w < (SELECT_TYPE)0; + SELECT_TYPE cond_w_e = (SELECT_TYPE)idx_in_w >= (SELECT_TYPE)(SRC_WIDTH - 1); + SELECT_TYPE cond_h_s = (SELECT_TYPE)idx_in_h < (SELECT_TYPE)0; + SELECT_TYPE cond_h_e = (SELECT_TYPE)idx_in_h >= (SELECT_TYPE)(SRC_HEIGHT - 1); + // Make invalid the values loaded if the x or y coordinate was clamped (out-of-bound) - data1 = select(data1, (VEC_DATA_TYPE(ACC_DATA_TYPE, VEC_SIZE))INITIAL_VALUE, (SELECT_DATA_TYPE(ACC_DATA_TYPE, VEC_SIZE))(pool_x_e == pool_x_s)); - data2 = select(data2, (VEC_DATA_TYPE(ACC_DATA_TYPE, VEC_SIZE))INITIAL_VALUE, (SELECT_DATA_TYPE(ACC_DATA_TYPE, VEC_SIZE))(pool_y_e == pool_y_s)); - data3 = select(data3, (VEC_DATA_TYPE(ACC_DATA_TYPE, VEC_SIZE))INITIAL_VALUE, (SELECT_DATA_TYPE(ACC_DATA_TYPE, VEC_SIZE))((pool_x_e == pool_x_s) || (pool_y_e == pool_y_s))); + data0 = select(data0, (VEC_DATA_TYPE(ACC_DATA_TYPE, VEC_SIZE))INITIAL_VALUE, (SELECT_TYPE)(cond_w_s | cond_h_s)); + data1 = select(data1, (VEC_DATA_TYPE(ACC_DATA_TYPE, VEC_SIZE))INITIAL_VALUE, (SELECT_TYPE)(cond_w_e | cond_h_s)); + data2 = select(data2, (VEC_DATA_TYPE(ACC_DATA_TYPE, VEC_SIZE))INITIAL_VALUE, (SELECT_TYPE)(cond_w_s | cond_h_e)); + data3 = select(data3, (VEC_DATA_TYPE(ACC_DATA_TYPE, VEC_SIZE))INITIAL_VALUE, (SELECT_TYPE)(cond_w_e | cond_h_e)); } #endif // !defined(POOL_MAX) |