From b57be0da77370e5e71fe82dfa281f528279e8127 Mon Sep 17 00:00:00 2001 From: Michele Di Giorgio Date: Fri, 31 Aug 2018 16:26:25 +0100 Subject: COMPMID-1330: Add support for NormalizePlanarYUV operator in CL Change-Id: Id0754b9e2bc3ef7ff2c4c21c3b89709588c41bd3 Reviewed-on: https://eu-gerrit-1.euhpc.arm.com/146637 Tested-by: Jenkins Reviewed-by: Georgios Pinitas Reviewed-by: Giorgio Arena --- arm_compute/core/CL/CLKernels.h | 1 + .../CL/kernels/CLNormalizePlanarYUVLayerKernel.h | 83 ++++++++++ .../kernels/GCNormalizePlanarYUVLayerKernel.h | 26 +++- arm_compute/runtime/CL/CLFunctions.h | 1 + .../CL/functions/CLNormalizePlanarYUVLayer.h | 75 +++++++++ .../functions/GCNormalizePlanarYUVLayer.h | 26 +++- src/core/CL/CLKernelLibrary.cpp | 6 + .../CL/cl_kernels/normalize_planar_yuv_layer.cl | 134 ++++++++++++++++ .../CL/kernels/CLNormalizePlanarYUVLayerKernel.cpp | 173 +++++++++++++++++++++ .../kernels/GCNormalizePlanarYUVLayerKernel.cpp | 93 ++++++++--- .../CL/functions/CLNormalizePlanarYUVLayer.cpp | 55 +++++++ .../functions/GCNormalizePlanarYUVLayer.cpp | 12 +- tests/datasets/NormalizePlanarYUVLayerDataset.h | 4 +- .../RandomNormalizePlanarYUVLayerDataset.h | 4 +- tests/validation/CL/NormalizePlanarYUVLayer.cpp | 142 +++++++++++++++++ .../GLES_COMPUTE/NormalizePlanarYUVLayer.cpp | 42 ++++- .../fixtures/NormalizePlanarYUVLayerFixture.h | 39 +++-- .../reference/NormalizePlanarYUVLayer.cpp | 10 +- .../validation/reference/NormalizePlanarYUVLayer.h | 4 +- 19 files changed, 860 insertions(+), 70 deletions(-) create mode 100644 arm_compute/core/CL/kernels/CLNormalizePlanarYUVLayerKernel.h create mode 100644 arm_compute/runtime/CL/functions/CLNormalizePlanarYUVLayer.h create mode 100644 src/core/CL/cl_kernels/normalize_planar_yuv_layer.cl create mode 100644 src/core/CL/kernels/CLNormalizePlanarYUVLayerKernel.cpp create mode 100644 src/runtime/CL/functions/CLNormalizePlanarYUVLayer.cpp create mode 100644 tests/validation/CL/NormalizePlanarYUVLayer.cpp diff --git a/arm_compute/core/CL/CLKernels.h b/arm_compute/core/CL/CLKernels.h index da2d3166a4..4750031603 100644 --- a/arm_compute/core/CL/CLKernels.h +++ b/arm_compute/core/CL/CLKernels.h @@ -96,6 +96,7 @@ #include "arm_compute/core/CL/kernels/CLNonLinearFilterKernel.h" #include "arm_compute/core/CL/kernels/CLNonMaximaSuppression3x3Kernel.h" #include "arm_compute/core/CL/kernels/CLNormalizationLayerKernel.h" +#include "arm_compute/core/CL/kernels/CLNormalizePlanarYUVLayerKernel.h" #include "arm_compute/core/CL/kernels/CLPermuteKernel.h" #include "arm_compute/core/CL/kernels/CLPixelWiseMultiplicationKernel.h" #include "arm_compute/core/CL/kernels/CLPoolingLayerKernel.h" diff --git a/arm_compute/core/CL/kernels/CLNormalizePlanarYUVLayerKernel.h b/arm_compute/core/CL/kernels/CLNormalizePlanarYUVLayerKernel.h new file mode 100644 index 0000000000..5418d31a0c --- /dev/null +++ b/arm_compute/core/CL/kernels/CLNormalizePlanarYUVLayerKernel.h @@ -0,0 +1,83 @@ +/* + * Copyright (c) 2018 ARM Limited. + * + * SPDX-License-Identifier: MIT + * + * Permission is hereby granted, free of charge, to any person obtaining a copy + * of this software and associated documentation files (the "Software"), to + * deal in the Software without restriction, including without limitation the + * rights to use, copy, modify, merge, publish, distribute, sublicense, and/or + * sell copies of the Software, and to permit persons to whom the Software is + * furnished to do so, subject to the following conditions: + * + * The above copyright notice and this permission notice shall be included in all + * copies or substantial portions of the Software. + * + * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR + * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, + * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE + * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER + * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, + * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE + * SOFTWARE. + */ +#ifndef __ARM_COMPUTE_CLNORMALIZEPLANARYUVLAYERKERNEL_H__ +#define __ARM_COMPUTE_CLNORMALIZEPLANARYUVLAYERKERNEL_H__ + +#include "arm_compute/core/CL/ICLKernel.h" + +namespace arm_compute +{ +class ICLTensor; + +/** Interface for the NormalizePlanarYUV layer kernel. */ +class CLNormalizePlanarYUVLayerKernel : public ICLKernel +{ +public: + /** Constructor */ + CLNormalizePlanarYUVLayerKernel(); + /** Prevent instances of this class from being copied (As this class contains pointers) */ + CLNormalizePlanarYUVLayerKernel(const CLNormalizePlanarYUVLayerKernel &) = delete; + /** Prevent instances of this class from being copied (As this class contains pointers) */ + CLNormalizePlanarYUVLayerKernel &operator=(const CLNormalizePlanarYUVLayerKernel &) = delete; + /** Default Move Constructor. */ + CLNormalizePlanarYUVLayerKernel(CLNormalizePlanarYUVLayerKernel &&) = default; + /** Default move assignment operator */ + CLNormalizePlanarYUVLayerKernel &operator=(CLNormalizePlanarYUVLayerKernel &&) = default; + /** Default destructor */ + ~CLNormalizePlanarYUVLayerKernel() = default; + + /** Set the input and output tensors. + * + * @param[in] input Source tensor. 3 lower dimensions represent a single input with dimensions [width, height, channels]. + * Data types supported: F16/F32. + * @param[out] output Destination tensor. Data type supported: same as @p input + * @param[in] mean Mean values tensor. 1 dimension with size equal to the number of input channels. Data types supported: same as @p input + * @param[in] std Standard deviation values tensor. 1 dimension with size equal to the number of input channels. + * Data types supported: same as @p input + */ + void configure(const ICLTensor *input, ICLTensor *output, const ICLTensor *mean, const ICLTensor *std); + /** Static function to check if given info will lead to a valid configuration of @ref CLNormalizePlanarYUVLayerKernel + * + * @param[in] input Source tensor info. 3 lower dimensions represent a single input with dimensions [width, height, channels]. + * Data types supported: F16/F32. + * @param[out] output Destination tensor info. Data type supported: same as @p input + * @param[in] mean Mean values tensor info. 1 dimension with size equal to the number of input channels. Data types supported: same as @p input + * @param[in] std Standard deviation values tensor info. 1 dimension with size equal to the number of input channels. + * Data types supported: same as @p input + * + * @return a status + */ + static Status validate(const ITensorInfo *input, const ITensorInfo *output, const ITensorInfo *mean, const ITensorInfo *std); + + // Inherited methods overridden: + void run(const Window &window, cl::CommandQueue &queue) override; + +private: + const ICLTensor *_input; + ICLTensor *_output; + const ICLTensor *_mean; + const ICLTensor *_std; +}; +} // namespace arm_compute +#endif /*__ARM_COMPUTE_CLNORMALIZEPLANARYUVLAYERKERNEL_H__ */ diff --git a/arm_compute/core/GLES_COMPUTE/kernels/GCNormalizePlanarYUVLayerKernel.h b/arm_compute/core/GLES_COMPUTE/kernels/GCNormalizePlanarYUVLayerKernel.h index 0d785ca0d4..7ffe5b20df 100644 --- a/arm_compute/core/GLES_COMPUTE/kernels/GCNormalizePlanarYUVLayerKernel.h +++ b/arm_compute/core/GLES_COMPUTE/kernels/GCNormalizePlanarYUVLayerKernel.h @@ -50,14 +50,26 @@ public: /** Set the input and output tensors. * - * @param[in] input Source tensor. 3 lower dimensions represent a single input with dimensions [width, height, FM]. + * @param[in] input Source tensor. 3 lower dimensions represent a single input with dimensions [width, height, channels]. * Data types supported: F16. - * @param[out] output Destination tensor. Output will have the same number of dimensions as input. Data type supported: same as @p input - * @param[in] mean Mean values tensor. 1 dimension with size equal to the feature maps [FM]. Data types supported: Same as @p input - * @param[in] sd Standard deviation values tensor. 1 dimension with size equal to the feature maps [FM]. - * Data types supported: Same as @p input + * @param[out] output Destination tensor. Data type supported: same as @p input + * @param[in] mean Mean values tensor. 1 dimension with size equal to the number of input channels. Data types supported: same as @p input + * @param[in] std Standard deviation values tensor. 1 dimension with size equal to the feature maps [FM]. + * Data types supported: same as @p input */ - void configure(const IGCTensor *input, IGCTensor *output, const IGCTensor *mean, const IGCTensor *sd); + void configure(const IGCTensor *input, IGCTensor *output, const IGCTensor *mean, const IGCTensor *std); + /** Static function to check if given info will lead to a valid configuration of @ref GCNormalizePlanarYUVLayerKernel + * + * @param[in] input Source tensor info. 3 lower dimensions represent a single input with dimensions [width, height, channels]. + * Data types supported: F16. + * @param[out] output Destination tensor info. Data type supported: same as @p input + * @param[in] mean Mean values tensor info. 1 dimension with size equal to the number of input channels. Data types supported: same as @p input + * @param[in] std Standard deviation values tensor info. 1 dimension with size equal to the number of input channels. + * Data types supported: same as @p input + * + * @return a status + */ + static Status validate(const ITensorInfo *input, const ITensorInfo *output, const ITensorInfo *mean, const ITensorInfo *std); // Inherited methods overridden: void run(const Window &window) override; @@ -66,7 +78,7 @@ private: const IGCTensor *_input; IGCTensor *_output; const IGCTensor *_mean; - const IGCTensor *_sd; + const IGCTensor *_std; }; } #endif /*__ARM_COMPUTE_GCNORMALIZEPLANARYUVLAYERKERNEL_H__ */ diff --git a/arm_compute/runtime/CL/CLFunctions.h b/arm_compute/runtime/CL/CLFunctions.h index d2bfdfd7cb..02a4dab6f1 100644 --- a/arm_compute/runtime/CL/CLFunctions.h +++ b/arm_compute/runtime/CL/CLFunctions.h @@ -93,6 +93,7 @@ #include "arm_compute/runtime/CL/functions/CLNonLinearFilter.h" #include "arm_compute/runtime/CL/functions/CLNonMaximaSuppression3x3.h" #include "arm_compute/runtime/CL/functions/CLNormalizationLayer.h" +#include "arm_compute/runtime/CL/functions/CLNormalizePlanarYUVLayer.h" #include "arm_compute/runtime/CL/functions/CLOpticalFlow.h" #include "arm_compute/runtime/CL/functions/CLPermute.h" #include "arm_compute/runtime/CL/functions/CLPhase.h" diff --git a/arm_compute/runtime/CL/functions/CLNormalizePlanarYUVLayer.h b/arm_compute/runtime/CL/functions/CLNormalizePlanarYUVLayer.h new file mode 100644 index 0000000000..85f7d93ddf --- /dev/null +++ b/arm_compute/runtime/CL/functions/CLNormalizePlanarYUVLayer.h @@ -0,0 +1,75 @@ +/* + * Copyright (c) 2018 ARM Limited. + * + * SPDX-License-Identifier: MIT + * + * Permission is hereby granted, free of charge, to any person obtaining a copy + * of this software and associated documentation files (the "Software"), to + * deal in the Software without restriction, including without limitation the + * rights to use, copy, modify, merge, publish, distribute, sublicense, and/or + * sell copies of the Software, and to permit persons to whom the Software is + * furnished to do so, subject to the following conditions: + * + * The above copyright notice and this permission notice shall be included in all + * copies or substantial portions of the Software. + * + * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR + * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, + * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE + * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER + * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, + * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE + * SOFTWARE. + */ +#ifndef __ARM_COMPUTE_CLNORMALIZEPLANARYUVLAYER_H__ +#define __ARM_COMPUTE_CLNORMALIZEPLANARYUVLAYER_H__ + +#include "arm_compute/runtime/IFunction.h" + +#include "arm_compute/core/CL/kernels/CLNormalizePlanarYUVLayerKernel.h" +#include "arm_compute/core/Types.h" + +namespace arm_compute +{ +class ICLTensor; + +/** Basic function to run @ref CLNormalizePlanarYUVLayerKernel + * + * @note The function simulates a NormalizePlanarYUV layer. + */ +class CLNormalizePlanarYUVLayer : public IFunction +{ +public: + /** Default constructor */ + CLNormalizePlanarYUVLayer(); + /** Set the input and output tensors. + * + * @param[in] input Source tensor. 3 lower dimensions represent a single input with dimensions [width, height, channels]. + * Data types supported: F16/F32. + * @param[out] output Destinationfeature tensor. Data type supported: same as @p input + * @param[in] mean Mean values tensor. 1 dimension with size equal to the number of input channels. Data types supported: Same as @p input + * @param[in] std Standard deviation values tensor. 1 dimension with size equal to the number of input channels. + * Data types supported: Same as @p input + */ + void configure(const ICLTensor *input, ICLTensor *output, const ICLTensor *mean, const ICLTensor *std); + /** Static function to check if given info will lead to a valid configuration of @ref CLNormalizePlanarYUVLayer + * + * @param[in] input Source tensor info. 3 lower dimensions represent a single input with dimensions [width, height, FM]. + * Data types supported: F16/F32. + * @param[out] output Destination tensor info. Data type supported: same as @p input + * @param[in] mean Mean values tensor info. 1 dimension with size equal to the number of input channels. Data types supported: Same as @p input + * @param[in] std Standard deviation values tensor info. 1 dimension with size equal to the number of input channels. + * Data types supported: Same as @p input + * + * @return a status + */ + static Status validate(const ITensorInfo *input, const ITensorInfo *output, const ITensorInfo *mean, const ITensorInfo *std); + + // Inherited methods overridden: + void run() override; + +private: + CLNormalizePlanarYUVLayerKernel _norm_kernel; /**< NormalizePlanarYUV layer kernel to run */ +}; +} // namespace arm_compute +#endif /* __ARM_COMPUTE_CLNORMALIZEPLANARYUVLAYER_H__ */ diff --git a/arm_compute/runtime/GLES_COMPUTE/functions/GCNormalizePlanarYUVLayer.h b/arm_compute/runtime/GLES_COMPUTE/functions/GCNormalizePlanarYUVLayer.h index 2862eeb9cd..d6cf4d0803 100644 --- a/arm_compute/runtime/GLES_COMPUTE/functions/GCNormalizePlanarYUVLayer.h +++ b/arm_compute/runtime/GLES_COMPUTE/functions/GCNormalizePlanarYUVLayer.h @@ -1,5 +1,5 @@ /* - * Copyright (c) 2017 ARM Limited. + * Copyright (c) 2017-2018 ARM Limited. * * SPDX-License-Identifier: MIT * @@ -44,14 +44,26 @@ public: GCNormalizePlanarYUVLayer(); /** Set the input and output tensors. * - * @param[in] input Source tensor. 3 lower dimensions represent a single input with dimensions [width, height, FM]. + * @param[in] input Source tensor. 3 lower dimensions represent a single input with dimensions [width, height, channels]. * Data types supported: F16. - * @param[out] output Destination tensor. Output will have the same number of dimensions as input. Data type supported: same as @p input - * @param[in] mean Mean values tensor. 1 dimension with size equal to the feature maps [FM]. Data types supported: Same as @p input - * @param[in] sd Standard deviation values tensor. 1 dimension with size equal to the feature maps [FM]. - * Data types supported: Same as @p input + * @param[out] output Destination tensor. Data type supported: same as @p input + * @param[in] mean Mean values tensor. 1 dimension with size equal to the number of input channels. Data types supported: same as @p input + * @param[in] std Standard deviation values tensor. 1 dimension with size equal to the number of input channels. + * Data types supported: same as @p input */ - void configure(const IGCTensor *input, IGCTensor *output, const IGCTensor *mean, const IGCTensor *sd); + void configure(const IGCTensor *input, IGCTensor *output, const IGCTensor *mean, const IGCTensor *std); + /** Static function to check if given info will lead to a valid configuration of @ref CLNormalizePlanarYUVLayer + * + * @param[in] input Source tensor info. 3 lower dimensions represent a single input with dimensions [width, height, channels]. + * Data types supported: F16/F32. + * @param[out] output Destination tensor info. Data type supported: same as @p input + * @param[in] mean Mean values tensor info. 1 dimension with size equal to the number of input channels. Data types supported: same as @p input + * @param[in] std Standard deviation values tensor info. 1 dimension with size equal to the number of input channels. + * Data types supported: same as @p input + * + * @return a status + */ + static Status validate(const ITensorInfo *input, const ITensorInfo *output, const ITensorInfo *mean, const ITensorInfo *std); // Inherited methods overridden: void run() override; diff --git a/src/core/CL/CLKernelLibrary.cpp b/src/core/CL/CLKernelLibrary.cpp index 0cc6e320bf..4af2b09530 100644 --- a/src/core/CL/CLKernelLibrary.cpp +++ b/src/core/CL/CLKernelLibrary.cpp @@ -308,6 +308,8 @@ const std::map CLKernelLibrary::_kernel_program_map = { "non_max_suppression", "nonmax.cl" }, { "normalization_layer_cross_map", "normalization_layer.cl" }, { "normalization_layer_in_map", "normalization_layer.cl" }, + { "normalize_planar_yuv_layer_nchw", "normalize_planar_yuv_layer.cl" }, + { "normalize_planar_yuv_layer_nhwc", "normalize_planar_yuv_layer.cl" }, { "NV12_to_IYUV_bt709", "color_convert.cl" }, { "NV12_to_RGB888_bt709", "color_convert.cl" }, { "NV12_to_RGBA8888_bt709", "color_convert.cl" }, @@ -672,6 +674,10 @@ const std::map CLKernelLibrary::_program_source_map = { "normalization_layer.cl", #include "./cl_kernels/normalization_layer.clembed" + }, + { + "normalize_planar_yuv_layer.cl", +#include "./cl_kernels/normalize_planar_yuv_layer.clembed" }, { "batchnormalization_layer.cl", diff --git a/src/core/CL/cl_kernels/normalize_planar_yuv_layer.cl b/src/core/CL/cl_kernels/normalize_planar_yuv_layer.cl new file mode 100644 index 0000000000..dc6652449e --- /dev/null +++ b/src/core/CL/cl_kernels/normalize_planar_yuv_layer.cl @@ -0,0 +1,134 @@ +/* + * Copyright (c) 2018 ARM Limited. + * + * SPDX-License-Identifier: MIT + * + * Permission is hereby granted, free of charge, to any person obtaining a copy + * of this software and associated documentation files (the "Software"), to + * deal in the Software without restriction, including without limitation the + * rights to use, copy, modify, merge, publish, distribute, sublicense, and/or + * sell copies of the Software, and to permit persons to whom the Software is + * furnished to do so, subject to the following conditions: + * + * The above copyright notice and this permission notice shall be included in all + * copies or substantial portions of the Software. + * + * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR + * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, + * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE + * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER + * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, + * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE + * SOFTWARE. + */ +#include "helpers.h" + +#if defined(DATA_TYPE) && defined(VEC_SIZE) + +#define TYPE VEC_DATA_TYPE(DATA_TYPE, VEC_SIZE) + +/** Apply normalize_planar_yuv layer on tensors with NCHW format. + * + * @note Data type should be given as a preprocessor argument using -DDATA_TYPE=type. e.g. -DDATA_TYPE=float + * @note Vector size should be given as a preprocessor argument using -DVEC_SIZE e.g. -DVEC_SIZE=8 + * @note The depth of the input tensor should be given as a preprocessor argument using -DNUM_CHANNELS e.g. -DNUM_CHANNELS=8 + * + * @param[in] src_ptr Pointer to the first source tensor. Supported data types: F16/F32 + * @param[in] src_stride_x Stride of the first source tensor in X dimension (in bytes) + * @param[in] src_step_x input_stride_x * number of elements along X processed per workitem(in bytes) + * @param[in] src_stride_y Stride of the first source tensor in Y dimension (in bytes) + * @param[in] src_step_y input_stride_y * number of elements along Y processed per workitem(in bytes) + * @param[in] src_stride_z Stride of the first source tensor in Z dimension (in bytes) + * @param[in] src_step_z input_stride_z * number of elements along Z processed per workitem(in bytes) + * @param[in] src_offset_first_element_in_bytes The offset of the first element in the first source tensor + * @param[out] dst_ptr Pointer to the destination tensor. Supported data types: same as @p src_ptr + * @param[in] dst_stride_x Stride of the destination tensor in X dimension (in bytes) + * @param[in] dst_step_x output_stride_x * number of elements along X processed per workitem(in bytes) + * @param[in] dst_stride_y Stride of the destination tensor in Y dimension (in bytes) + * @param[in] dst_step_y output_stride_y * number of elements along Y processed per workitem(in bytes) + * @param[in] dst_stride_z Stride of the destination tensor in Z dimension (in bytes) + * @param[in] dst_step_z output_stride_z * number of elements along Z processed per workitem(in bytes) + * @param[in] dst_offset_first_element_in_bytes The offset of the first element in the destination tensor + * @param[in] mean_ptr Pointer to the mean source tensor. Supported data types: same as @p src_ptr + * @param[in] mean_stride_x Stride of the mean source tensor in X dimension (in bytes) + * @param[in] mean_step_x mean_stride_x * number of elements along X processed per workitem(in bytes) + * @param[in] mean_offset_first_element_in_bytes The offset of the first element in the mean source tensor + * @param[in] std_ptr Pointer to the std tensor. Supported data types: same as @p src_ptr + * @param[in] std_stride_x Stride of the std tensor in X dimension (in bytes) + * @param[in] std_step_x std_stride_x * number of elements along X processed per workitem(in bytes) + * @param[in] std_offset_first_element_in_bytes The offset of the first element in the var source tensor + */ +__kernel void normalize_planar_yuv_layer_nchw(TENSOR3D_DECLARATION(src), + TENSOR3D_DECLARATION(dst), + VECTOR_DECLARATION(mean), + VECTOR_DECLARATION(std)) +{ + Tensor3D src = CONVERT_TO_TENSOR3D_STRUCT(src); + Tensor3D dst = CONVERT_TO_TENSOR3D_STRUCT(dst); + Vector mean = CONVERT_TO_VECTOR_STRUCT(mean); + Vector std = CONVERT_TO_VECTOR_STRUCT(std); + + const uint current_slice = get_global_id(2) % NUM_CHANNELS; + + const DATA_TYPE curr_mean = *((__global DATA_TYPE *)(mean.ptr + current_slice * mean.stride_x)); + const DATA_TYPE curr_std = *((__global DATA_TYPE *)(std.ptr + current_slice * std.stride_x)); + + TYPE data = VLOAD(VEC_SIZE)(0, (__global DATA_TYPE *)src.ptr); + TYPE res = (data - curr_mean) / curr_std; + + VSTORE(VEC_SIZE) + (res, 0, (__global DATA_TYPE *)dst.ptr); +} + +/** Apply normalize_planar_yuv layer on tensors with NHWC format. + * + * @note Data type should be given as a preprocessor argument using -DDATA_TYPE=type. e.g. -DDATA_TYPE=float + * @note Vector size should be given as a preprocessor argument using -DVEC_SIZE e.g. -DVEC_SIZE=8 + * + * @param[in] src_ptr Pointer to the first source tensor. Supported data types: F16/F32 + * @param[in] src_stride_x Stride of the first source tensor in X dimension (in bytes) + * @param[in] src_step_x input_stride_x * number of elements along X processed per workitem(in bytes) + * @param[in] src_stride_y Stride of the first source tensor in Y dimension (in bytes) + * @param[in] src_step_y input_stride_y * number of elements along Y processed per workitem(in bytes) + * @param[in] src_stride_z Stride of the first source tensor in Z dimension (in bytes) + * @param[in] src_step_z input_stride_z * number of elements along Z processed per workitem(in bytes) + * @param[in] src_offset_first_element_in_bytes The offset of the first element in the first source tensor + * @param[out] dst_ptr Pointer to the destination tensor. Supported data types: same as @p src_ptr + * @param[in] dst_stride_x Stride of the destination tensor in X dimension (in bytes) + * @param[in] dst_step_x output_stride_x * number of elements along X processed per workitem(in bytes) + * @param[in] dst_stride_y Stride of the destination tensor in Y dimension (in bytes) + * @param[in] dst_step_y output_stride_y * number of elements along Y processed per workitem(in bytes) + * @param[in] dst_stride_z Stride of the destination tensor in Z dimension (in bytes) + * @param[in] dst_step_z output_stride_z * number of elements along Z processed per workitem(in bytes) + * @param[in] dst_offset_first_element_in_bytes The offset of the first element in the destination tensor + * @param[in] mean_ptr Pointer to the mean source tensor. Supported data types: same as @p src_ptr + * @param[in] mean_stride_x Stride of the mean source tensor in X dimension (in bytes) + * @param[in] mean_step_x mean_stride_x * number of elements along X processed per workitem(in bytes) + * @param[in] mean_offset_first_element_in_bytes The offset of the first element in the mean source tensor + * @param[in] std_ptr Pointer to the std tensor. Supported data types: same as @p src_ptr + * @param[in] std_stride_x Stride of the std tensor in X dimension (in bytes) + * @param[in] std_step_x std_stride_x * number of elements along X processed per workitem(in bytes) + * @param[in] std_offset_first_element_in_bytes The offset of the first element in the var source tensor + */ +__kernel void normalize_planar_yuv_layer_nhwc(TENSOR3D_DECLARATION(src), + TENSOR3D_DECLARATION(dst), + VECTOR_DECLARATION(mean), + VECTOR_DECLARATION(std)) +{ + Tensor3D src = CONVERT_TO_TENSOR3D_STRUCT(src); + Tensor3D dst = CONVERT_TO_TENSOR3D_STRUCT(dst); + Vector mean = CONVERT_TO_VECTOR_STRUCT(mean); + Vector std = CONVERT_TO_VECTOR_STRUCT(std); + + const uint current_slice = get_global_id(0); + + const TYPE curr_mean = VLOAD(VEC_SIZE)(0, (__global DATA_TYPE *)(mean.ptr + current_slice * VEC_SIZE * mean.stride_x)); + const TYPE curr_std = VLOAD(VEC_SIZE)(0, (__global DATA_TYPE *)(std.ptr + current_slice * VEC_SIZE * std.stride_x)); + + TYPE data = VLOAD(VEC_SIZE)(0, (__global DATA_TYPE *)src.ptr); + TYPE res = (data - curr_mean) / curr_std; + + VSTORE(VEC_SIZE) + (res, 0, (__global DATA_TYPE *)dst.ptr); +} +#endif // defined(DATA_TYPE) && defined(VEC_SIZE) diff --git a/src/core/CL/kernels/CLNormalizePlanarYUVLayerKernel.cpp b/src/core/CL/kernels/CLNormalizePlanarYUVLayerKernel.cpp new file mode 100644 index 0000000000..31451ef422 --- /dev/null +++ b/src/core/CL/kernels/CLNormalizePlanarYUVLayerKernel.cpp @@ -0,0 +1,173 @@ +/* + * Copyright (c) 2018 ARM Limited. + * + * SPDX-License-Identifier: MIT + * + * Permission is hereby granted, free of charge, to any person obtaining a copy + * of this software and associated documentation files (the "Software"), to + * deal in the Software without restriction, including without limitation the + * rights to use, copy, modify, merge, publish, distribute, sublicense, and/or + * sell copies of the Software, and to permit persons to whom the Software is + * furnished to do so, subject to the following conditions: + * + * The above copyright notice and this permission notice shall be included in all + * copies or substantial portions of the Software. + * + * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR + * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, + * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE + * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER + * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, + * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE + * SOFTWARE. + */ +#include "arm_compute/core/CL/kernels/CLNormalizePlanarYUVLayerKernel.h" + +#include "arm_compute/core/AccessWindowStatic.h" +#include "arm_compute/core/CL/CLHelpers.h" +#include "arm_compute/core/CL/CLKernelLibrary.h" +#include "arm_compute/core/CL/CLValidate.h" +#include "arm_compute/core/CL/ICLTensor.h" +#include "arm_compute/core/Helpers.h" +#include "arm_compute/core/TensorInfo.h" +#include "arm_compute/core/Utils.h" +#include "arm_compute/core/Window.h" + +#include "support/ToolchainSupport.h" + +using namespace arm_compute; + +namespace +{ +Status validate_arguments(const ITensorInfo *input, const ITensorInfo *output, const ITensorInfo *mean, const ITensorInfo *std) +{ + ARM_COMPUTE_RETURN_ERROR_ON_F16_UNSUPPORTED(input); + ARM_COMPUTE_RETURN_ERROR_ON_DATA_TYPE_CHANNEL_NOT_IN(input, 1, DataType::F16, DataType::F32); + ARM_COMPUTE_RETURN_ERROR_ON_NULLPTR(output); + + ARM_COMPUTE_RETURN_ERROR_ON_MISMATCHING_DATA_TYPES(input, mean, std); + ARM_COMPUTE_RETURN_ERROR_ON_MISMATCHING_SHAPES(mean, std); + ARM_COMPUTE_RETURN_ERROR_ON_MSG(mean->num_dimensions() > 1, "mean and std must be vectors"); + + const unsigned int channel_idx = get_data_layout_dimension_index(input->data_layout(), DataLayoutDimension::CHANNEL); + ARM_COMPUTE_RETURN_ERROR_ON(input->dimension(channel_idx) != mean->dimension(0)); + + // Checks performed when output is configured + if(output->total_size() != 0) + { + ARM_COMPUTE_RETURN_ERROR_ON_MISMATCHING_DATA_TYPES(input, output); + ARM_COMPUTE_RETURN_ERROR_ON_MISMATCHING_SHAPES(input, output); + } + + return Status{}; +} + +std::pair validate_and_configure_window(ITensorInfo *input, ITensorInfo *output, ITensorInfo *mean, ITensorInfo *std) +{ + // Output tensor auto initialization if not yet initialized + auto_init_if_empty(*output, *input->clone()); + + const unsigned int num_elems_processed_per_iteration = 16 / input->element_size(); + + Window win = calculate_max_window(*input, Steps(num_elems_processed_per_iteration)); + + AccessWindowHorizontal input_access(input, 0, num_elems_processed_per_iteration); + AccessWindowHorizontal output_access(output, 0, num_elems_processed_per_iteration); + + bool window_changed = update_window_and_padding(win, input_access, output_access); + output_access.set_valid_region(win, input->valid_region()); + + if(input->data_layout() == DataLayout::NHWC) + { + AccessWindowHorizontal mean_access(mean, 0, num_elems_processed_per_iteration); + AccessWindowHorizontal std_access(std, 0, num_elems_processed_per_iteration); + window_changed = window_changed || update_window_and_padding(win, mean_access, std_access); + } + + Status err = (window_changed) ? ARM_COMPUTE_CREATE_ERROR(ErrorCode::RUNTIME_ERROR, "Insufficient Padding!") : Status{}; + return std::make_pair(err, win); +} +} // namespace + +CLNormalizePlanarYUVLayerKernel::CLNormalizePlanarYUVLayerKernel() + : _input(nullptr), _output(nullptr), _mean(nullptr), _std(nullptr) +{ +} + +void CLNormalizePlanarYUVLayerKernel::configure(const ICLTensor *input, ICLTensor *output, const ICLTensor *mean, const ICLTensor *std) +{ + ARM_COMPUTE_ERROR_ON_NULLPTR(input, output, mean, std); + + // Output tensor auto initialization if not yet initialized + auto_init_if_empty(*output->info(), *input->info()->clone()); + + // Perform validation step + ARM_COMPUTE_ERROR_THROW_ON(validate_arguments(input->info(), output->info(), mean->info(), std->info())); + + _input = input; + _output = output; + _mean = mean; + _std = std; + + const unsigned int num_elems_processed_per_iteration = 16 / input->info()->element_size(); + const unsigned int channel_idx = get_data_layout_dimension_index(input->info()->data_layout(), DataLayoutDimension::CHANNEL); + + // Set build options + CLBuildOptions build_opts; + build_opts.add_option(("-DDATA_TYPE=" + get_cl_type_from_data_type(input->info()->data_type()))); + build_opts.add_option(("-DVEC_SIZE=" + support::cpp11::to_string(num_elems_processed_per_iteration))); + build_opts.add_option(("-DNUM_CHANNELS=" + support::cpp11::to_string(input->info()->dimension(channel_idx)))); + + // Create kernel + _kernel = static_cast(CLKernelLibrary::get().create_kernel("normalize_planar_yuv_layer_" + lower_string(string_from_data_layout(input->info()->data_layout())), build_opts.options())); + + // Configure kernel window + auto win_config = validate_and_configure_window(input->info(), output->info(), mean->info(), std->info()); + ARM_COMPUTE_ERROR_THROW_ON(win_config.first); + ICLKernel::configure_internal(win_config.second); + + // Set config_id for enabling LWS tuning + _config_id = "normalize_planar_yuv_layer_"; + _config_id += lower_string(string_from_data_layout(input->info()->data_layout())); + _config_id += "_"; + _config_id += lower_string(string_from_data_type(input->info()->data_type())); + _config_id += "_"; + _config_id += support::cpp11::to_string(input->info()->dimension(0)); + _config_id += "_"; + _config_id += support::cpp11::to_string(input->info()->dimension(1)); + _config_id += "_"; + _config_id += support::cpp11::to_string(input->info()->dimension(2)); +} + +Status CLNormalizePlanarYUVLayerKernel::validate(const ITensorInfo *input, const ITensorInfo *output, const ITensorInfo *mean, const ITensorInfo *std) +{ + ARM_COMPUTE_RETURN_ON_ERROR(validate_arguments(input, output, mean, std)); + ARM_COMPUTE_RETURN_ON_ERROR(validate_and_configure_window(input->clone().get(), output->clone().get(), mean->clone().get(), std->clone().get()).first); + + return Status{}; +} + +void CLNormalizePlanarYUVLayerKernel::run(const Window &window, cl::CommandQueue &queue) +{ + ARM_COMPUTE_ERROR_ON_UNCONFIGURED_KERNEL(this); + ARM_COMPUTE_ERROR_ON_INVALID_SUBWINDOW(IKernel::window(), window); + + Window collapsed = window.collapse_if_possible(ICLKernel::window(), Window::DimZ); + Window slice = collapsed.first_slice_window_3D(); + + Window slice_in = collapsed.first_slice_window_1D(); + slice_in.set(Window::DimX, Window::Dimension(0, 0, 0)); + + unsigned int idx = 2 * num_arguments_per_3D_tensor(); + add_1D_tensor_argument(idx, _mean, slice_in); + add_1D_tensor_argument(idx, _std, slice_in); + + do + { + idx = 0; + add_3D_tensor_argument(idx, _input, slice); + add_3D_tensor_argument(idx, _output, slice); + enqueue(queue, *this, slice, lws_hint()); + } + while(collapsed.slide_window_slice_3D(slice)); +} diff --git a/src/core/GLES_COMPUTE/kernels/GCNormalizePlanarYUVLayerKernel.cpp b/src/core/GLES_COMPUTE/kernels/GCNormalizePlanarYUVLayerKernel.cpp index fac29024e3..03463b2552 100644 --- a/src/core/GLES_COMPUTE/kernels/GCNormalizePlanarYUVLayerKernel.cpp +++ b/src/core/GLES_COMPUTE/kernels/GCNormalizePlanarYUVLayerKernel.cpp @@ -36,26 +36,75 @@ using namespace arm_compute; +namespace +{ +Status validate_arguments(const ITensorInfo *input, const ITensorInfo *output, const ITensorInfo *mean, const ITensorInfo *std) +{ + ARM_COMPUTE_RETURN_ERROR_ON_DATA_TYPE_CHANNEL_NOT_IN(input, 1, DataType::F16); + ARM_COMPUTE_RETURN_ERROR_ON_NULLPTR(output); + ARM_COMPUTE_RETURN_ERROR_ON(input->data_layout() != DataLayout::NCHW); + + ARM_COMPUTE_RETURN_ERROR_ON_MISMATCHING_DATA_TYPES(input, mean, std); + ARM_COMPUTE_RETURN_ERROR_ON_MISMATCHING_SHAPES(mean, std); + ARM_COMPUTE_RETURN_ERROR_ON_MSG(mean->num_dimensions() > 1, "mean and std must be vectors"); + + const unsigned int channel_idx = get_data_layout_dimension_index(input->data_layout(), DataLayoutDimension::CHANNEL); + ARM_COMPUTE_RETURN_ERROR_ON(input->dimension(channel_idx) != mean->dimension(0)); + + // Checks performed when output is configured + if(output->total_size() != 0) + { + ARM_COMPUTE_RETURN_ERROR_ON_MISMATCHING_DATA_TYPES(input, output); + ARM_COMPUTE_RETURN_ERROR_ON_MISMATCHING_SHAPES(input, output); + } + + return Status{}; +} + +std::pair validate_and_configure_window(ITensorInfo *input, ITensorInfo *output, ITensorInfo *mean, ITensorInfo *std) +{ + // Output tensor auto initialization if not yet initialized + auto_init_if_empty(*output, *input->clone()); + + const unsigned int num_elems_processed_per_iteration = 4; + + // Configure kernel window + Window win = calculate_max_window(*input, Steps(num_elems_processed_per_iteration)); + + AccessWindowHorizontal input_access(input, 0, num_elems_processed_per_iteration); + AccessWindowHorizontal output_access(output, 0, num_elems_processed_per_iteration); + const int mean_padding = ceil_to_multiple(mean->dimension(0), num_elems_processed_per_iteration) - mean->dimension(0); + const int std_padding = ceil_to_multiple(std->dimension(0), num_elems_processed_per_iteration) - std->dimension(0); + AccessWindowStatic mean_access(mean, 0, 0, mean->dimension(0) + mean_padding, mean->dimension(1)); + AccessWindowStatic std_access(std, 0, 0, std->dimension(0) + std_padding, std->dimension(1)); + + const bool window_changed = update_window_and_padding(win, input_access, output_access, mean_access, std_access); + output_access.set_valid_region(win, input->valid_region()); + + Status err = (window_changed) ? ARM_COMPUTE_CREATE_ERROR(ErrorCode::RUNTIME_ERROR, "Insufficient Padding!") : Status{}; + return std::make_pair(err, win); +} +} // namespace + GCNormalizePlanarYUVLayerKernel::GCNormalizePlanarYUVLayerKernel() - : _input(nullptr), _output(nullptr), _mean(nullptr), _sd(nullptr) + : _input(nullptr), _output(nullptr), _mean(nullptr), _std(nullptr) { } -void GCNormalizePlanarYUVLayerKernel::configure(const IGCTensor *input, IGCTensor *output, const IGCTensor *mean, const IGCTensor *sd) +void GCNormalizePlanarYUVLayerKernel::configure(const IGCTensor *input, IGCTensor *output, const IGCTensor *mean, const IGCTensor *std) { - ARM_COMPUTE_ERROR_ON_DATA_TYPE_CHANNEL_NOT_IN(input, 1, DataType::F16); - ARM_COMPUTE_ERROR_ON_NULLPTR(output); - ARM_COMPUTE_ERROR_ON_MISMATCHING_DATA_TYPES(input, output, mean, sd); - ARM_COMPUTE_ERROR_ON_MISMATCHING_SHAPES(input, output); - ARM_COMPUTE_ERROR_ON_MISMATCHING_SHAPES(mean, sd); - ARM_COMPUTE_ERROR_ON(input->info()->dimension(2) != mean->info()->dimension(0)); + ARM_COMPUTE_ERROR_ON_NULLPTR(input, output, mean, std); + + // Output tensor auto initialization if not yet initialized + auto_init_if_empty(*output->info(), *input->info()->clone()); + + // Perform validation step + ARM_COMPUTE_ERROR_THROW_ON(validate_arguments(input->info(), output->info(), mean->info(), std->info())); _input = input; _output = output; _mean = mean; - _sd = sd; - - const unsigned int num_elems_processed_per_iteration = 4; + _std = std; // Set build options std::set build_opts; @@ -67,19 +116,17 @@ void GCNormalizePlanarYUVLayerKernel::configure(const IGCTensor *input, IGCTenso _kernel = static_cast(GCKernelLibrary::get().create_kernel("normalize_planar_yuv_layer", build_opts)); // Configure kernel window - Window win = calculate_max_window(*input->info(), Steps(num_elems_processed_per_iteration)); + auto win_config = validate_and_configure_window(input->info(), output->info(), mean->info(), std->info()); + ARM_COMPUTE_ERROR_THROW_ON(std::get<0>(win_config)); - AccessWindowHorizontal input_access(input->info(), 0, num_elems_processed_per_iteration); - AccessWindowHorizontal output_access(output->info(), 0, num_elems_processed_per_iteration); - const int mean_padding = ceil_to_multiple(mean->info()->dimension(0), num_elems_processed_per_iteration) - mean->info()->dimension(0); - const int sd_padding = ceil_to_multiple(sd->info()->dimension(0), num_elems_processed_per_iteration) - sd->info()->dimension(0); - AccessWindowStatic mean_access(mean->info(), 0, 0, mean->info()->dimension(0) + mean_padding, mean->info()->dimension(1)); - AccessWindowStatic sd_access(sd->info(), 0, 0, sd->info()->dimension(0) + sd_padding, sd->info()->dimension(1)); - - update_window_and_padding(win, input_access, output_access, mean_access, sd_access); - output_access.set_valid_region(win, input->info()->valid_region()); + IGCKernel::configure(std::get<1>(win_config)); +} - IGCKernel::configure(win); +Status GCNormalizePlanarYUVLayerKernel::validate(const ITensorInfo *input, const ITensorInfo *output, const ITensorInfo *mean, const ITensorInfo *std) +{ + ARM_COMPUTE_RETURN_ON_ERROR(validate_arguments(input, output, mean, std)); + ARM_COMPUTE_RETURN_ON_ERROR(std::get<0>(validate_and_configure_window(input->clone().get(), output->clone().get(), mean->clone().get(), std->clone().get()))); + return Status{}; } void GCNormalizePlanarYUVLayerKernel::run(const Window &window) @@ -100,7 +147,7 @@ void GCNormalizePlanarYUVLayerKernel::run(const Window &window) unsigned int idx = 2 * num_arguments_per_3D_tensor(); add_1D_tensor_argument(idx, _mean, 3, slice_in); - add_1D_tensor_argument(idx, _sd, 4, slice_in); + add_1D_tensor_argument(idx, _std, 4, slice_in); slice_in = window.first_slice_window_3D(); diff --git a/src/runtime/CL/functions/CLNormalizePlanarYUVLayer.cpp b/src/runtime/CL/functions/CLNormalizePlanarYUVLayer.cpp new file mode 100644 index 0000000000..11d70e31fb --- /dev/null +++ b/src/runtime/CL/functions/CLNormalizePlanarYUVLayer.cpp @@ -0,0 +1,55 @@ +/* + * Copyright (c) 2018 ARM Limited. + * + * SPDX-License-Identifier: MIT + * + * Permission is hereby granted, free of charge, to any person obtaining a copy + * of this software and associated documentation files (the "Software"), to + * deal in the Software without restriction, including without limitation the + * rights to use, copy, modify, merge, publish, distribute, sublicense, and/or + * sell copies of the Software, and to permit persons to whom the Software is + * furnished to do so, subject to the following conditions: + * + * The above copyright notice and this permission notice shall be included in all + * copies or substantial portions of the Software. + * + * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR + * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, + * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE + * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER + * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, + * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE + * SOFTWARE. + */ + +#include "arm_compute/runtime/CL/functions/CLNormalizePlanarYUVLayer.h" + +#include "arm_compute/core/Error.h" +#include "arm_compute/core/TensorInfo.h" +#include "arm_compute/core/Types.h" +#include "arm_compute/core/Validate.h" +#include "arm_compute/runtime/CL/CLScheduler.h" + +namespace arm_compute +{ +CLNormalizePlanarYUVLayer::CLNormalizePlanarYUVLayer() + : _norm_kernel() +{ +} + +void CLNormalizePlanarYUVLayer::configure(const ICLTensor *input, ICLTensor *output, const ICLTensor *mean, const ICLTensor *std) +{ + _norm_kernel.configure(input, output, mean, std); +} + +Status CLNormalizePlanarYUVLayer::validate(const ITensorInfo *input, const ITensorInfo *output, + const ITensorInfo *mean, const ITensorInfo *std) +{ + return CLNormalizePlanarYUVLayerKernel::validate(input, output, mean, std); +} + +void CLNormalizePlanarYUVLayer::run() +{ + CLScheduler::get().enqueue(_norm_kernel, true); +} +} // namespace arm_compute diff --git a/src/runtime/GLES_COMPUTE/functions/GCNormalizePlanarYUVLayer.cpp b/src/runtime/GLES_COMPUTE/functions/GCNormalizePlanarYUVLayer.cpp index 5fb971c154..19fdc3d7c0 100755 --- a/src/runtime/GLES_COMPUTE/functions/GCNormalizePlanarYUVLayer.cpp +++ b/src/runtime/GLES_COMPUTE/functions/GCNormalizePlanarYUVLayer.cpp @@ -1,5 +1,5 @@ /* - * Copyright (c) 2017 ARM Limited. + * Copyright (c) 2017-2018 ARM Limited. * * SPDX-License-Identifier: MIT * @@ -37,9 +37,15 @@ GCNormalizePlanarYUVLayer::GCNormalizePlanarYUVLayer() { } -void GCNormalizePlanarYUVLayer::configure(const IGCTensor *input, IGCTensor *output, const IGCTensor *mean, const IGCTensor *sd) +void GCNormalizePlanarYUVLayer::configure(const IGCTensor *input, IGCTensor *output, const IGCTensor *mean, const IGCTensor *std) { - _norm_kernel.configure(input, output, mean, sd); + _norm_kernel.configure(input, output, mean, std); +} + +Status GCNormalizePlanarYUVLayer::validate(const ITensorInfo *input, const ITensorInfo *output, + const ITensorInfo *mean, const ITensorInfo *std) +{ + return GCNormalizePlanarYUVLayerKernel::validate(input, output, mean, std); } void GCNormalizePlanarYUVLayer::run() diff --git a/tests/datasets/NormalizePlanarYUVLayerDataset.h b/tests/datasets/NormalizePlanarYUVLayerDataset.h index 2d71a56a30..1a97e68a92 100644 --- a/tests/datasets/NormalizePlanarYUVLayerDataset.h +++ b/tests/datasets/NormalizePlanarYUVLayerDataset.h @@ -1,5 +1,5 @@ /* - * Copyright (c) 2017 ARM Limited. + * Copyright (c) 2017-2018 ARM Limited. * * SPDX-License-Identifier: MIT * @@ -55,7 +55,7 @@ public: description << "In=" << *_tensor_it << ":"; description << "Out=" << *_tensor_it << ":"; description << "Mean=" << *_param_it << ":"; - description << "Sd=" << *_param_it << ":"; + description << "Std=" << *_param_it << ":"; return description.str(); } diff --git a/tests/datasets/RandomNormalizePlanarYUVLayerDataset.h b/tests/datasets/RandomNormalizePlanarYUVLayerDataset.h index 5693004070..56eb604cca 100644 --- a/tests/datasets/RandomNormalizePlanarYUVLayerDataset.h +++ b/tests/datasets/RandomNormalizePlanarYUVLayerDataset.h @@ -1,5 +1,5 @@ /* - * Copyright (c) 2017 ARM Limited. + * Copyright (c) 2017-2018 ARM Limited. * * SPDX-License-Identifier: MIT * @@ -46,6 +46,8 @@ public: add_config(TensorShape(21U, 11U, 12U, 1U), TensorShape(12U)); add_config(TensorShape(7U, 3U, 6U, 1U), TensorShape(6U)); add_config(TensorShape(7U, 2U, 3U, 1U), TensorShape(3U)); + add_config(TensorShape(7U, 2U, 3U, 3U), TensorShape(3U)); + add_config(TensorShape(21U, 11U, 12U, 3U), TensorShape(12U)); } }; } // namespace datasets diff --git a/tests/validation/CL/NormalizePlanarYUVLayer.cpp b/tests/validation/CL/NormalizePlanarYUVLayer.cpp new file mode 100644 index 0000000000..aa1a00e106 --- /dev/null +++ b/tests/validation/CL/NormalizePlanarYUVLayer.cpp @@ -0,0 +1,142 @@ +/* + * Copyright (c) 2018 ARM Limited. + * + * SPDX-License-Identifier: MIT + * + * Permission is hereby granted, free of charge, to any person obtaining a copy + * of this software and associated documentation files (the "Software"), to + * deal in the Software without restriction, including without limitation the + * rights to use, copy, modify, merge, publish, distribute, sublicense, and/or + * sell copies of the Software, and to permit persons to whom the Software is + * furnished to do so, subject to the following conditions: + * + * The above copyright notice and this permission notice shall be included in all + * copies or substantial portions of the Software. + * + * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR + * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, + * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE + * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER + * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, + * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE + * SOFTWARE. + */ +#include "arm_compute/core/Types.h" +#include "arm_compute/runtime/CL/CLTensor.h" +#include "arm_compute/runtime/CL/CLTensorAllocator.h" +#include "arm_compute/runtime/CL/functions/CLNormalizePlanarYUVLayer.h" +#include "tests/CL/CLAccessor.h" +#include "tests/PaddingCalculator.h" +#include "tests/datasets/RandomNormalizePlanarYUVLayerDataset.h" +#include "tests/datasets/ShapeDatasets.h" +#include "tests/framework/Asserts.h" +#include "tests/framework/Macros.h" +#include "tests/framework/datasets/Datasets.h" +#include "tests/validation/Validation.h" +#include "tests/validation/fixtures/NormalizePlanarYUVLayerFixture.h" + +namespace arm_compute +{ +namespace test +{ +namespace validation +{ +namespace +{ +constexpr RelativeTolerance tolerance_f16(0.001f); /**< Tolerance value for comparing reference's output against implementation's output for DataType::F16 */ +constexpr RelativeTolerance tolerance_f32(0.001f); /**< Tolerance value for comparing reference's output against implementation's output for DataType::F32 */ +} // namespace + +TEST_SUITE(CL) +TEST_SUITE(NormalizePlanarYUVLayer) + +template +using CLNormalizePlanarYUVLayerFixture = NormalizePlanarYUVLayerValidationFixture; + +DATA_TEST_CASE(Configuration, framework::DatasetMode::ALL, combine(combine(datasets::RandomNormalizePlanarYUVLayerDataset(), framework::dataset::make("DataType", { DataType::F16 })), + framework::dataset::make("DataLayout", { DataLayout::NCHW, DataLayout::NHWC })), + shape0, shape1, dt, data_layout) +{ + TensorShape src_dst_shapes = shape0; + if(data_layout == DataLayout::NHWC) + { + permute(src_dst_shapes, PermutationVector(2U, 0U, 1U)); + } + + // Create tensors + CLTensor src = create_tensor(src_dst_shapes, dt, 1, QuantizationInfo(), data_layout); + CLTensor dst = create_tensor(src_dst_shapes, dt, 1, QuantizationInfo(), data_layout); + CLTensor mean = create_tensor(shape1, dt, 1); + CLTensor sd = create_tensor(shape1, dt, 1); + + // Create and Configure function + CLNormalizePlanarYUVLayer norm; + norm.configure(&src, &dst, &mean, &sd); + + // Validate valid region + const ValidRegion valid_region = shape_to_valid_region(src_dst_shapes); + validate(dst.info()->valid_region(), valid_region); +} + +// *INDENT-OFF* +// clang-format off +DATA_TEST_CASE(Validate, framework::DatasetMode::ALL, zip(zip(zip( + framework::dataset::make("InputInfo", { TensorInfo(TensorShape(27U, 13U, 2U), 1, DataType::F32), // Mismatching data types + TensorInfo(TensorShape(27U, 13U, 2U), 1, DataType::F16), // Window shrink + TensorInfo(TensorShape(32U, 13U, 2U), 1, DataType::U8), // Unsupported data type + TensorInfo(TensorShape(32U, 16U, 8U), 1, DataType::F16), + TensorInfo(TensorShape(32U, 13U, 2U), 1, DataType::F16), // Mismatching mean and sd shapes + TensorInfo(TensorShape(27U, 13U, 2U), 1, DataType::F32), // Mismatching shapes + }), + framework::dataset::make("OutputInfo",{ TensorInfo(TensorShape(27U, 13U, 2U), 1, DataType::F16), + TensorInfo(TensorShape(27U, 13U, 2U), 1, DataType::F16), + TensorInfo(TensorShape(32U, 13U, 2U), 1, DataType::U8), + TensorInfo(TensorShape(32U, 16U, 8U), 1, DataType::F16), + TensorInfo(TensorShape(32U, 13U, 2U), 1, DataType::F16), + TensorInfo(TensorShape(30U, 11U, 2U), 1, DataType::F32), + })), + framework::dataset::make("MSTDInfo",{ TensorInfo(TensorShape(2U), 1, DataType::F16), + TensorInfo(TensorShape(2U), 1, DataType::F16), + TensorInfo(TensorShape(2U), 1, DataType::U8), + TensorInfo(TensorShape(8U), 1, DataType::F16), + TensorInfo(TensorShape(6U), 1, DataType::F16), + TensorInfo(TensorShape(2U), 1, DataType::F32), + })), + framework::dataset::make("Expected", { false, false, false, true, false, false })), + input_info, output_info, msd_info, expected) +{ + const auto &mean_info = msd_info; + const auto &sd_info = msd_info; + bool has_error = bool(CLNormalizePlanarYUVLayer::validate(&input_info.clone()->set_is_resizable(false), &output_info.clone()->set_is_resizable(false), &mean_info.clone()->set_is_resizable(false), &sd_info.clone()->set_is_resizable(false))); + ARM_COMPUTE_EXPECT(has_error == expected, framework::LogLevel::ERRORS); +} +// clang-format on +// *INDENT-ON* + +TEST_SUITE(Float) +TEST_SUITE(FP16) +FIXTURE_DATA_TEST_CASE(Random, CLNormalizePlanarYUVLayerFixture, framework::DatasetMode::PRECOMMIT, combine(combine(datasets::RandomNormalizePlanarYUVLayerDataset(), + framework::dataset::make("DataType", DataType::F16)), + framework::dataset::make("DataLayout", { DataLayout::NCHW, DataLayout::NHWC }))) +{ + // Validate output + validate(CLAccessor(_target), _reference, tolerance_f16, 0); +} +TEST_SUITE_END() + +TEST_SUITE(FP32) +FIXTURE_DATA_TEST_CASE(Random, CLNormalizePlanarYUVLayerFixture, framework::DatasetMode::PRECOMMIT, combine(combine(datasets::RandomNormalizePlanarYUVLayerDataset(), + framework::dataset::make("DataType", DataType::F32)), + framework::dataset::make("DataLayout", { DataLayout::NCHW, DataLayout::NHWC }))) +{ + // Validate output + validate(CLAccessor(_target), _reference, tolerance_f32); +} +TEST_SUITE_END() +TEST_SUITE_END() + +TEST_SUITE_END() +TEST_SUITE_END() +} // namespace validation +} // namespace test +} // namespace arm_compute diff --git a/tests/validation/GLES_COMPUTE/NormalizePlanarYUVLayer.cpp b/tests/validation/GLES_COMPUTE/NormalizePlanarYUVLayer.cpp index e06b19cfea..540a2be143 100644 --- a/tests/validation/GLES_COMPUTE/NormalizePlanarYUVLayer.cpp +++ b/tests/validation/GLES_COMPUTE/NormalizePlanarYUVLayer.cpp @@ -1,5 +1,5 @@ /* - * Copyright (c) 2017 ARM Limited. + * Copyright (c) 2017-2018 ARM Limited. * * SPDX-License-Identifier: MIT * @@ -70,10 +70,46 @@ DATA_TEST_CASE(Configuration, framework::DatasetMode::ALL, combine(datasets::Ran validate(dst.info()->valid_region(), valid_region); } +// *INDENT-OFF* +// clang-format off +DATA_TEST_CASE(Validate, framework::DatasetMode::ALL, zip(zip(zip( + framework::dataset::make("InputInfo", { TensorInfo(TensorShape(27U, 13U, 2U), 1, DataType::F16), // Mismatching data types + TensorInfo(TensorShape(27U, 13U, 2U), 1, DataType::F16), // Window shrink + TensorInfo(TensorShape(32U, 13U, 2U), 1, DataType::U8), // Unsupported data type + TensorInfo(TensorShape(32U, 16U, 8U), 1, DataType::F16), + TensorInfo(TensorShape(32U, 13U, 2U), 1, DataType::F16), // Mismatching mean and sd shapes + TensorInfo(TensorShape(27U, 13U, 2U), 1, DataType::F16), // Mismatching shapes + }), + framework::dataset::make("OutputInfo",{ TensorInfo(TensorShape(27U, 13U, 2U), 1, DataType::F32), + TensorInfo(TensorShape(27U, 13U, 2U), 1, DataType::F16), + TensorInfo(TensorShape(32U, 13U, 2U), 1, DataType::U8), + TensorInfo(TensorShape(32U, 16U, 8U), 1, DataType::F16), + TensorInfo(TensorShape(32U, 13U, 2U), 1, DataType::F16), + TensorInfo(TensorShape(30U, 11U, 2U), 1, DataType::F16), + })), + framework::dataset::make("MSTDInfo",{ TensorInfo(TensorShape(2U), 1, DataType::F16), + TensorInfo(TensorShape(2U), 1, DataType::F16), + TensorInfo(TensorShape(2U), 1, DataType::U8), + TensorInfo(TensorShape(8U), 1, DataType::F16), + TensorInfo(TensorShape(6U), 1, DataType::F16), + TensorInfo(TensorShape(2U), 1, DataType::F16), + })), + framework::dataset::make("Expected", { false, false, false, true, false, false })), + input_info, output_info, msd_info, expected) +{ + const auto &mean_info = msd_info; + const auto &sd_info = msd_info; + bool has_error = bool(GCNormalizePlanarYUVLayer::validate(&input_info.clone()->set_is_resizable(false), &output_info.clone()->set_is_resizable(false), &mean_info.clone()->set_is_resizable(false), &sd_info.clone()->set_is_resizable(false))); + ARM_COMPUTE_EXPECT(has_error == expected, framework::LogLevel::ERRORS); +} +// clang-format on +// *INDENT-ON* + TEST_SUITE(Float) TEST_SUITE(FP16) -FIXTURE_DATA_TEST_CASE(Random, GCNormalizePlanarYUVLayerFixture, framework::DatasetMode::PRECOMMIT, combine(datasets::RandomNormalizePlanarYUVLayerDataset(), - framework::dataset::make("DataType", DataType::F16))) +FIXTURE_DATA_TEST_CASE(Random, GCNormalizePlanarYUVLayerFixture, framework::DatasetMode::PRECOMMIT, combine(combine(datasets::RandomNormalizePlanarYUVLayerDataset(), + framework::dataset::make("DataType", DataType::F16)), + framework::dataset::make("DataLayout", { DataLayout::NCHW }))) { // Validate output validate(GCAccessor(_target), _reference, tolerance_f16, 0); diff --git a/tests/validation/fixtures/NormalizePlanarYUVLayerFixture.h b/tests/validation/fixtures/NormalizePlanarYUVLayerFixture.h index 09905cfef7..cc73e530ef 100644 --- a/tests/validation/fixtures/NormalizePlanarYUVLayerFixture.h +++ b/tests/validation/fixtures/NormalizePlanarYUVLayerFixture.h @@ -45,16 +45,16 @@ class NormalizePlanarYUVLayerValidationFixture : public framework::Fixture { public: template - void setup(TensorShape shape0, TensorShape shape1, DataType dt) + void setup(TensorShape shape0, TensorShape shape1, DataType dt, DataLayout data_layout) { _data_type = dt; - _target = compute_target(shape0, shape1, dt); + _target = compute_target(shape0, shape1, dt, data_layout); _reference = compute_reference(shape0, shape1, dt); } protected: template - void fill(U &&src_tensor, U &&mean_tensor, U &&sd_tensor) + void fill(U &&src_tensor, U &&mean_tensor, U &&std_tensor) { if(is_data_type_float(_data_type)) { @@ -62,43 +62,48 @@ protected: float max_bound = 0.f; std::tie(min_bound, max_bound) = get_normalize_planar_yuv_layer_test_bounds(); std::uniform_real_distribution<> distribution(min_bound, max_bound); - std::uniform_real_distribution<> distribution_sd(0.1, max_bound); + std::uniform_real_distribution<> distribution_std(0.1, max_bound); library->fill(src_tensor, distribution, 0); library->fill(mean_tensor, distribution, 1); - library->fill(sd_tensor, distribution_sd, 2); + library->fill(std_tensor, distribution_std, 2); } } - TensorType compute_target(const TensorShape &shape0, const TensorShape &shape1, DataType dt) + TensorType compute_target(TensorShape shape0, const TensorShape &shape1, DataType dt, DataLayout data_layout) { + if(data_layout == DataLayout::NHWC) + { + permute(shape0, PermutationVector(2U, 0U, 1U)); + } + // Create tensors - TensorType src = create_tensor(shape0, dt, 1); - TensorType dst = create_tensor(shape0, dt, 1); + TensorType src = create_tensor(shape0, dt, 1, QuantizationInfo(), data_layout); + TensorType dst = create_tensor(shape0, dt, 1, QuantizationInfo(), data_layout); TensorType mean = create_tensor(shape1, dt, 1); - TensorType sd = create_tensor(shape1, dt, 1); + TensorType std = create_tensor(shape1, dt, 1); // Create and configure function FunctionType norm; - norm.configure(&src, &dst, &mean, &sd); + norm.configure(&src, &dst, &mean, &std); ARM_COMPUTE_EXPECT(src.info()->is_resizable(), framework::LogLevel::ERRORS); ARM_COMPUTE_EXPECT(dst.info()->is_resizable(), framework::LogLevel::ERRORS); ARM_COMPUTE_EXPECT(mean.info()->is_resizable(), framework::LogLevel::ERRORS); - ARM_COMPUTE_EXPECT(sd.info()->is_resizable(), framework::LogLevel::ERRORS); + ARM_COMPUTE_EXPECT(std.info()->is_resizable(), framework::LogLevel::ERRORS); // Allocate tensors src.allocator()->allocate(); dst.allocator()->allocate(); mean.allocator()->allocate(); - sd.allocator()->allocate(); + std.allocator()->allocate(); ARM_COMPUTE_EXPECT(!src.info()->is_resizable(), framework::LogLevel::ERRORS); ARM_COMPUTE_EXPECT(!dst.info()->is_resizable(), framework::LogLevel::ERRORS); ARM_COMPUTE_EXPECT(!mean.info()->is_resizable(), framework::LogLevel::ERRORS); - ARM_COMPUTE_EXPECT(!sd.info()->is_resizable(), framework::LogLevel::ERRORS); + ARM_COMPUTE_EXPECT(!std.info()->is_resizable(), framework::LogLevel::ERRORS); // Fill tensors - fill(AccessorType(src), AccessorType(mean), AccessorType(sd)); + fill(AccessorType(src), AccessorType(mean), AccessorType(std)); // Compute function norm.run(); @@ -111,12 +116,12 @@ protected: // Create reference SimpleTensor ref_src{ shape0, dt, 1 }; SimpleTensor ref_mean{ shape1, dt, 1 }; - SimpleTensor ref_sd{ shape1, dt, 1 }; + SimpleTensor ref_std{ shape1, dt, 1 }; // Fill reference - fill(ref_src, ref_mean, ref_sd); + fill(ref_src, ref_mean, ref_std); - return reference::normalize_planar_yuv_layer(ref_src, ref_mean, ref_sd); + return reference::normalize_planar_yuv_layer(ref_src, ref_mean, ref_std); } TensorType _target{}; diff --git a/tests/validation/reference/NormalizePlanarYUVLayer.cpp b/tests/validation/reference/NormalizePlanarYUVLayer.cpp index 2442943bb4..afb899220d 100644 --- a/tests/validation/reference/NormalizePlanarYUVLayer.cpp +++ b/tests/validation/reference/NormalizePlanarYUVLayer.cpp @@ -1,5 +1,5 @@ /* - * Copyright (c) 2017 ARM Limited. + * Copyright (c) 2017-2018 ARM Limited. * * SPDX-License-Identifier: MIT * @@ -35,7 +35,7 @@ namespace reference { // NormalizePlanarYUV Layer for floating point type template ::value, int>::type *> -SimpleTensor normalize_planar_yuv_layer(const SimpleTensor &src, const SimpleTensor &mean, const SimpleTensor &sd) +SimpleTensor normalize_planar_yuv_layer(const SimpleTensor &src, const SimpleTensor &mean, const SimpleTensor &std) { SimpleTensor result(src.shape(), src.data_type()); @@ -53,7 +53,7 @@ SimpleTensor normalize_planar_yuv_layer(const SimpleTensor &src, const Sim for(int l = 0; l < cols; ++l) { const int pos = l + k * cols + i * rows * cols + r * cols * rows * depth; - result[pos] = (src[pos] - mean[i]) / sd[i]; + result[pos] = (src[pos] - mean[i]) / std[i]; } } } @@ -61,8 +61,8 @@ SimpleTensor normalize_planar_yuv_layer(const SimpleTensor &src, const Sim return result; } -template SimpleTensor normalize_planar_yuv_layer(const SimpleTensor &src, const SimpleTensor &mean, const SimpleTensor &sd); - +template SimpleTensor normalize_planar_yuv_layer(const SimpleTensor &src, const SimpleTensor &mean, const SimpleTensor &std); +template SimpleTensor normalize_planar_yuv_layer(const SimpleTensor &src, const SimpleTensor &mean, const SimpleTensor &std); } // namespace reference } // namespace validation } // namespace test diff --git a/tests/validation/reference/NormalizePlanarYUVLayer.h b/tests/validation/reference/NormalizePlanarYUVLayer.h index c8740a383b..41ce48630c 100644 --- a/tests/validation/reference/NormalizePlanarYUVLayer.h +++ b/tests/validation/reference/NormalizePlanarYUVLayer.h @@ -1,5 +1,5 @@ /* - * Copyright (c) 2017 ARM Limited. + * Copyright (c) 2017-2018 ARM Limited. * * SPDX-License-Identifier: MIT * @@ -36,7 +36,7 @@ namespace validation namespace reference { template ::value, int>::type * = nullptr> -SimpleTensor normalize_planar_yuv_layer(const SimpleTensor &src, const SimpleTensor &mean, const SimpleTensor &sd); +SimpleTensor normalize_planar_yuv_layer(const SimpleTensor &src, const SimpleTensor &mean, const SimpleTensor &std); } // namespace reference } // namespace validation } // namespace test -- cgit v1.2.1