From 18bdfaefd3aeb2557c2e484b67b06c515f9a120f Mon Sep 17 00:00:00 2001 From: SiCong Li Date: Sun, 8 Nov 2020 21:58:01 +0000 Subject: COMPMID-3951 LargeGraph_FLOAT32_Rank4_25 CTS failures in Android Q in CL Fix1 * Fix CLSpaceToBatchLayerKernel and NESpaceToBatchLayerKernel validation errors by using the correctly calculated output tensor shape Signed-off-by: SiCong Li Change-Id: I21d61f870e6a23a2e38dcb95c939b0bf08082b6f Reviewed-on: https://review.mlplatform.org/c/ml/ComputeLibrary/+/4347 Tested-by: Arm Jenkins Reviewed-by: Michele Di Giorgio Reviewed-by: TeresaARM Reviewed-by: Gian Marco Iodice Comments-Addressed: Arm Jenkins --- arm_compute/core/utils/misc/ShapeCalculator.h | 9 ++++--- .../runtime/CL/functions/CLSpaceToBatchLayer.h | 24 +++++++++---------- .../runtime/NEON/functions/NESpaceToBatchLayer.h | 16 ++++++------- src/core/CL/cl_kernels/space_to_batch.cl | 4 ---- src/core/CL/kernels/CLSpaceToBatchLayerKernel.cpp | 27 +++++++++------------ src/core/CL/kernels/CLSpaceToBatchLayerKernel.h | 28 +++++++++++----------- .../NEON/kernels/NESpaceToBatchLayerKernel.cpp | 26 ++++++++------------ src/core/NEON/kernels/NESpaceToBatchLayerKernel.h | 20 ++++++++-------- 8 files changed, 71 insertions(+), 83 deletions(-) diff --git a/arm_compute/core/utils/misc/ShapeCalculator.h b/arm_compute/core/utils/misc/ShapeCalculator.h index 72b7675749..5ed8aea277 100644 --- a/arm_compute/core/utils/misc/ShapeCalculator.h +++ b/arm_compute/core/utils/misc/ShapeCalculator.h @@ -1157,9 +1157,12 @@ inline TensorShape compute_space_to_batch_shape(const ITensorInfo *input, const const int idx_height = get_data_layout_dimension_index(data_layout, DataLayoutDimension::HEIGHT); const int idx_batch = get_data_layout_dimension_index(data_layout, DataLayoutDimension::BATCHES); - output_shape.set(idx_width, input->tensor_shape()[idx_width] * block_x + padding_left.x() + padding_right.x()); - output_shape.set(idx_height, input->tensor_shape()[idx_height] * block_y + padding_left.y() + padding_right.y()); - output_shape.set(idx_batch, input->tensor_shape()[idx_batch] / (block_x * block_y)); + ARM_COMPUTE_ERROR_ON((input->tensor_shape()[idx_width] + padding_left.x() + padding_right.x()) % block_x != 0); + ARM_COMPUTE_ERROR_ON((input->tensor_shape()[idx_height] + padding_left.y() + padding_right.y()) % block_y != 0); + + output_shape.set(idx_width, (input->tensor_shape()[idx_width] + padding_left.x() + padding_right.x()) / block_x); + output_shape.set(idx_height, (input->tensor_shape()[idx_height] + padding_left.y() + padding_right.y()) / block_y); + output_shape.set(idx_batch, input->tensor_shape()[idx_batch] * block_x * block_y); return output_shape; } diff --git a/arm_compute/runtime/CL/functions/CLSpaceToBatchLayer.h b/arm_compute/runtime/CL/functions/CLSpaceToBatchLayer.h index 1611aa8ed4..5c5e5bed9a 100644 --- a/arm_compute/runtime/CL/functions/CLSpaceToBatchLayer.h +++ b/arm_compute/runtime/CL/functions/CLSpaceToBatchLayer.h @@ -61,8 +61,8 @@ public: /** Set the input and output tensors. * * @param[in] input Tensor input. Supported tensor rank: 4. Data types supported: All. - * @param[in] block_shape 1-D tensor with shape [M]. Data types supported: S32 - * @param[in] paddings 2-D tensor with shape [2, M]. Data types supported: S32 + * @param[in] block_shape 1-D tensor with shape [M]. Supported M: 2. Data types supported: S32 + * @param[in] paddings 2-D tensor with shape [2, M] (First dimension is the fastest-changing dimension). Supported M: 2. Data types supported: S32 * @param[out] output Tensor output. Data types supported: same as @p input */ void configure(const ICLTensor *input, const ICLTensor *block_shape, const ICLTensor *paddings, ICLTensor *output); @@ -70,8 +70,8 @@ public: * * @param[in] compile_context The compile context to be used. * @param[in] input Tensor input. Supported tensor rank: 4. Data types supported: All. - * @param[in] block_shape 1-D tensor with shape [M]. Data types supported: S32 - * @param[in] paddings 2-D tensor with shape [2, M]. Data types supported: S32 + * @param[in] block_shape 1-D tensor with shape [M]. Supported M: 2. Data types supported: S32 + * @param[in] paddings 2-D tensor with shape [2, M] (First dimension is the fastest-changing dimension). Supported M: 2. Data types supported: S32 * @param[out] output Tensor output. Data types supported: same as @p input */ void configure(const CLCompileContext &compile_context, const ICLTensor *input, const ICLTensor *block_shape, const ICLTensor *paddings, ICLTensor *output); @@ -80,8 +80,8 @@ public: * @param[in] input Tensor input. Supported tensor rank: 4. Data types supported: All. * @param[in] block_shape_x Block shape x value. * @param[in] block_shape_y Block shape y value. - * @param[in] padding_left The left padding of the output tensor. - * @param[in] padding_right The right padding of the output tensor. + * @param[in] padding_left The padding at the beginning of every dimension of the output tensor. + * @param[in] padding_right The padding at the end of every dimension of the output tensor. * @param[out] output Tensor output. Data types supported: same as @p input */ void configure(const ICLTensor *input, const int block_shape_x, const int block_shape_y, const Size2D &padding_left, const Size2D &padding_right, ICLTensor *output); @@ -91,8 +91,8 @@ public: * @param[in] input Tensor input. Supported tensor rank: 4. Data types supported: All. * @param[in] block_shape_x Block shape x value. * @param[in] block_shape_y Block shape y value. - * @param[in] padding_left The left padding of the output tensor. - * @param[in] padding_right The right padding of the output tensor. + * @param[in] padding_left The padding at the beginning of every dimension of the output tensor. + * @param[in] padding_right The padding at the end of every dimension of the output tensor. * @param[out] output Tensor output. Data types supported: same as @p input */ void configure(const CLCompileContext &compile_context, const ICLTensor *input, const int block_shape_x, const int block_shape_y, const Size2D &padding_left, const Size2D &padding_right, @@ -100,8 +100,8 @@ public: /** Static function to check if given info will lead to a valid configuration of @ref CLSpaceToBatchLayer * * @param[in] input Tensor input info. Supported tensor rank: 4. Data types supported: All. - * @param[in] block_shape block shape tensor info with shape [M]. Data types supported: S32 - * @param[in] paddings paddings tensor info with shape [2, M]. Data types supported: S32 + * @param[in] block_shape block shape tensor info with shape [M]. Supported M: 2. Data types supported: S32 + * @param[in] paddings paddings tensor info with shape [2, M] (First dimension is the fastest-changing dimension). Supported M: 2. Data types supported: S32 * @param[out] output Tensor output info. Data types supported: same as @p input * * @return a status @@ -112,8 +112,8 @@ public: * @param[in] input Tensor input info. Supported tensor rank: 4. Data types supported: All. * @param[in] block_shape_x Block shape x value. * @param[in] block_shape_y Block shape y value. - * @param[in] padding_left The left padding of the output tensor. - * @param[in] padding_right The right padding of the output tensor. + * @param[in] padding_left The padding at the beginning of every dimension of the output tensor. + * @param[in] padding_right The padding at the end of every dimension of the output tensor. * @param[out] output Tensor output info. Data types supported: same as @p input * * @return a status diff --git a/arm_compute/runtime/NEON/functions/NESpaceToBatchLayer.h b/arm_compute/runtime/NEON/functions/NESpaceToBatchLayer.h index 6df06e87ec..62af092c40 100644 --- a/arm_compute/runtime/NEON/functions/NESpaceToBatchLayer.h +++ b/arm_compute/runtime/NEON/functions/NESpaceToBatchLayer.h @@ -59,8 +59,8 @@ public: /** Set the input and output tensors. * * @param[in] input Tensor input. Supported tensor rank: 4. Data types supported: All. - * @param[in] block_shape 1-D tensor with shape [M]. Data types supported: S32 - * @param[in] paddings 2-D tensor with shape [2, M]. Data types supported: S32 + * @param[in] block_shape 1-D tensor with shape [M]. Supported M: 2. Data types supported: S32 + * @param[in] paddings 2-D tensor with shape [2, M] (First dimension is the fastest-changing dimension). Supported M: 2. Data types supported: S32 * @param[out] output Tensor output. Data types supported: same as @p input */ void configure(const ITensor *input, const ITensor *block_shape, const ITensor *paddings, ITensor *output); @@ -69,16 +69,16 @@ public: * @param[in] input Tensor input. Supported tensor rank: 4. Data types supported: All. * @param[in] block_shape_x Block shape x value. * @param[in] block_shape_y Block shape y value. - * @param[in] padding_left The left padding of the output tensor. - * @param[in] padding_right The right padding of the output tensor. + * @param[in] padding_left The padding at the beginning of every dimension of the output tensor. + * @param[in] padding_right The padding at the end of every dimension of the output tensor. * @param[out] output Tensor output. Data types supported: same as @p input */ void configure(const ITensor *input, const int block_shape_x, const int block_shape_y, const Size2D &padding_left, const Size2D &padding_right, ITensor *output); /** Static function to check if given info will lead to a valid configuration of @ref NESpaceToBatchLayer * * @param[in] input Tensor input info. Supported tensor rank: 4. Data types supported: All. - * @param[in] block_shape block shape tensor info with shape [M]. Data types supported: S32 - * @param[in] paddings paddings tensor info with shape [2, M]. Data types supported: S32 + * @param[in] block_shape 1-D tensor with shape [M]. Supported M: 2. Data types supported: S32 + * @param[in] paddings 2-D tensor with shape [2, M] (First dimension is the fastest-changing dimension). Supported M: 2. Data types supported: S32 * @param[in] output Tensor output info. Data types supported: same as @p input * * @return a status @@ -89,8 +89,8 @@ public: * @param[in] input Tensor input info. Supported tensor rank: 4. Data types supported: All. * @param[in] block_shape_x Block shape x value. * @param[in] block_shape_y Block shape y value. - * @param[in] padding_left The left padding of the output tensor. - * @param[in] padding_right The right padding of the output tensor. + * @param[in] padding_left The padding at the beginning of every dimension of the output tensor. + * @param[in] padding_right The padding at the end of every dimension of the output tensor. * @param[in] output Tensor output info. Data types supported: same as @p input * * @return a status diff --git a/src/core/CL/cl_kernels/space_to_batch.cl b/src/core/CL/cl_kernels/space_to_batch.cl index 5ade9c5a7c..cb11786ac4 100644 --- a/src/core/CL/cl_kernels/space_to_batch.cl +++ b/src/core/CL/cl_kernels/space_to_batch.cl @@ -46,8 +46,6 @@ * @param[in] block_shape_ptr Pointer to the block shape tensor. Supported data types: S32 * @param[in] block_shape_stride_x Stride of the block shape tensor in X dimension (in bytes) * @param[in] block_shape_step_x block_shape_stride_x * number of elements along X processed per workitem(in bytes) - * @param[in] block_shape_stride_y Stride of the block shape tensor in Y dimension (in bytes) - * @param[in] block_shape_step_y block_shape_stride_y * number of elements along Y processed per workitem(in bytes) * @param[in] block_shape_offset_first_element_in_bytes The offset of the first element in the block shapetensor * @param[in] batch_id The output tensor batch id * @param[out] output_ptr Pointer to the destination tensor. Supported data types: same as @p input_ptr @@ -117,8 +115,6 @@ __kernel void space_to_batch_nchw( * @param[in] block_shape_ptr Pointer to the block shape tensor. Supported data types: S32 * @param[in] block_shape_stride_x Stride of the block shape tensor in X dimension (in bytes) * @param[in] block_shape_step_x block_shape_stride_x * number of elements along X processed per workitem(in bytes) - * @param[in] block_shape_stride_y Stride of the block shape tensor in Y dimension (in bytes) - * @param[in] block_shape_step_y block_shape_stride_y * number of elements along Y processed per workitem(in bytes) * @param[in] block_shape_offset_first_element_in_bytes The offset of the first element in the block shapetensor * @param[in] batch_id The output tensor batch id * @param[out] output_ptr Pointer to the destination tensor. Supported data types: same as @p input_ptr diff --git a/src/core/CL/kernels/CLSpaceToBatchLayerKernel.cpp b/src/core/CL/kernels/CLSpaceToBatchLayerKernel.cpp index 91b889a10a..7af0071025 100644 --- a/src/core/CL/kernels/CLSpaceToBatchLayerKernel.cpp +++ b/src/core/CL/kernels/CLSpaceToBatchLayerKernel.cpp @@ -37,15 +37,16 @@ namespace arm_compute { namespace { -Status validate_arguments(const ITensorInfo *input, const ITensorInfo *block_info, const ITensorInfo *padddings, const ITensorInfo *output) +Status validate_arguments(const ITensorInfo *input, const ITensorInfo *block_info, const ITensorInfo *paddings, const ITensorInfo *output) { - ARM_COMPUTE_RETURN_ERROR_ON_NULLPTR(input, block_info, padddings, output); + ARM_COMPUTE_RETURN_ERROR_ON_NULLPTR(input, block_info, paddings, output); ARM_COMPUTE_RETURN_ERROR_ON_DATA_TYPE_CHANNEL_NOT_IN(block_info, 1, DataType::S32); ARM_COMPUTE_RETURN_ERROR_ON(input->data_type() == DataType::UNKNOWN); ARM_COMPUTE_RETURN_ERROR_ON(input->num_dimensions() > 4); ARM_COMPUTE_RETURN_ERROR_ON(block_info->num_dimensions() > 1); - ARM_COMPUTE_RETURN_ERROR_ON(padddings->num_dimensions() > 2); - ARM_COMPUTE_RETURN_ERROR_ON(padddings->tensor_shape()[1] != block_info->tensor_shape()[0]); + ARM_COMPUTE_RETURN_ERROR_ON_MISMATCHING_DIMENSIONS(block_info->tensor_shape(), TensorShape{ 2 }); + ARM_COMPUTE_RETURN_ERROR_ON(paddings->num_dimensions() > 2); + ARM_COMPUTE_RETURN_ERROR_ON_MISMATCHING_DIMENSIONS(paddings->tensor_shape(), TensorShape{ 2, 2 }); // Validate output if initialized if(output->total_size() != 0) @@ -54,6 +55,7 @@ Status validate_arguments(const ITensorInfo *input, const ITensorInfo *block_inf const int idx_channel = get_data_layout_dimension_index(data_layout, DataLayoutDimension::CHANNEL); ARM_COMPUTE_RETURN_ERROR_ON(input->tensor_shape()[idx_channel] != output->tensor_shape()[idx_channel]); ARM_COMPUTE_RETURN_ERROR_ON_MISMATCHING_DATA_TYPES(input, output); + ARM_COMPUTE_RETURN_ERROR_ON_MISMATCHING_QUANTIZATION_INFO(input, output); } return Status{}; @@ -62,22 +64,15 @@ Status validate_arguments_static(const ITensorInfo *input, const int block_shape const ITensorInfo *output) { ARM_COMPUTE_RETURN_ERROR_ON_NULLPTR(input, output); - ARM_COMPUTE_RETURN_ERROR_ON(block_shape_x < 1 || block_shape_y < 1); + ARM_COMPUTE_RETURN_ERROR_ON(input->data_type() == DataType::UNKNOWN); ARM_COMPUTE_RETURN_ERROR_ON(input->num_dimensions() > 4); + ARM_COMPUTE_RETURN_ERROR_ON(block_shape_x < 1 || block_shape_y < 1); // Validate output if initialized if(output->total_size() != 0) { - const DataLayout data_layout = input->data_layout(); - const int idx_width = get_data_layout_dimension_index(data_layout, DataLayoutDimension::WIDTH); - const int idx_height = get_data_layout_dimension_index(data_layout, DataLayoutDimension::HEIGHT); - const int idx_channel = get_data_layout_dimension_index(data_layout, DataLayoutDimension::CHANNEL); - const int idx_batch = get_data_layout_dimension_index(data_layout, DataLayoutDimension::BATCHES); - ARM_COMPUTE_RETURN_ERROR_ON(output->tensor_shape()[idx_width] < padding_left.x() + padding_right.y()); - ARM_COMPUTE_RETURN_ERROR_ON((input->tensor_shape()[idx_width] + padding_left.x() + padding_right.x()) % block_shape_x != 0); - ARM_COMPUTE_RETURN_ERROR_ON((input->tensor_shape()[idx_height] + padding_left.y() + padding_right.y()) % block_shape_y != 0); - ARM_COMPUTE_RETURN_ERROR_ON(input->tensor_shape()[idx_channel] != output->tensor_shape()[idx_channel]); - ARM_COMPUTE_RETURN_ERROR_ON(output->tensor_shape()[idx_batch] % (block_shape_x * block_shape_y) != 0); + TensorShape expected_output_shape = misc::shape_calculator::compute_space_to_batch_shape(input, block_shape_x, block_shape_y, padding_left, padding_right); + ARM_COMPUTE_RETURN_ERROR_ON_MISMATCHING_DIMENSIONS(output->tensor_shape(), expected_output_shape); ARM_COMPUTE_RETURN_ERROR_ON_MISMATCHING_DATA_TYPES(input, output); ARM_COMPUTE_RETURN_ERROR_ON_MISMATCHING_QUANTIZATION_INFO(input, output); } @@ -98,7 +93,7 @@ void CLSpaceToBatchLayerKernel::configure(const ICLTensor *input, const ICLTenso void CLSpaceToBatchLayerKernel::configure(const CLCompileContext &compile_context, const ICLTensor *input, const ICLTensor *block_shape, const ICLTensor *paddings, ICLTensor *output) { - ARM_COMPUTE_ERROR_ON_NULLPTR(input, output); + ARM_COMPUTE_ERROR_ON_NULLPTR(input, block_shape, paddings, output); ARM_COMPUTE_ERROR_THROW_ON(validate_arguments(input->info(), block_shape->info(), paddings->info(), output->info())); _input = input; diff --git a/src/core/CL/kernels/CLSpaceToBatchLayerKernel.h b/src/core/CL/kernels/CLSpaceToBatchLayerKernel.h index 4819c80fce..4817cfeef2 100644 --- a/src/core/CL/kernels/CLSpaceToBatchLayerKernel.h +++ b/src/core/CL/kernels/CLSpaceToBatchLayerKernel.h @@ -50,8 +50,8 @@ public: /** Initialise the kernel's inputs and output. * * @param[in] input Tensor input. Supported tensor rank: 4. Data types supported: All. - * @param[in] block_shape 1-D tensor with shape [M]. Data types supported: S32 - * @param[in] paddings 2-D tensor with shape [2, M]. Data types supported: S32 + * @param[in] block_shape 1-D tensor with shape [M]. Supported M: 2. Data types supported: S32 + * @param[in] paddings 2-D tensor with shape [2, M] (First dimension is the fastest-changing dimension). Supported M: 2. Data types supported: S32 * @param[out] output Tensor output. Data types supported: same as @p input */ void configure(const ICLTensor *input, const ICLTensor *block_shape, const ICLTensor *paddings, ICLTensor *output); @@ -59,8 +59,8 @@ public: * * @param[in] compile_context The compile context to be used. * @param[in] input Tensor input. Supported tensor rank: 4. Data types supported: All. - * @param[in] block_shape 1-D tensor with shape [M]. Data types supported: S32 - * @param[in] paddings 2-D tensor with shape [2, M]. Data types supported: S32 + * @param[in] block_shape 1-D tensor with shape [M]. Supported M: 2. Data types supported: S32 + * @param[in] paddings 2-D tensor with shape [2, M] (First dimension is the fastest-changing dimension). Supported M: 2. Data types supported: S32 * @param[out] output Tensor output. Data types supported: same as @p input */ void configure(const CLCompileContext &compile_context, const ICLTensor *input, const ICLTensor *block_shape, const ICLTensor *paddings, ICLTensor *output); @@ -69,8 +69,8 @@ public: * @param[in] input Tensor input. Supported tensor rank: 4. Data types supported: All. * @param[in] block_shape_x Block shape x value. * @param[in] block_shape_y Block shape y value. - * @param[in] padding_left The left padding of the output tensor. - * @param[in] padding_right The right padding of the output tensor. + * @param[in] padding_left The padding at the beginning of every dimension of the output tensor. + * @param[in] padding_right The padding at the end of every dimension of the output tensor. * @param[out] output Tensor output. Data types supported: same as @p input */ void configure(const ICLTensor *input, const int block_shape_x, const int block_shape_y, const Size2D &padding_left, const Size2D &padding_right, ICLTensor *output); @@ -80,8 +80,8 @@ public: * @param[in] input Tensor input. Supported tensor rank: 4. Data types supported: All. * @param[in] block_shape_x Block shape x value. * @param[in] block_shape_y Block shape y value. - * @param[in] padding_left The left padding of the output tensor. - * @param[in] padding_right The right padding of the output tensor. + * @param[in] padding_left The padding at the beginning of every dimension of the output tensor. + * @param[in] padding_right The padding at the end of every dimension of the output tensor. * @param[out] output Tensor output. Data types supported: same as @p input */ void configure(const CLCompileContext &compile_context, const ICLTensor *input, const int block_shape_x, const int block_shape_y, const Size2D &padding_left, const Size2D &padding_right, @@ -89,8 +89,8 @@ public: /** Static function to check if given info will lead to a valid configuration of @ref CLSpaceToBatchLayerKernel * * @param[in] input Tensor input. Supported tensor rank: 4. Data types supported: All. - * @param[in] block_shape 1-D tensor with shape [M]. Data types supported: S32 - * @param[in] paddings 2-D tensor with shape [2, M]. Data types supported: S32 + * @param[in] block_shape 1-D tensor with shape [M]. Supported M: 2. Data types supported: S32 + * @param[in] paddings 2-D tensor with shape [2, M] (First dimension is the fastest-changing dimension). Supported M: 2. Data types supported: S32 * @param[in] output Tensor output. Data types supported: same as @p input * * @return a status @@ -101,8 +101,8 @@ public: * @param[in] input Tensor input. Supported tensor rank: 4. Data types supported: All. * @param[in] block_shape_x Block shape x value. * @param[in] block_shape_y Block shape y value. - * @param[in] padding_left The left padding of the output tensor. - * @param[in] padding_right The right padding of the output tensor. + * @param[in] padding_left The padding at the beginning of every dimension of the output tensor. + * @param[in] padding_right The padding at the end of every dimension of the output tensor. * @param[in] output Tensor output. Data types supported: same as @p input * * @return a status @@ -114,8 +114,8 @@ public: private: const ICLTensor *_input; /**< Source tensor */ - const ICLTensor *_block_shape; /**< Block shape tensor */ - const ICLTensor *_paddings; /**< Paddings tensor */ + const ICLTensor *_block_shape; /**< Block shape tensor for dynamic evaluation */ + const ICLTensor *_paddings; /**< Paddings tensor for dynamic evaluation */ ICLTensor *_output; /**< Destination tensor */ }; } // namespace arm_compute diff --git a/src/core/NEON/kernels/NESpaceToBatchLayerKernel.cpp b/src/core/NEON/kernels/NESpaceToBatchLayerKernel.cpp index 27b3154298..673eace3c1 100644 --- a/src/core/NEON/kernels/NESpaceToBatchLayerKernel.cpp +++ b/src/core/NEON/kernels/NESpaceToBatchLayerKernel.cpp @@ -41,15 +41,16 @@ namespace arm_compute { namespace { -Status validate_arguments(const ITensorInfo *input, const ITensorInfo *block_info, const ITensorInfo *padddings, const ITensorInfo *output) +Status validate_arguments(const ITensorInfo *input, const ITensorInfo *block_info, const ITensorInfo *paddings, const ITensorInfo *output) { - ARM_COMPUTE_RETURN_ERROR_ON_NULLPTR(input, block_info, padddings, output); + ARM_COMPUTE_RETURN_ERROR_ON_NULLPTR(input, block_info, paddings, output); ARM_COMPUTE_RETURN_ERROR_ON(input->data_type() == DataType::UNKNOWN); ARM_COMPUTE_RETURN_ERROR_ON_DATA_TYPE_CHANNEL_NOT_IN(block_info, 1, DataType::S32); ARM_COMPUTE_RETURN_ERROR_ON(input->num_dimensions() > 4); ARM_COMPUTE_RETURN_ERROR_ON(block_info->num_dimensions() > 1); - ARM_COMPUTE_RETURN_ERROR_ON(padddings->num_dimensions() > 2); - ARM_COMPUTE_RETURN_ERROR_ON(padddings->tensor_shape()[1] != block_info->tensor_shape()[0]); + ARM_COMPUTE_RETURN_ERROR_ON_MISMATCHING_DIMENSIONS(block_info->tensor_shape(), TensorShape{ 2 }); + ARM_COMPUTE_RETURN_ERROR_ON(paddings->num_dimensions() > 2); + ARM_COMPUTE_RETURN_ERROR_ON_MISMATCHING_DIMENSIONS(paddings->tensor_shape(), TensorShape{ 2, 2 }); // Validate output if initialized if(output->total_size() != 0) @@ -58,6 +59,7 @@ Status validate_arguments(const ITensorInfo *input, const ITensorInfo *block_inf const int idx_channel = get_data_layout_dimension_index(data_layout, DataLayoutDimension::CHANNEL); ARM_COMPUTE_RETURN_ERROR_ON(input->tensor_shape()[idx_channel] != output->tensor_shape()[idx_channel]); ARM_COMPUTE_RETURN_ERROR_ON_MISMATCHING_DATA_TYPES(input, output); + ARM_COMPUTE_RETURN_ERROR_ON_MISMATCHING_QUANTIZATION_INFO(input, output); } return Status{}; @@ -67,22 +69,14 @@ Status validate_arguments_static(const ITensorInfo *input, const int block_shape { ARM_COMPUTE_RETURN_ERROR_ON_NULLPTR(input, output); ARM_COMPUTE_RETURN_ERROR_ON(input->data_type() == DataType::UNKNOWN); - ARM_COMPUTE_RETURN_ERROR_ON(block_shape_x < 1 || block_shape_y < 1); ARM_COMPUTE_RETURN_ERROR_ON(input->num_dimensions() > 4); + ARM_COMPUTE_RETURN_ERROR_ON(block_shape_x < 1 || block_shape_y < 1); // Validate output if initialized if(output->total_size() != 0) { - const DataLayout data_layout = input->data_layout(); - const int idx_width = get_data_layout_dimension_index(data_layout, DataLayoutDimension::WIDTH); - const int idx_height = get_data_layout_dimension_index(data_layout, DataLayoutDimension::HEIGHT); - const int idx_channel = get_data_layout_dimension_index(data_layout, DataLayoutDimension::CHANNEL); - const int idx_batch = get_data_layout_dimension_index(data_layout, DataLayoutDimension::BATCHES); - ARM_COMPUTE_RETURN_ERROR_ON(output->tensor_shape()[idx_width] < padding_left.x() + padding_right.y()); - ARM_COMPUTE_RETURN_ERROR_ON((input->tensor_shape()[idx_width] + padding_left.x() + padding_right.x()) % block_shape_x != 0); - ARM_COMPUTE_RETURN_ERROR_ON((input->tensor_shape()[idx_height] + padding_left.y() + padding_right.y()) % block_shape_y != 0); - ARM_COMPUTE_RETURN_ERROR_ON(input->tensor_shape()[idx_channel] != output->tensor_shape()[idx_channel]); - ARM_COMPUTE_RETURN_ERROR_ON(output->tensor_shape()[idx_batch] % (block_shape_x * block_shape_y) != 0); + TensorShape expected_output_shape = misc::shape_calculator::compute_space_to_batch_shape(input, block_shape_x, block_shape_y, padding_left, padding_right); + ARM_COMPUTE_RETURN_ERROR_ON_MISMATCHING_DIMENSIONS(output->tensor_shape(), expected_output_shape); ARM_COMPUTE_RETURN_ERROR_ON_MISMATCHING_DATA_TYPES(input, output); ARM_COMPUTE_RETURN_ERROR_ON_MISMATCHING_QUANTIZATION_INFO(input, output); } @@ -98,7 +92,7 @@ NESpaceToBatchLayerKernel::NESpaceToBatchLayerKernel() void NESpaceToBatchLayerKernel::configure(const ITensor *input, const ITensor *block_shape, const ITensor *paddings, ITensor *output) { - ARM_COMPUTE_ERROR_ON_NULLPTR(input, output); + ARM_COMPUTE_ERROR_ON_NULLPTR(input, block_shape, paddings, output); ARM_COMPUTE_ERROR_THROW_ON(validate_arguments(input->info(), block_shape->info(), paddings->info(), output->info())); _input = input; diff --git a/src/core/NEON/kernels/NESpaceToBatchLayerKernel.h b/src/core/NEON/kernels/NESpaceToBatchLayerKernel.h index 627724580b..44b8cbb514 100644 --- a/src/core/NEON/kernels/NESpaceToBatchLayerKernel.h +++ b/src/core/NEON/kernels/NESpaceToBatchLayerKernel.h @@ -55,8 +55,8 @@ public: /** Initialise the kernel's inputs and output. * * @param[in] input Tensor input. Supported tensor rank: 4. Data types supported: All. - * @param[in] block_shape 1-D tensor with shape [M]. Data types supported: S32 - * @param[in] paddings 2-D tensor with shape [2, M]. Data types supported: S32 + * @param[in] block_shape 1-D tensor with shape [M]. Supported M: 2. Data types supported: S32 + * @param[in] paddings 2-D tensor with shape [2, M] (First dimension is the fastest-changing dimension). Supported M: 2. Data types supported: S32 * @param[out] output Tensor output. Data types supported: same as @p input */ void configure(const ITensor *input, const ITensor *block_shape, const ITensor *paddings, ITensor *output); @@ -65,16 +65,16 @@ public: * @param[in] input Tensor input. Supported tensor rank: 4. Data types supported: All. * @param[in] block_shape_x Block shape x value. * @param[in] block_shape_y Block shape y value. - * @param[in] padding_left The left padding of the output tensor. - * @param[in] padding_right The right padding of the output tensor. + * @param[in] padding_left The padding at the beginning of every dimension of the output tensor. + * @param[in] padding_right The padding at the end of every dimension of the output tensor. * @param[out] output Tensor output. Data types supported: same as @p input */ void configure(const ITensor *input, const int block_shape_x, const int block_shape_y, const Size2D &padding_left, const Size2D &padding_right, ITensor *output); /** Static function to check if given info will lead to a valid configuration of @ref NESpaceToBatchLayerKernel * * @param[in] input Tensor input. Supported tensor rank: 4. Data types supported: All. - * @param[in] block_shape 1-D tensor with shape [M]. Data types supported: S32 - * @param[in] paddings 2-D tensor with shape [2, M]. Data types supported: S32 + * @param[in] block_shape 1-D tensor with shape [M]. Supported M: 2. Data types supported: S32 + * @param[in] paddings 2-D tensor with shape [2, M] (First dimension is the fastest-changing dimension). Supported M: 2. Data types supported: S32 * @param[in] output Tensor output. Data types supported: same as @p input * * @return a status @@ -85,8 +85,8 @@ public: * @param[in] input Tensor input. Supported tensor rank: 4. Data types supported: All. * @param[in] block_shape_x Block shape x value. * @param[in] block_shape_y Block shape y value. - * @param[in] padding_left The left padding of the output tensor. - * @param[in] padding_right The right padding of the output tensor. + * @param[in] padding_left The padding at the beginning of every dimension of the output tensor. + * @param[in] padding_right The padding at the end of every dimension of the output tensor. * @param[in] output Tensor output. Data types supported: same as @p input * * @return a status @@ -98,8 +98,8 @@ public: private: const ITensor *_input; /**< Source tensor */ - const ITensor *_block_shape; /**< Block shape tensor */ - const ITensor *_paddings; /**< Paddings tensor */ + const ITensor *_block_shape; /**< Block shape tensor for dynamic evaluation */ + const ITensor *_paddings; /**< Paddings tensor for dynamic evaluation */ ITensor *_output; /**< Destination tensor */ DataLayout _data_layout; /**< Data layout to be used at run-time */ -- cgit v1.2.1