From 5c8e05c6c2726739d9c7ce0b7f6533be5be5fabd Mon Sep 17 00:00:00 2001 From: Michalis Spyrou Date: Thu, 22 Mar 2018 11:56:01 +0000 Subject: COMPMID-1019 Implement copy function CL Change-Id: Ica17528bf6c812d9caf9d66c612c11434ec1dc69 Reviewed-on: https://eu-gerrit-1.euhpc.arm.com/125542 Tested-by: Jenkins Reviewed-by: Georgios Pinitas --- arm_compute/core/CL/CLKernels.h | 1 + arm_compute/core/CL/kernels/CLCopyKernel.h | 63 +++++++++++++++++ arm_compute/runtime/CL/CLFunctions.h | 1 + arm_compute/runtime/CL/functions/CLCopy.h | 48 +++++++++++++ src/core/CL/CLKernelLibrary.cpp | 5 ++ src/core/CL/cl_kernels/copy_tensor.cl | 48 +++++++++++++ src/core/CL/kernels/CLCopyKernel.cpp | 87 +++++++++++++++++++++++ src/runtime/CL/functions/CLCopy.cpp | 43 +++++++++++ tests/validation/CL/Copy.cpp | 98 +++++++++++++++++++++++++ tests/validation/fixtures/CopyFixture.h | 110 +++++++++++++++++++++++++++++ tests/validation/reference/Copy.cpp | 57 +++++++++++++++ tests/validation/reference/Copy.h | 43 +++++++++++ 12 files changed, 604 insertions(+) create mode 100644 arm_compute/core/CL/kernels/CLCopyKernel.h create mode 100644 arm_compute/runtime/CL/functions/CLCopy.h create mode 100644 src/core/CL/cl_kernels/copy_tensor.cl create mode 100644 src/core/CL/kernels/CLCopyKernel.cpp create mode 100644 src/runtime/CL/functions/CLCopy.cpp create mode 100644 tests/validation/CL/Copy.cpp create mode 100644 tests/validation/fixtures/CopyFixture.h create mode 100644 tests/validation/reference/Copy.cpp create mode 100644 tests/validation/reference/Copy.h diff --git a/arm_compute/core/CL/CLKernels.h b/arm_compute/core/CL/CLKernels.h index 6f5c61523f..9481e144df 100644 --- a/arm_compute/core/CL/CLKernels.h +++ b/arm_compute/core/CL/CLKernels.h @@ -42,6 +42,7 @@ #include "arm_compute/core/CL/kernels/CLCol2ImKernel.h" #include "arm_compute/core/CL/kernels/CLColorConvertKernel.h" #include "arm_compute/core/CL/kernels/CLConvolutionKernel.h" +#include "arm_compute/core/CL/kernels/CLCopyKernel.h" #include "arm_compute/core/CL/kernels/CLDeconvolutionLayerUpsampleKernel.h" #include "arm_compute/core/CL/kernels/CLDepthConcatenateLayerKernel.h" #include "arm_compute/core/CL/kernels/CLDepthConvertLayerKernel.h" diff --git a/arm_compute/core/CL/kernels/CLCopyKernel.h b/arm_compute/core/CL/kernels/CLCopyKernel.h new file mode 100644 index 0000000000..40b8203543 --- /dev/null +++ b/arm_compute/core/CL/kernels/CLCopyKernel.h @@ -0,0 +1,63 @@ +/* + * 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_CLCOPYKERNEL_H__ +#define __ARM_COMPUTE_CLCOPYKERNEL_H__ + +#include "arm_compute/core/CL/ICLKernel.h" +#include "arm_compute/core/Types.h" + +namespace arm_compute +{ +class ICLTensor; + +/** OpenCL kernel to perform a copy between two tensors */ +class CLCopyKernel : public ICLKernel +{ +public: + /** Default constructor */ + CLCopyKernel(); + /** Prevent instances of this class from being copied (As this class contains pointers). */ + CLCopyKernel(const CLCopyKernel &) = delete; + /** Prevent instances of this class from being copied (As this class contains pointers). */ + CLCopyKernel &operator=(const CLCopyKernel &) = delete; + /** Allow instances of this class to be moved */ + CLCopyKernel(CLCopyKernel &&) = default; + /** Allow instances of this class to be moved */ + CLCopyKernel &operator=(CLCopyKernel &&) = default; + /** Initialize the kernel's input, output. + * + * @param[in] input Source tensor. Data types supported: U8. + * @param[out] output Destination tensor. Data types supported: U8. + */ + void configure(const ICLTensor *input, ICLTensor *output); + + // Inherited methods overridden: + void run(const Window &window, cl::CommandQueue &queue) override; + +private: + const ICLTensor *_input; + ICLTensor *_output; +}; +} // namespace arm_compute +#endif /*__ARM_COMPUTE_CLCOPYKERNEL_H__ */ diff --git a/arm_compute/runtime/CL/CLFunctions.h b/arm_compute/runtime/CL/CLFunctions.h index adf240e642..a63afaac39 100644 --- a/arm_compute/runtime/CL/CLFunctions.h +++ b/arm_compute/runtime/CL/CLFunctions.h @@ -42,6 +42,7 @@ #include "arm_compute/runtime/CL/functions/CLColorConvert.h" #include "arm_compute/runtime/CL/functions/CLConvolution.h" #include "arm_compute/runtime/CL/functions/CLConvolutionLayer.h" +#include "arm_compute/runtime/CL/functions/CLCopy.h" #include "arm_compute/runtime/CL/functions/CLDeconvolutionLayer.h" #include "arm_compute/runtime/CL/functions/CLDeconvolutionLayerUpsample.h" #include "arm_compute/runtime/CL/functions/CLDepthConcatenateLayer.h" diff --git a/arm_compute/runtime/CL/functions/CLCopy.h b/arm_compute/runtime/CL/functions/CLCopy.h new file mode 100644 index 0000000000..d76f0702af --- /dev/null +++ b/arm_compute/runtime/CL/functions/CLCopy.h @@ -0,0 +1,48 @@ +/* + * 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_CLCOPY_H__ +#define __ARM_COMPUTE_CLCOPY_H__ + +#include "arm_compute/core/Types.h" +#include "arm_compute/runtime/CL/ICLSimpleFunction.h" + +#include + +namespace arm_compute +{ +class ICLTensor; + +class CLCopy : public ICLSimpleFunction +{ +public: + /** Initialise the function's source and destination. + * + * @param[in] input Source tensor. Data types supported: U8/S8/QS8/QASYMM8/U16/S16/QS16/F16/U32/S32/F32. + * @param[out] output Output tensor. Data types supported: Same as @p input. + * + */ + void configure(ICLTensor *input, ICLTensor *output); +}; +} // namespace arm_compute +#endif /*__ARM_COMPUTE_CLCOPY_H__ */ diff --git a/src/core/CL/CLKernelLibrary.cpp b/src/core/CL/CLKernelLibrary.cpp index 740a98bbac..45149c29b8 100644 --- a/src/core/CL/CLKernelLibrary.cpp +++ b/src/core/CL/CLKernelLibrary.cpp @@ -184,6 +184,7 @@ const std::map CLKernelLibrary::_kernel_program_map = { "convolution_separable9x1_static", "convolution9x9.cl" }, { "convert_depth_down", "depth_convert.cl" }, { "convert_depth_up", "depth_convert.cl" }, + { "copy_tensor", "copy_tensor.cl" }, { "copy_plane", "channel_extract.cl" }, { "copy_planes_3p", "channel_combine.cl" }, { "copy_to_keypoint", "fast_corners.cl" }, @@ -436,6 +437,10 @@ const std::map CLKernelLibrary::_program_source_map = { "convolution_rectangle.cl", #include "./cl_kernels/convolution_rectangle.clembed" + }, + { + "copy_tensor.cl", +#include "./cl_kernels/copy_tensor.clembed" }, { "deconvolution_layer.cl", diff --git a/src/core/CL/cl_kernels/copy_tensor.cl b/src/core/CL/cl_kernels/copy_tensor.cl new file mode 100644 index 0000000000..4b37decf21 --- /dev/null +++ b/src/core/CL/cl_kernels/copy_tensor.cl @@ -0,0 +1,48 @@ +/* + * 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" + +/** Performs a copy of input tensor to the output tensor. + * + * @param[in] in_ptr Pointer to the source image. Supported data types: U8. + * @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: U8. + * @param[in] out_stride_x Stride of the destination image in X dimension (in bytes) + * @param[in] out_step_x out_stride_x * number of elements along X 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 copy_tensor( + VECTOR_DECLARATION(in), + VECTOR_DECLARATION(out)) +{ + Vector in = CONVERT_TO_VECTOR_STRUCT(in); + Vector out = CONVERT_TO_VECTOR_STRUCT(out); + + VEC_DATA_TYPE(DATA_TYPE, 16) + data = vload16(0, (__global DATA_TYPE *)in.ptr); + + vstore16(data, 0, (__global DATA_TYPE *)out.ptr); +} \ No newline at end of file diff --git a/src/core/CL/kernels/CLCopyKernel.cpp b/src/core/CL/kernels/CLCopyKernel.cpp new file mode 100644 index 0000000000..4f00ef9eef --- /dev/null +++ b/src/core/CL/kernels/CLCopyKernel.cpp @@ -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 "arm_compute/core/CL/kernels/CLCopyKernel.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/ICLTensor.h" +#include "arm_compute/core/Error.h" +#include "arm_compute/core/Helpers.h" +#include "arm_compute/core/TensorInfo.h" +#include "arm_compute/core/Validate.h" +#include "arm_compute/core/Window.h" + +#include + +using namespace arm_compute; + +CLCopyKernel::CLCopyKernel() + : _input(nullptr), _output(nullptr) +{ +} + +void CLCopyKernel::configure(const ICLTensor *input, ICLTensor *output) +{ + ARM_COMPUTE_ERROR_ON_MISMATCHING_DIMENSIONS(input->info()->tensor_shape(), output->info()->tensor_shape()); + ARM_COMPUTE_ERROR_ON_MISMATCHING_DATA_TYPES(input, output); + + _input = input; + _output = output; + + // Create kernel + CLBuildOptions build_opts; + build_opts.add_option("-DDATA_TYPE=" + get_cl_type_from_data_type(input->info()->data_type())); + _kernel = static_cast(CLKernelLibrary::get().create_kernel("copy_tensor", build_opts.options())); + + // Configure window + constexpr unsigned int num_elems_processed_per_iteration = 16; + + Window win = calculate_max_window(*input->info(), Steps(num_elems_processed_per_iteration)); + + AccessWindowHorizontal input_access(input->info(), 0, num_elems_processed_per_iteration); + AccessWindowHorizontal output_access(output->info(), 0, num_elems_processed_per_iteration); + + update_window_and_padding(win, input_access, output_access); + + ICLKernel::configure(win); +} + +void CLCopyKernel::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); + Window slice = collapsed.first_slice_window_1D(); + + do + { + unsigned int idx = 0; + add_1D_tensor_argument(idx, _input, slice); + add_1D_tensor_argument(idx, _output, slice); + enqueue(queue, *this, slice); + } + while(collapsed.slide_window_slice_1D(slice)); +} diff --git a/src/runtime/CL/functions/CLCopy.cpp b/src/runtime/CL/functions/CLCopy.cpp new file mode 100644 index 0000000000..3442e3781f --- /dev/null +++ b/src/runtime/CL/functions/CLCopy.cpp @@ -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. + */ +#include "arm_compute/runtime/CL/functions/CLCopy.h" + +#include "arm_compute/core/CL/ICLTensor.h" +#include "arm_compute/core/CL/kernels/CLCopyKernel.h" +#include "arm_compute/core/Error.h" +#include "arm_compute/core/PixelValue.h" +#include "arm_compute/core/TensorInfo.h" +#include "arm_compute/core/Validate.h" +#include "support/ToolchainSupport.h" + +#include + +using namespace arm_compute; + +void CLCopy::configure(ICLTensor *input, ICLTensor *output) +{ + auto k = arm_compute::support::cpp14::make_unique(); + k->configure(input, output); + _kernel = std::move(k); +} diff --git a/tests/validation/CL/Copy.cpp b/tests/validation/CL/Copy.cpp new file mode 100644 index 0000000000..033f7a65a2 --- /dev/null +++ b/tests/validation/CL/Copy.cpp @@ -0,0 +1,98 @@ +/* + * 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/CLCopy.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/CopyFixture.h" + +namespace arm_compute +{ +namespace test +{ +namespace validation +{ +TEST_SUITE(CL) +TEST_SUITE(Copy) + +template +using CLCopyFixture = CopyFixture; + +TEST_SUITE(F32) +FIXTURE_DATA_TEST_CASE(RunSmall, CLCopyFixture, framework::DatasetMode::PRECOMMIT, combine(zip(datasets::SmallShapes(), datasets::SmallShapes()), framework::dataset::make("DataType", + DataType::F32))) +{ + // Validate output + validate(CLAccessor(_target), _reference); +} +FIXTURE_DATA_TEST_CASE(RunLarge, CLCopyFixture, framework::DatasetMode::NIGHTLY, combine(zip(datasets::LargeShapes(), datasets::LargeShapes()), framework::dataset::make("DataType", + DataType::F32))) +{ + // Validate output + validate(CLAccessor(_target), _reference); +} +TEST_SUITE_END() // F32 + +TEST_SUITE(U8) +FIXTURE_DATA_TEST_CASE(RunSmall, CLCopyFixture, framework::DatasetMode::PRECOMMIT, combine(zip(datasets::SmallShapes(), datasets::SmallShapes()), framework::dataset::make("DataType", + DataType::U8))) +{ + // Validate output + validate(CLAccessor(_target), _reference); +} +FIXTURE_DATA_TEST_CASE(RunLarge, CLCopyFixture, framework::DatasetMode::NIGHTLY, combine(zip(datasets::LargeShapes(), datasets::LargeShapes()), framework::dataset::make("DataType", + DataType::U8))) +{ + // Validate output + validate(CLAccessor(_target), _reference); +} +TEST_SUITE_END() // U8 + +TEST_SUITE(U16) +FIXTURE_DATA_TEST_CASE(RunSmall, CLCopyFixture, framework::DatasetMode::PRECOMMIT, combine(zip(datasets::SmallShapes(), datasets::SmallShapes()), framework::dataset::make("DataType", + DataType::U16))) +{ + // Validate output + validate(CLAccessor(_target), _reference); +} +FIXTURE_DATA_TEST_CASE(RunLarge, CLCopyFixture, framework::DatasetMode::NIGHTLY, combine(zip(datasets::LargeShapes(), datasets::LargeShapes()), framework::dataset::make("DataType", + DataType::U16))) +{ + // Validate output + validate(CLAccessor(_target), _reference); +} +TEST_SUITE_END() // U16 + +TEST_SUITE_END() // Copy +TEST_SUITE_END() // CL +} // namespace validation +} // namespace test +} // namespace arm_compute diff --git a/tests/validation/fixtures/CopyFixture.h b/tests/validation/fixtures/CopyFixture.h new file mode 100644 index 0000000000..911d908e53 --- /dev/null +++ b/tests/validation/fixtures/CopyFixture.h @@ -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. + */ +#ifndef ARM_COMPUTE_TEST_COPY_FIXTURE +#define ARM_COMPUTE_TEST_COPY_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/Copy.h" + +namespace arm_compute +{ +namespace test +{ +namespace validation +{ +template +class CopyFixture : public framework::Fixture +{ +public: + template + void setup(TensorShape input_shape, TensorShape output_shape, DataType data_type) + { + _target = compute_target(input_shape, output_shape, data_type); + _reference = compute_reference(input_shape, output_shape, data_type); + } + +protected: + template + void fill(U &&tensor, int i) + { + library->fill_tensor_uniform(tensor, i); + } + + TensorType compute_target(const TensorShape &input_shape, const TensorShape &output_shape, DataType data_type) + { + // Check if indeed the input shape can be reshape to the output one + ARM_COMPUTE_EXPECT(input_shape.total_size() == output_shape.total_size(), framework::LogLevel::ERRORS); + + // Create tensors + TensorType src = create_tensor(input_shape, data_type); + TensorType dst = create_tensor(output_shape, data_type); + + // Create and configure function + FunctionType copy; + + copy.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 + copy.run(); + + return dst; + } + + SimpleTensor compute_reference(const TensorShape &input_shape, const TensorShape &output_shape, DataType data_type) + { + // Create reference + SimpleTensor src{ input_shape, data_type }; + + // Fill reference + fill(src, 0); + + return reference::copy(src, output_shape); + } + + TensorType _target{}; + SimpleTensor _reference{}; +}; +} // namespace validation +} // namespace test +} // namespace arm_compute +#endif /* ARM_COMPUTE_TEST_COPY_FIXTURE */ diff --git a/tests/validation/reference/Copy.cpp b/tests/validation/reference/Copy.cpp new file mode 100644 index 0000000000..dc519a4e45 --- /dev/null +++ b/tests/validation/reference/Copy.cpp @@ -0,0 +1,57 @@ +/* + * 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 "Copy.h" + +#include "arm_compute/core/Types.h" + +namespace arm_compute +{ +namespace test +{ +namespace validation +{ +namespace reference +{ +template +SimpleTensor copy(const SimpleTensor &src, const TensorShape &output_shape) +{ + ARM_COMPUTE_ERROR_ON_MISMATCHING_DIMENSIONS(src.shape(), output_shape); + + SimpleTensor dst(output_shape, src.data_type()); + std::copy_n(src.data(), src.num_elements(), dst.data()); + return dst; +} + +template SimpleTensor copy(const SimpleTensor &src, const TensorShape &output_shape); +template SimpleTensor copy(const SimpleTensor &src, const TensorShape &output_shape); +template SimpleTensor copy(const SimpleTensor &src, const TensorShape &output_shape); +template SimpleTensor copy(const SimpleTensor &src, const TensorShape &output_shape); +template SimpleTensor copy(const SimpleTensor &src, const TensorShape &output_shape); +template SimpleTensor copy(const SimpleTensor &src, const TensorShape &output_shape); +template SimpleTensor copy(const SimpleTensor &src, const TensorShape &output_shape); +template SimpleTensor copy(const SimpleTensor &src, const TensorShape &output_shape); +} // namespace reference +} // namespace validation +} // namespace test +} // namespace arm_compute diff --git a/tests/validation/reference/Copy.h b/tests/validation/reference/Copy.h new file mode 100644 index 0000000000..362af0319e --- /dev/null +++ b/tests/validation/reference/Copy.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_COPY_H__ +#define __ARM_COMPUTE_TEST_COPY_H__ + +#include "tests/SimpleTensor.h" + +namespace arm_compute +{ +namespace test +{ +namespace validation +{ +namespace reference +{ +template +SimpleTensor copy(const SimpleTensor &src, const TensorShape &output_shape); +} // namespace reference +} // namespace validation +} // namespace test +} // namespace arm_compute +#endif /* __ARM_COMPUTE_TEST_COPY_H__ */ -- cgit v1.2.1