diff options
| -rw-r--r-- | arm_compute/core/CL/CLKernels.h | 1 | ||||
| -rw-r--r-- | arm_compute/core/CL/kernels/CLElementWiseUnaryLayerKernel.h | 60 | ||||
| -rw-r--r-- | arm_compute/core/Types.h | 7 | ||||
| -rw-r--r-- | arm_compute/runtime/CL/CLFunctions.h | 1 | ||||
| -rw-r--r-- | arm_compute/runtime/CL/functions/CLElementWiseUnaryLayer.h | 73 | ||||
| -rw-r--r-- | src/core/CL/CLKernelLibrary.cpp | 5 | ||||
| -rw-r--r-- | src/core/CL/cl_kernels/elementwise_unary.cl | 87 | ||||
| -rw-r--r-- | src/core/CL/kernels/CLElementWiseUnaryLayerKernel.cpp | 117 | ||||
| -rw-r--r-- | src/runtime/CL/functions/CLElementWiseUnaryLayer.cpp | 54 | ||||
| -rw-r--r-- | tests/validation/CL/ExpLayer.cpp | 110 | ||||
| -rw-r--r-- | tests/validation/CL/RsqrtLayer.cpp | 110 | ||||
| -rw-r--r-- | tests/validation/fixtures/ElementWiseUnaryFixture.h | 146 | ||||
| -rw-r--r-- | tests/validation/reference/ElementWiseUnary.cpp | 62 | ||||
| -rw-r--r-- | tests/validation/reference/ElementWiseUnary.h | 43 | ||||
| -rw-r--r-- | utils/TypePrinter.h | 37 |
15 files changed, 913 insertions, 0 deletions
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<std::string, std::string> 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" }, @@ -651,6 +652,10 @@ const std::map<std::string, std::string> CLKernelLibrary::_program_source_map = #include "./cl_kernels/elementwise_operation_quantized.clembed" }, { + "elementwise_unary.cl", +#include "./cl_kernels/elementwise_unary.clembed" + }, + { "erode.cl", #include "./cl_kernels/erode.clembed" }, 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<int>(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<cl::Kernel>(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 <utility> + +namespace arm_compute +{ +void CLRsqrtLayer::configure(const ICLTensor *input, ICLTensor *output) +{ + auto k = arm_compute::support::cpp14::make_unique<CLElementWiseUnaryLayerKernel>(); + 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<CLElementWiseUnaryLayerKernel>(); + 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<float> tolerance_fp32(0.000001f); +RelativeTolerance<float> 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<CLTensor>(shape, data_type); + CLTensor dst = create_tensor<CLTensor>(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 <typename T> +using CLExpLayerFixture = ExpValidationFixture<CLTensor, CLAccessor, CLExpLayer, T>; + +TEST_SUITE(Float) +TEST_SUITE(FP16) +FIXTURE_DATA_TEST_CASE(RunSmall, CLExpLayerFixture<half>, 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<half>, 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<float>, 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<float>, 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<float> tolerance_fp32(0.000001f); +RelativeTolerance<float> 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<CLTensor>(shape, data_type); + CLTensor dst = create_tensor<CLTensor>(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 <typename T> +using CLRsqrtLayerFixture = RsqrtValidationFixture<CLTensor, CLAccessor, CLRsqrtLayer, T>; + +TEST_SUITE(Float) +TEST_SUITE(FP16) +FIXTURE_DATA_TEST_CASE(RunSmall, CLRsqrtLayerFixture<half>, 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<half>, 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<float>, 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<float>, 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 <typename TensorType, typename AccessorType, typename FunctionType, typename T> +class ElementWiseUnaryValidationFixture : public framework::Fixture +{ +public: + template <typename...> + 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 <typename U> + 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<TensorType>(shape, data_type); + TensorType dst = create_tensor<TensorType>(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<T> compute_reference(const TensorShape &shape, DataType data_type) + { + // Create reference + SimpleTensor<T> src{ shape, data_type }; + + // Fill reference + fill(src, 0); + + return reference::elementwise_unary<T>(src, _op); + } + + TensorType _target{}; + SimpleTensor<T> _reference{}; + ElementWiseUnary _op{}; +}; + +template <typename TensorType, typename AccessorType, typename FunctionType, typename T> +class RsqrtValidationFixture : public ElementWiseUnaryValidationFixture<TensorType, AccessorType, FunctionType, T> +{ +public: + template <typename...> + void setup(const TensorShape &shape, DataType data_type) + { + ElementWiseUnaryValidationFixture<TensorType, AccessorType, FunctionType, T>::setup(shape, data_type, ElementWiseUnary::RSQRT); + } +}; + +template <typename TensorType, typename AccessorType, typename FunctionType, typename T> +class ExpValidationFixture : public ElementWiseUnaryValidationFixture<TensorType, AccessorType, FunctionType, T> +{ +public: + template <typename...> + void setup(const TensorShape &shape, DataType data_type) + { + ElementWiseUnaryValidationFixture<TensorType, AccessorType, FunctionType, T>::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 <typename T> +SimpleTensor<T> elementwise_unary(const SimpleTensor<T> &src, ElementWiseUnary op) +{ + SimpleTensor<T> 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<float> elementwise_unary(const SimpleTensor<float> &src, ElementWiseUnary op); +template SimpleTensor<half> elementwise_unary(const SimpleTensor<half> &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 <typename T> +SimpleTensor<T> elementwise_unary(const SimpleTensor<T> &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. |