From e93626222041917617cbe1ccccbb472b47a2358c Mon Sep 17 00:00:00 2001 From: Michalis Spyrou Date: Fri, 23 Nov 2018 17:41:37 +0000 Subject: COMPMID-1719 CL: Implement RSqrt, Exp Change-Id: I827b26239043a9e90d26c2583122648d2a45303a Reviewed-on: https://review.mlplatform.org/317 Reviewed-by: Georgios Pinitas Tested-by: Arm Jenkins --- arm_compute/core/CL/CLKernels.h | 1 + .../CL/kernels/CLElementWiseUnaryLayerKernel.h | 60 +++++++++ arm_compute/core/Types.h | 7 + arm_compute/runtime/CL/CLFunctions.h | 1 + .../runtime/CL/functions/CLElementWiseUnaryLayer.h | 73 +++++++++++ src/core/CL/CLKernelLibrary.cpp | 5 + src/core/CL/cl_kernels/elementwise_unary.cl | 87 ++++++++++++ .../CL/kernels/CLElementWiseUnaryLayerKernel.cpp | 117 +++++++++++++++++ .../CL/functions/CLElementWiseUnaryLayer.cpp | 54 ++++++++ tests/validation/CL/ExpLayer.cpp | 110 ++++++++++++++++ tests/validation/CL/RsqrtLayer.cpp | 110 ++++++++++++++++ .../validation/fixtures/ElementWiseUnaryFixture.h | 146 +++++++++++++++++++++ tests/validation/reference/ElementWiseUnary.cpp | 62 +++++++++ tests/validation/reference/ElementWiseUnary.h | 43 ++++++ utils/TypePrinter.h | 37 ++++++ 15 files changed, 913 insertions(+) create mode 100644 arm_compute/core/CL/kernels/CLElementWiseUnaryLayerKernel.h create mode 100644 arm_compute/runtime/CL/functions/CLElementWiseUnaryLayer.h create mode 100644 src/core/CL/cl_kernels/elementwise_unary.cl create mode 100644 src/core/CL/kernels/CLElementWiseUnaryLayerKernel.cpp create mode 100644 src/runtime/CL/functions/CLElementWiseUnaryLayer.cpp create mode 100644 tests/validation/CL/ExpLayer.cpp create mode 100644 tests/validation/CL/RsqrtLayer.cpp create mode 100644 tests/validation/fixtures/ElementWiseUnaryFixture.h create mode 100644 tests/validation/reference/ElementWiseUnary.cpp create mode 100644 tests/validation/reference/ElementWiseUnary.h diff --git a/arm_compute/core/CL/CLKernels.h b/arm_compute/core/CL/CLKernels.h index 201679347a..9d7ec71a32 100644 --- a/arm_compute/core/CL/CLKernels.h +++ b/arm_compute/core/CL/CLKernels.h @@ -59,6 +59,7 @@ #include "arm_compute/core/CL/kernels/CLDilateKernel.h" #include "arm_compute/core/CL/kernels/CLDirectConvolutionLayerKernel.h" #include "arm_compute/core/CL/kernels/CLDirectConvolutionLayerOutputStageKernel.h" +#include "arm_compute/core/CL/kernels/CLElementWiseUnaryLayerKernel.h" #include "arm_compute/core/CL/kernels/CLElementwiseOperationKernel.h" #include "arm_compute/core/CL/kernels/CLErodeKernel.h" #include "arm_compute/core/CL/kernels/CLFastCornersKernel.h" diff --git a/arm_compute/core/CL/kernels/CLElementWiseUnaryLayerKernel.h b/arm_compute/core/CL/kernels/CLElementWiseUnaryLayerKernel.h new file mode 100644 index 0000000000..0a3511c1f1 --- /dev/null +++ b/arm_compute/core/CL/kernels/CLElementWiseUnaryLayerKernel.h @@ -0,0 +1,60 @@ +/* + * 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_CLELEMENTWISEUNARYLAYERKERNEL_H__ +#define __ARM_COMPUTE_CLELEMENTWISEUNARYLAYERKERNEL_H__ + +#include "arm_compute/core/CL/ICLKernel.h" +#include "arm_compute/core/CL/ICLSimpleKernel.h" +#include "arm_compute/core/Types.h" + +namespace arm_compute +{ +class ICLTensor; + +/** Interface for the elementwise unary operator */ +class CLElementWiseUnaryLayerKernel : public ICLSimpleKernel +{ +public: + /** Initialise the kernel's inputs, output. + * + * @param[in] input First tensor input. Data types supported: F16/F32. + * @param[out] output Output tensor. Data types supported: Same as @p input. + * @param[in] op Element wise unary operation to perform. + */ + void configure(const ICLTensor *input, ICLTensor *output, const ElementWiseUnary &op); + /** Static function to check if given info will lead to a valid configuration of @ref CLElementWiseUnaryLayerKernel + * + * @param[in] input First tensor input info. Data types supported: F16/F32. + * @param[in] output Output tensor info. Data types supported: Same as @p input. + * @param[in] op Element wise unary operation to perform. + * + * @return a status + */ + static Status validate(const ITensorInfo *input, const ITensorInfo *output, const ElementWiseUnary &op); + + // Inherited methods overridden: + void run(const Window &window, cl::CommandQueue &queue) override; +}; +} // namespace arm_compute +#endif /* __ARM_COMPUTE_CLELEMENTWISEUNARYLAYERKERNEL_H__ */ diff --git a/arm_compute/core/Types.h b/arm_compute/core/Types.h index d46c93247c..6833a66cd9 100644 --- a/arm_compute/core/Types.h +++ b/arm_compute/core/Types.h @@ -563,6 +563,13 @@ enum class ArithmeticOperation SQUARED_DIFF, /**< (x - y)^2 */ }; +/** Available element wise unary operations */ +enum class ElementWiseUnary +{ + RSQRT, /**< Reverse square root */ + EXP, /**< Exponential */ +}; + /** The normalization type used for the normalization layer */ enum class NormType { diff --git a/arm_compute/runtime/CL/CLFunctions.h b/arm_compute/runtime/CL/CLFunctions.h index 3f2df0efff..d4827af88a 100644 --- a/arm_compute/runtime/CL/CLFunctions.h +++ b/arm_compute/runtime/CL/CLFunctions.h @@ -60,6 +60,7 @@ #include "arm_compute/runtime/CL/functions/CLDerivative.h" #include "arm_compute/runtime/CL/functions/CLDilate.h" #include "arm_compute/runtime/CL/functions/CLDirectConvolutionLayer.h" +#include "arm_compute/runtime/CL/functions/CLElementWiseUnaryLayer.h" #include "arm_compute/runtime/CL/functions/CLElementwiseOperations.h" #include "arm_compute/runtime/CL/functions/CLEqualizeHistogram.h" #include "arm_compute/runtime/CL/functions/CLErode.h" diff --git a/arm_compute/runtime/CL/functions/CLElementWiseUnaryLayer.h b/arm_compute/runtime/CL/functions/CLElementWiseUnaryLayer.h new file mode 100644 index 0000000000..37061a647c --- /dev/null +++ b/arm_compute/runtime/CL/functions/CLElementWiseUnaryLayer.h @@ -0,0 +1,73 @@ +/* + * 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_CLELEMENTWISEUNARYLAYER_H__ +#define __ARM_COMPUTE_CLELEMENTWISEUNARYLAYER_H__ + +#include "arm_compute/runtime/CL/ICLSimpleFunction.h" + +namespace arm_compute +{ +class ICLTensor; + +/** Basic function to perform inverse square root on an input tensor. */ +class CLRsqrtLayer : public ICLSimpleFunction +{ +public: + /** Initialize the function + * + * @param[in] input Input tensor. Data types supported: F16/F32. + * @param[out] output Output tensor. Data types supported: same as @p input. + */ + void configure(const ICLTensor *input, ICLTensor *output); + /** Static function to check if given info will lead to a valid configuration of @ref CLRsqrtLayer + * + * @param[in] input First tensor input info. Data types supported: F16/F32. + * @param[in] output Output tensor info. Data types supported: Same as @p input. + * + * @return a status + */ + static Status validate(const ITensorInfo *input, const ITensorInfo *output); +}; + +/** Basic function to perform exponential on an input tensor. */ +class CLExpLayer : public ICLSimpleFunction +{ +public: + /** Initialize the function + * + * @param[in] input Input tensor. Data types supported: F16/F32. + * @param[out] output Output tensor. Data types supported: same as @p input. + */ + void configure(const ICLTensor *input, ICLTensor *output); + /** Static function to check if given info will lead to a valid configuration of @ref CLExpLayer + * + * @param[in] input First tensor input info. Data types supported: F16/F32. + * @param[in] output Output tensor info. Data types supported: Same as @p input. + * + * @return a status + */ + static Status validate(const ITensorInfo *input, const ITensorInfo *output); +}; +} // namespace arm_compute +#endif /* __ARM_COMPUTE_CLELEMENTWISEUNARYLAYER_H__ */ diff --git a/src/core/CL/CLKernelLibrary.cpp b/src/core/CL/CLKernelLibrary.cpp index 1361d02f74..da85472005 100644 --- a/src/core/CL/CLKernelLibrary.cpp +++ b/src/core/CL/CLKernelLibrary.cpp @@ -253,6 +253,7 @@ const std::map CLKernelLibrary::_kernel_program_map = { "elementwise_operation_MIN_quantized", "elementwise_operation_quantized.cl" }, { "elementwise_operation_DIV_quantized", "elementwise_operation_quantized.cl" }, { "elementwise_operation_SQUARED_DIFF_quantized", "elementwise_operation_quantized.cl" }, + { "elementwise_unary", "elementwise_unary.cl" }, { "erode", "erode.cl" }, { "fast_corners", "fast_corners.cl" }, { "flatten", "flatten.cl" }, @@ -649,6 +650,10 @@ const std::map CLKernelLibrary::_program_source_map = { "elementwise_operation_quantized.cl", #include "./cl_kernels/elementwise_operation_quantized.clembed" + }, + { + "elementwise_unary.cl", +#include "./cl_kernels/elementwise_unary.clembed" }, { "erode.cl", diff --git a/src/core/CL/cl_kernels/elementwise_unary.cl b/src/core/CL/cl_kernels/elementwise_unary.cl new file mode 100644 index 0000000000..bccb47ed1f --- /dev/null +++ b/src/core/CL/cl_kernels/elementwise_unary.cl @@ -0,0 +1,87 @@ +/* + * 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" +#include "warp_helpers.h" + +#if defined(DATA_TYPE) && defined(VEC_SIZE) && defined(OPERATION) +/** Calculate reverse square root + * + * @param[in] input Pointer to the first element. + * + * @return reverse square root + */ +inline VEC_DATA_TYPE(DATA_TYPE, VEC_SIZE) inverse_sqrt(const VEC_DATA_TYPE(DATA_TYPE, VEC_SIZE) input) +{ + return rsqrt(input); +} + +/** Calculate exponential + * + * @param[in] input Pointer to the first element. + * + * @return exponential + */ +inline VEC_DATA_TYPE(DATA_TYPE, VEC_SIZE) exponential(const VEC_DATA_TYPE(DATA_TYPE, VEC_SIZE) input) +{ + return exp(input); +} + +/** Applies element wise unary operator in a tensor. + * + * @param[in] in_ptr Pointer to the source image. Supported data types: F16/32. + * @param[in] in_stride_x Stride of the source image in X dimension (in bytes) + * @param[in] in_step_x in_stride_x * number of elements along X processed per work item (in bytes) + * @param[in] in_offset_first_element_in_bytes Offset of the first element in the source image + * @param[out] out_ptr Pointer to the destination image. Supported data types: F16/32. + * @param[in] out_stride_x Stride of the destination image in X dimension (in bytes) + * @param[in] out_step_y out_stride_y * number of elements along Y processed per work item (in bytes) + * @param[in] out_offset_first_element_in_bytes Offset of the first element in the destination image + */ +__kernel void elementwise_unary( + VECTOR_DECLARATION(in), + VECTOR_DECLARATION(out)) +{ + Vector in = CONVERT_TO_VECTOR_STRUCT(in); + Vector out = CONVERT_TO_VECTOR_STRUCT(out); + +#if defined(VEC_SIZE) && defined(LAST_ACCESSED_X) + // Check if access on width gets out of bounds + // If it does shift access vector to access elements within bounds + const int xi = (int)(get_global_id(0) * VEC_SIZE); + in.ptr -= max(xi - (int)LAST_ACCESSED_X, 0) * in_stride_x; + out.ptr -= max(xi - (int)LAST_ACCESSED_X, 0) * out_stride_x; + + VEC_DATA_TYPE(DATA_TYPE, VEC_SIZE) + data = VLOAD(VEC_SIZE)(0, (__global DATA_TYPE *)in.ptr); + + VSTORE(VEC_SIZE) + (OPERATION(data), 0, (__global DATA_TYPE *)out.ptr); +#else // !defined(VEC_SIZE) || !defined(LAST_ACCESSED_X) + VEC_DATA_TYPE(DATA_TYPE, VEC_SIZE) + data = VLOAD(VEC_SIZE)(0, (__global DATA_TYPE *)in.ptr); + VSTORE(VEC_SIZE) + (OPERATION(data), 0, (__global DATA_TYPE *)out.ptr); +#endif // defined(VEC_SIZE) && defined(LAST_ACCESSED_X) +} +#endif // defined(DATA_TYPE) && defined(VEC_SIZE) && defined(OPERATION) diff --git a/src/core/CL/kernels/CLElementWiseUnaryLayerKernel.cpp b/src/core/CL/kernels/CLElementWiseUnaryLayerKernel.cpp new file mode 100644 index 0000000000..6d2105f3b1 --- /dev/null +++ b/src/core/CL/kernels/CLElementWiseUnaryLayerKernel.cpp @@ -0,0 +1,117 @@ +/* + * 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/CLElementWiseUnaryLayerKernel.h" + +#include "arm_compute/core/CL/CLHelpers.h" +#include "arm_compute/core/CL/CLValidate.h" +#include "arm_compute/core/CL/ICLTensor.h" + +using namespace arm_compute; + +namespace +{ +Status validate_arguments(const ITensorInfo &input, const ITensorInfo &output) +{ + ARM_COMPUTE_RETURN_ERROR_ON_F16_UNSUPPORTED(&input); + ARM_COMPUTE_RETURN_ERROR_ON_DATA_TYPE_CHANNEL_NOT_IN(&input, 1, DataType::F16, DataType::F32); + + // Validate in case of configured output + if(output.total_size() > 0) + { + ARM_COMPUTE_RETURN_ERROR_ON_F16_UNSUPPORTED(&output); + ARM_COMPUTE_RETURN_ERROR_ON_DATA_TYPE_CHANNEL_NOT_IN(&output, 1, DataType::F16, DataType::F32); + } + + return Status{}; +} +} // namespace + +void CLElementWiseUnaryLayerKernel::configure(const ICLTensor *input, ICLTensor *output, const ElementWiseUnary &op) +{ + ARM_COMPUTE_ERROR_ON_NULLPTR(input, output); + ARM_COMPUTE_ERROR_THROW_ON(validate_arguments(*input->info(), *output->info())); + + // Configure kernel window + _input = input; + _output = output; + + const std::string kernel_name = "elementwise_unary"; + const int vec_size_x = 16 / output->info()->element_size(); + const int output_width_x = output->info()->tensor_shape().x(); + const bool multi_access_x = (output_width_x / vec_size_x > 0); + + Window win = calculate_max_window(*output->info()); + if(multi_access_x) + { + win.set(Window::DimX, + Window::Dimension(win.x().start(), ceil_to_multiple(win.x().end(), vec_size_x), vec_size_x)); + } + ICLKernel::configure_internal(win); + + // Set kernel 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_if(multi_access_x, "-DVEC_SIZE=" + support::cpp11::to_string(vec_size_x)); + build_opts.add_option_if(multi_access_x, "-DLAST_ACCESSED_X=" + support::cpp11::to_string(std::max(output_width_x - vec_size_x, 0))); + switch(op) + { + case ElementWiseUnary::RSQRT: + build_opts.add_option("-DOPERATION=inverse_sqrt"); + break; + case ElementWiseUnary::EXP: + build_opts.add_option("-DOPERATION=exponential"); + break; + default: + ARM_COMPUTE_ERROR("Not implemented"); + } + + // Create kernel + _kernel = static_cast(CLKernelLibrary::get().create_kernel(kernel_name, build_opts.options())); +} + +Status CLElementWiseUnaryLayerKernel::validate(const ITensorInfo *input, const ITensorInfo *output, const ElementWiseUnary &op) +{ + ARM_COMPUTE_UNUSED(op); + ARM_COMPUTE_RETURN_ERROR_ON_NULLPTR(input, output); + ARM_COMPUTE_RETURN_ON_ERROR(validate_arguments(*input, *output)); + + return Status{}; +} + +void CLElementWiseUnaryLayerKernel::run(const Window &window, cl::CommandQueue &queue) +{ + ARM_COMPUTE_ERROR_ON_UNCONFIGURED_KERNEL(this); + ARM_COMPUTE_ERROR_ON_INVALID_SUBWINDOW(ICLKernel::window(), window); + + Window collapsed = window.collapse_if_possible(ICLKernel::window(), Window::DimX); + + do + { + unsigned int idx = 0; + add_1D_tensor_argument(idx, _input, collapsed); + add_1D_tensor_argument(idx, _output, collapsed); + enqueue(queue, *this, collapsed); + } + while(window.slide_window_slice_1D(collapsed)); +} \ No newline at end of file diff --git a/src/runtime/CL/functions/CLElementWiseUnaryLayer.cpp b/src/runtime/CL/functions/CLElementWiseUnaryLayer.cpp new file mode 100644 index 0000000000..b7e9a68680 --- /dev/null +++ b/src/runtime/CL/functions/CLElementWiseUnaryLayer.cpp @@ -0,0 +1,54 @@ +/* + * 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/CLElementWiseUnaryLayer.h" + +#include "arm_compute/core/CL/kernels/CLElementWiseUnaryLayerKernel.h" +#include "support/ToolchainSupport.h" + +#include + +namespace arm_compute +{ +void CLRsqrtLayer::configure(const ICLTensor *input, ICLTensor *output) +{ + auto k = arm_compute::support::cpp14::make_unique(); + k->configure(input, output, ElementWiseUnary::RSQRT); + _kernel = std::move(k); +} +Status CLRsqrtLayer::validate(const ITensorInfo *input, const ITensorInfo *output) +{ + return CLElementWiseUnaryLayerKernel::validate(input, output, ElementWiseUnary::RSQRT); +} + +void CLExpLayer::configure(const ICLTensor *input, ICLTensor *output) +{ + auto k = arm_compute::support::cpp14::make_unique(); + k->configure(input, output, ElementWiseUnary::EXP); + _kernel = std::move(k); +} +Status CLExpLayer::validate(const ITensorInfo *input, const ITensorInfo *output) +{ + return CLElementWiseUnaryLayerKernel::validate(input, output, ElementWiseUnary::EXP); +} +} // namespace arm_compute diff --git a/tests/validation/CL/ExpLayer.cpp b/tests/validation/CL/ExpLayer.cpp new file mode 100644 index 0000000000..08d1f33a65 --- /dev/null +++ b/tests/validation/CL/ExpLayer.cpp @@ -0,0 +1,110 @@ +/* + * 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/functions/CLElementWiseUnaryLayer.h" +#include "arm_compute/runtime/Tensor.h" +#include "arm_compute/runtime/TensorAllocator.h" +#include "tests/CL/CLAccessor.h" +#include "tests/PaddingCalculator.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/ElementWiseUnaryFixture.h" + +namespace arm_compute +{ +namespace test +{ +namespace validation +{ +namespace +{ +RelativeTolerance tolerance_fp32(0.000001f); +RelativeTolerance tolerance_fp16(0.001f); +} // namespace +TEST_SUITE(CL) +TEST_SUITE(ExpLayer) + +DATA_TEST_CASE(Configuration, framework::DatasetMode::ALL, combine(concat(datasets::SmallShapes(), datasets::LargeShapes()), framework::dataset::make("DataType", DataType::F32)), shape, data_type) +{ + // Create tensors + CLTensor src = create_tensor(shape, data_type); + CLTensor dst = create_tensor(shape, data_type); + + ARM_COMPUTE_EXPECT(src.info()->is_resizable(), framework::LogLevel::ERRORS); + ARM_COMPUTE_EXPECT(dst.info()->is_resizable(), framework::LogLevel::ERRORS); + + // Create and configure function + CLExpLayer exp_layer; + exp_layer.configure(&src, &dst); + + // Validate valid region + const ValidRegion valid_region = shape_to_valid_region(shape); + validate(src.info()->valid_region(), valid_region); + validate(dst.info()->valid_region(), valid_region); +} + +template +using CLExpLayerFixture = ExpValidationFixture; + +TEST_SUITE(Float) +TEST_SUITE(FP16) +FIXTURE_DATA_TEST_CASE(RunSmall, CLExpLayerFixture, framework::DatasetMode::PRECOMMIT, combine(datasets::SmallShapes(), framework::dataset::make("DataType", + DataType::F16))) +{ + // Validate output + validate(CLAccessor(_target), _reference, tolerance_fp16); +} +FIXTURE_DATA_TEST_CASE(RunLarge, CLExpLayerFixture, framework::DatasetMode::NIGHTLY, combine(datasets::LargeShapes(), framework::dataset::make("DataType", + DataType::F16))) +{ + // Validate output + validate(CLAccessor(_target), _reference, tolerance_fp16); +} + +TEST_SUITE_END() // FP16 +TEST_SUITE(FP32) +FIXTURE_DATA_TEST_CASE(RunSmall, CLExpLayerFixture, framework::DatasetMode::PRECOMMIT, combine(datasets::SmallShapes(), framework::dataset::make("DataType", + DataType::F32))) +{ + // Validate output + validate(CLAccessor(_target), _reference, tolerance_fp32); +} +FIXTURE_DATA_TEST_CASE(RunLarge, CLExpLayerFixture, framework::DatasetMode::NIGHTLY, combine(datasets::LargeShapes(), framework::dataset::make("DataType", + DataType::F32))) +{ + // Validate output + validate(CLAccessor(_target), _reference, tolerance_fp32); +} + +TEST_SUITE_END() // FP32 +TEST_SUITE_END() // Float + +TEST_SUITE_END() // ExpLayer +TEST_SUITE_END() // CL +} // namespace validation +} // namespace test +} // namespace arm_compute diff --git a/tests/validation/CL/RsqrtLayer.cpp b/tests/validation/CL/RsqrtLayer.cpp new file mode 100644 index 0000000000..ee9e9363b3 --- /dev/null +++ b/tests/validation/CL/RsqrtLayer.cpp @@ -0,0 +1,110 @@ +/* + * 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/functions/CLElementWiseUnaryLayer.h" +#include "arm_compute/runtime/Tensor.h" +#include "arm_compute/runtime/TensorAllocator.h" +#include "tests/CL/CLAccessor.h" +#include "tests/PaddingCalculator.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/ElementWiseUnaryFixture.h" + +namespace arm_compute +{ +namespace test +{ +namespace validation +{ +namespace +{ +RelativeTolerance tolerance_fp32(0.000001f); +RelativeTolerance tolerance_fp16(0.001f); +} // namespace +TEST_SUITE(CL) +TEST_SUITE(RsqrtLayer) + +DATA_TEST_CASE(Configuration, framework::DatasetMode::ALL, combine(concat(datasets::SmallShapes(), datasets::LargeShapes()), framework::dataset::make("DataType", DataType::F32)), shape, data_type) +{ + // Create tensors + CLTensor src = create_tensor(shape, data_type); + CLTensor dst = create_tensor(shape, data_type); + + ARM_COMPUTE_EXPECT(src.info()->is_resizable(), framework::LogLevel::ERRORS); + ARM_COMPUTE_EXPECT(dst.info()->is_resizable(), framework::LogLevel::ERRORS); + + // Create and configure function + CLRsqrtLayer exp_layer; + exp_layer.configure(&src, &dst); + + // Validate valid region + const ValidRegion valid_region = shape_to_valid_region(shape); + validate(src.info()->valid_region(), valid_region); + validate(dst.info()->valid_region(), valid_region); +} + +template +using CLRsqrtLayerFixture = RsqrtValidationFixture; + +TEST_SUITE(Float) +TEST_SUITE(FP16) +FIXTURE_DATA_TEST_CASE(RunSmall, CLRsqrtLayerFixture, framework::DatasetMode::PRECOMMIT, combine(datasets::SmallShapes(), framework::dataset::make("DataType", + DataType::F16))) +{ + // Validate output + validate(CLAccessor(_target), _reference, tolerance_fp16); +} +FIXTURE_DATA_TEST_CASE(RunLarge, CLRsqrtLayerFixture, framework::DatasetMode::NIGHTLY, combine(datasets::LargeShapes(), framework::dataset::make("DataType", + DataType::F16))) +{ + // Validate output + validate(CLAccessor(_target), _reference, tolerance_fp16); +} + +TEST_SUITE_END() // FP16 +TEST_SUITE(FP32) +FIXTURE_DATA_TEST_CASE(RunSmall, CLRsqrtLayerFixture, framework::DatasetMode::PRECOMMIT, combine(datasets::SmallShapes(), framework::dataset::make("DataType", + DataType::F32))) +{ + // Validate output + validate(CLAccessor(_target), _reference, tolerance_fp32); +} +FIXTURE_DATA_TEST_CASE(RunLarge, CLRsqrtLayerFixture, framework::DatasetMode::NIGHTLY, combine(datasets::LargeShapes(), framework::dataset::make("DataType", + DataType::F32))) +{ + // Validate output + validate(CLAccessor(_target), _reference, tolerance_fp32); +} + +TEST_SUITE_END() // FP32 +TEST_SUITE_END() // Float + +TEST_SUITE_END() // RsqrtLayer +TEST_SUITE_END() // CL +} // namespace validation +} // namespace test +} // namespace arm_compute diff --git a/tests/validation/fixtures/ElementWiseUnaryFixture.h b/tests/validation/fixtures/ElementWiseUnaryFixture.h new file mode 100644 index 0000000000..f508bc1d34 --- /dev/null +++ b/tests/validation/fixtures/ElementWiseUnaryFixture.h @@ -0,0 +1,146 @@ +/* + * 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_TEST_ELEMENTWISE_UNARY_FIXTURE +#define ARM_COMPUTE_TEST_ELEMENTWISE_UNARY_FIXTURE + +#include "arm_compute/core/TensorShape.h" +#include "arm_compute/core/Types.h" +#include "tests/AssetsLibrary.h" +#include "tests/Globals.h" +#include "tests/IAccessor.h" +#include "tests/framework/Asserts.h" +#include "tests/framework/Fixture.h" +#include "tests/validation/reference/ElementWiseUnary.h" + +namespace arm_compute +{ +namespace test +{ +namespace validation +{ +template +class ElementWiseUnaryValidationFixture : public framework::Fixture +{ +public: + template + void setup(TensorShape shape, DataType data_type, ElementWiseUnary op) + { + _op = op; + _target = compute_target(shape, data_type); + _reference = compute_reference(shape, data_type); + } + +protected: + template + void fill(U &&tensor, int i) + { + switch(_op) + { + case ElementWiseUnary::EXP: + { + std::uniform_real_distribution<> distribution(-1.0f, 1.0f); + library->fill(tensor, distribution, i); + break; + } + case ElementWiseUnary::RSQRT: + { + std::uniform_real_distribution<> distribution(1.0f, 2.0f); + library->fill(tensor, distribution, i); + break; + } + default: + ARM_COMPUTE_ERROR("Not implemented"); + } + } + + TensorType compute_target(const TensorShape &shape, DataType data_type) + { + // Create tensors + TensorType src = create_tensor(shape, data_type); + TensorType dst = create_tensor(shape, data_type); + + // Create and configure function + FunctionType elwiseunary_layer; + + elwiseunary_layer.configure(&src, &dst); + + ARM_COMPUTE_EXPECT(src.info()->is_resizable(), framework::LogLevel::ERRORS); + ARM_COMPUTE_EXPECT(dst.info()->is_resizable(), framework::LogLevel::ERRORS); + + // Allocate tensors + src.allocator()->allocate(); + dst.allocator()->allocate(); + ARM_COMPUTE_EXPECT(!src.info()->is_resizable(), framework::LogLevel::ERRORS); + ARM_COMPUTE_EXPECT(!dst.info()->is_resizable(), framework::LogLevel::ERRORS); + + // Fill tensors + fill(AccessorType(src), 0); + + // Compute function + elwiseunary_layer.run(); + + return dst; + } + + SimpleTensor compute_reference(const TensorShape &shape, DataType data_type) + { + // Create reference + SimpleTensor src{ shape, data_type }; + + // Fill reference + fill(src, 0); + + return reference::elementwise_unary(src, _op); + } + + TensorType _target{}; + SimpleTensor _reference{}; + ElementWiseUnary _op{}; +}; + +template +class RsqrtValidationFixture : public ElementWiseUnaryValidationFixture +{ +public: + template + void setup(const TensorShape &shape, DataType data_type) + { + ElementWiseUnaryValidationFixture::setup(shape, data_type, ElementWiseUnary::RSQRT); + } +}; + +template +class ExpValidationFixture : public ElementWiseUnaryValidationFixture +{ +public: + template + void setup(const TensorShape &shape, DataType data_type) + { + ElementWiseUnaryValidationFixture::setup(shape, data_type, ElementWiseUnary::EXP); + } +}; +} // namespace validation +} // namespace test +} // namespace arm_compute +#endif /* ARM_COMPUTE_TEST_ELEMENTWISE_UNARY_FIXTURE */ diff --git a/tests/validation/reference/ElementWiseUnary.cpp b/tests/validation/reference/ElementWiseUnary.cpp new file mode 100644 index 0000000000..ae7f256339 --- /dev/null +++ b/tests/validation/reference/ElementWiseUnary.cpp @@ -0,0 +1,62 @@ +/* + * 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 "ElementWiseUnary.h" + +namespace arm_compute +{ +namespace test +{ +namespace validation +{ +namespace reference +{ +template +SimpleTensor elementwise_unary(const SimpleTensor &src, ElementWiseUnary op) +{ + SimpleTensor dst(src.shape(), src.data_type()); + + for(int i = 0; i < src.num_elements(); ++i) + { + switch(op) + { + case ElementWiseUnary::RSQRT: + dst[i] = 1.f / std::sqrt(src[i]); + break; + case ElementWiseUnary::EXP: + dst[i] = std::exp(src[i]); + break; + default: + ARM_COMPUTE_ERROR("Not implemented"); + } + } + + return dst; +} + +template SimpleTensor elementwise_unary(const SimpleTensor &src, ElementWiseUnary op); +template SimpleTensor elementwise_unary(const SimpleTensor &src, ElementWiseUnary op); +} // namespace reference +} // namespace validation +} // namespace test +} // namespace arm_compute diff --git a/tests/validation/reference/ElementWiseUnary.h b/tests/validation/reference/ElementWiseUnary.h new file mode 100644 index 0000000000..9c6fe14e2b --- /dev/null +++ b/tests/validation/reference/ElementWiseUnary.h @@ -0,0 +1,43 @@ +/* + * 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_TEST_ELEMENTWISE_UNARY_H__ +#define __ARM_COMPUTE_TEST_ELEMENTWISE_UNARY_H__ + +#include "tests/SimpleTensor.h" + +namespace arm_compute +{ +namespace test +{ +namespace validation +{ +namespace reference +{ +template +SimpleTensor elementwise_unary(const SimpleTensor &src, ElementWiseUnary op); +} // namespace reference +} // namespace validation +} // namespace test +} // namespace arm_compute +#endif /* __ARM_COMPUTE_TEST_ELEMENTWISE_UNARY_H__ */ diff --git a/utils/TypePrinter.h b/utils/TypePrinter.h index 27560e6b07..6850ea018a 100644 --- a/utils/TypePrinter.h +++ b/utils/TypePrinter.h @@ -1453,6 +1453,30 @@ inline ::std::ostream &operator<<(::std::ostream &os, const ComparisonOperation return os; } +/** Formatted output of the Elementwise unary Operations. + * + * @param[out] os Output stream. + * @param[in] op Type to output. + * + * @return Modified output stream. + */ +inline ::std::ostream &operator<<(::std::ostream &os, const ElementWiseUnary &op) +{ + switch(op) + { + case ElementWiseUnary::RSQRT: + os << "RSQRT"; + break; + case ElementWiseUnary::EXP: + os << "EXP"; + break; + default: + ARM_COMPUTE_ERROR("NOT_SUPPORTED!"); + } + + return os; +} + /** Formatted output of the Comparison Operations. * * @param[in] op Type to output. @@ -1466,6 +1490,19 @@ inline std::string to_string(const ComparisonOperation &op) return str.str(); } +/** Formatted output of the Elementwise unary Operations. + * + * @param[in] op Type to output. + * + * @return Formatted string. + */ +inline std::string to_string(const ElementWiseUnary &op) +{ + std::stringstream str; + str << op; + return str.str(); +} + /** Formatted output of the Norm Type. * * @param[in] type Type to output. -- cgit v1.2.1