From 13fc22c3c3f609489e53ec706026a2a7991bf367 Mon Sep 17 00:00:00 2001 From: Georgios Pinitas Date: Thu, 19 Oct 2017 18:35:59 +0100 Subject: COMPMID-556: Fix CLPoolingLayer checks Change-Id: Ib76554adf00fb3c1943da634dc948089843f0e78 Reviewed-on: http://mpd-gerrit.cambridge.arm.com/92439 Tested-by: Kaizen Reviewed-by: Anthony Barbier --- src/core/CL/cl_kernels/pooling_layer.cl | 4 ++-- src/core/CL/kernels/CLPoolingLayerKernel.cpp | 3 +-- 2 files changed, 3 insertions(+), 4 deletions(-) (limited to 'src/core/CL') diff --git a/src/core/CL/cl_kernels/pooling_layer.cl b/src/core/CL/cl_kernels/pooling_layer.cl index 99d7e6e01b..5304c71adf 100644 --- a/src/core/CL/cl_kernels/pooling_layer.cl +++ b/src/core/CL/cl_kernels/pooling_layer.cl @@ -515,7 +515,7 @@ __kernel void pooling_layer_7( /** Performs a pooling function of pool size equal to N * - * @note Datatype must be passed using -DDATA_TYPE e.g. -DDATA_TYPE=float. Supported data types are F16/F32; + * @note Datatype must be passed using -DDATA_TYPE e.g. -DDATA_TYPE=float. Supported data types are QS8/QS16/F16/F32; * @note -DFP16 must be passed at compile time if half float data type is used * @note Pool size must be passed using -DPOOL_SIZE e.g. -DPOOL_SIZE=13; * @note In case of average pooling the following information must be passed at compile time: @@ -524,7 +524,7 @@ __kernel void pooling_layer_7( * -DSTRIDE_X and -DSTRIDE_Y which are the steps of the window along the x and y directions * -DPAD_X and -DPAD_Y which are the pooling paddings in x and y dimension * - * @param[in] input_ptr Pointer to the source image. Supported data types: F16/F32 + * @param[in] input_ptr Pointer to the source image. Supported data types: QS8/QS16/F16/F32 * @param[in] input_stride_x Stride of the source image in X dimension (in bytes) * @param[in] input_step_x input_stride_x * number of elements along X processed per workitem(in bytes) * @param[in] input_stride_y Stride of the source image in Y dimension (in bytes) diff --git a/src/core/CL/kernels/CLPoolingLayerKernel.cpp b/src/core/CL/kernels/CLPoolingLayerKernel.cpp index 497e87b2b5..542d5dcf9f 100644 --- a/src/core/CL/kernels/CLPoolingLayerKernel.cpp +++ b/src/core/CL/kernels/CLPoolingLayerKernel.cpp @@ -67,7 +67,6 @@ void CLPoolingLayerKernel::configure(const ICLTensor *input, ICLTensor *output, ARM_COMPUTE_ERROR_ON_DATA_TYPE_CHANNEL_NOT_IN(input, 1, DataType::QS8, DataType::QS16, DataType::F16, DataType::F32); ARM_COMPUTE_ERROR_ON_NULLPTR(output); ARM_COMPUTE_ERROR_ON(pool_pad_x >= pool_size || pool_pad_y >= pool_size); - ARM_COMPUTE_ERROR_ON(pool_size > 7 && is_data_type_fixed_point(input->info()->data_type())); // Check output dimensions std::tie(pooled_w, pooled_h) = scaled_dimensions(input->info()->dimension(0), @@ -118,7 +117,7 @@ void CLPoolingLayerKernel::configure(const ICLTensor *input, ICLTensor *output, } // Create kernel - if(pool_size <= 7) + if((pool_size == 2) || (pool_size == 3) || (pool_size == 7)) { // Check if we have pool3x3 with stride_x less equal than 3. In these cases, run an optimized OpenCL kernel where // each thread computes 4 output elements -- cgit v1.2.1