From 17a51732acfd46017c3d45ee2583e405cb8b5034 Mon Sep 17 00:00:00 2001 From: =?utf8?q?Prasanna=20R/System=20SW=20/SRI-Bangalore/Engineer/?= =?utf8?q?=EC=82=BC=EC=84=B1=EC=A0=84=EC=9E=90?= Date: Wed, 17 Oct 2018 16:06:03 +0530 Subject: [PATCH] Add CL Kernels for Equal op (#3029) This patch adds CL Kernels for Equal op. Signed-off-by: prasannar --- .../arm_compute/core/CL/kernels/CLEqualKernel.h | 58 +++++++++++ .../arm_compute/runtime/CL/functions/CLEqual.h | 44 +++++++++ libs/ARMComputeEx/src/core/CL/CLKernelLibrary.cpp | 10 ++ libs/ARMComputeEx/src/core/CL/cl_kernels/equal.cl | 64 ++++++++++++ .../src/core/CL/cl_kernels/equal_quantized.cl | 75 ++++++++++++++ .../src/core/CL/kernels/CLEqualKernel.cpp | 108 +++++++++++++++++++++ .../src/runtime/CL/functions/CLEqual.cpp | 30 ++++++ 7 files changed, 389 insertions(+) create mode 100644 libs/ARMComputeEx/arm_compute/core/CL/kernels/CLEqualKernel.h create mode 100644 libs/ARMComputeEx/arm_compute/runtime/CL/functions/CLEqual.h create mode 100644 libs/ARMComputeEx/src/core/CL/cl_kernels/equal.cl create mode 100644 libs/ARMComputeEx/src/core/CL/cl_kernels/equal_quantized.cl create mode 100644 libs/ARMComputeEx/src/core/CL/kernels/CLEqualKernel.cpp create mode 100644 libs/ARMComputeEx/src/runtime/CL/functions/CLEqual.cpp diff --git a/libs/ARMComputeEx/arm_compute/core/CL/kernels/CLEqualKernel.h b/libs/ARMComputeEx/arm_compute/core/CL/kernels/CLEqualKernel.h new file mode 100644 index 0000000..7e366d2 --- /dev/null +++ b/libs/ARMComputeEx/arm_compute/core/CL/kernels/CLEqualKernel.h @@ -0,0 +1,58 @@ +/* + * Copyright (c) 2018 Samsung Electronics Co., Ltd. All Rights Reserved + * Copyright (c) 2016-2018 ARM Limited. + * + * Licensed under the Apache License, Version 2.0 (the "License"); + * you may not use this file except in compliance with the License. + * You may obtain a copy of the License at + * + * http://www.apache.org/licenses/LICENSE-2.0 + * + * Unless required by applicable law or agreed to in writing, software + * distributed under the License is distributed on an "AS IS" BASIS, + * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. + * See the License for the specific language governing permissions and + * limitations under the License. + */ +#ifndef __ARM_COMPUTE_CLEQUALKERNEL_H__ +#define __ARM_COMPUTE_CLEQUALKERNEL_H__ + +#include "arm_compute/core/CL/ICLKernel.h" +#include "arm_compute/core/Types.h" + +namespace arm_compute +{ +class ICLTensor; + +/** OpenCL kernel to check if values in both tensors are equal*/ +class CLEqualKernel : public ICLKernel +{ +public: + /** Default constructor */ + CLEqualKernel(); + /** Prevent instances of this class from being copied (As this class contains pointers). */ + CLEqualKernel(const CLEqualKernel &) = delete; + /** Prevent instances of this class from being copied (As this class contains pointers). */ + CLEqualKernel &operator=(const CLEqualKernel &) = delete; + /** Allow instances of this class to be moved */ + CLEqualKernel(CLEqualKernel &&) = default; + /** Allow instances of this class to be moved */ + CLEqualKernel &operator=(CLEqualKernel &&) = default; + /** Initialize the kernel's input, output. + * + * @param[in] input1 Source tensor1. + * @param[in] input2 Source tensor2. + * @param[out] output Output tensor. + */ + void configure(const ICLTensor *input1, const ICLTensor *input2, ICLTensor *output); + + // Inherited methods overridden: + void run(const Window &window, cl::CommandQueue &queue) override; + +private: + const ICLTensor *_input1; + const ICLTensor *_input2; + ICLTensor *_output; +}; +} // namespace arm_compute +#endif /*__ARM_COMPUTE_CLEQUALKERNEL_H__ */ diff --git a/libs/ARMComputeEx/arm_compute/runtime/CL/functions/CLEqual.h b/libs/ARMComputeEx/arm_compute/runtime/CL/functions/CLEqual.h new file mode 100644 index 0000000..e49bbac --- /dev/null +++ b/libs/ARMComputeEx/arm_compute/runtime/CL/functions/CLEqual.h @@ -0,0 +1,44 @@ +/* + * Copyright (c) 2018 Samsung Electronics Co., Ltd. All Rights Reserved + * Copyright (c) 2016-2018 ARM Limited. + * + * Licensed under the Apache License, Version 2.0 (the "License"); + * you may not use this file except in compliance with the License. + * You may obtain a copy of the License at + * + * http://www.apache.org/licenses/LICENSE-2.0 + * + * Unless required by applicable law or agreed to in writing, software + * distributed under the License is distributed on an "AS IS" BASIS, + * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. + * See the License for the specific language governing permissions and + * limitations under the License. + */ +#ifndef __ARM_COMPUTE_CLEQUAL_H__ +#define __ARM_COMPUTE_CLEQUAL_H__ + +#include "arm_compute/core/Types.h" +#include "arm_compute/runtime/CL/ICLSimpleFunction.h" + +#include + +namespace arm_compute +{ +class ICLTensor; + +class CLEqual : public ICLSimpleFunction +{ +public: + /** Initialise the function's source and destination. + * + * @param[in] input1 Source tensor1. Data types supported: + * U8/S8/QS8/QASYMM8/U16/S16/QS16/F16/U32/S32/F32. + * @param[in] input2 Source tensor2. 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 *input1, ICLTensor *input2, ICLTensor *output); +}; +} // namespace arm_compute +#endif /*__ARM_COMPUTE_CLEQUAL_H__ */ diff --git a/libs/ARMComputeEx/src/core/CL/CLKernelLibrary.cpp b/libs/ARMComputeEx/src/core/CL/CLKernelLibrary.cpp index ccda8fc..8b3f6a0 100644 --- a/libs/ARMComputeEx/src/core/CL/CLKernelLibrary.cpp +++ b/libs/ARMComputeEx/src/core/CL/CLKernelLibrary.cpp @@ -117,6 +117,8 @@ const std::map CLKernelLibraryEx::_kernel_program_map {"direct_convolution5x5", "direct_convolution5x5.cl"}, {"direct_convolution5x5_f32_bifrost", "direct_convolution5x5.cl"}, {"direct_convolution_1x1_3x3_5x5_quantized", "direct_convolution_1x1_3x3_5x5_quantized.cl"}, + {"equal", "equal.cl"}, + {"equal_quantized", "equal_quantized.cl"}, {"erode", "erode.cl"}, {"exp_layer", "exp.cl"}, {"fast_corners", "fast_corners.cl"}, @@ -319,6 +321,14 @@ const std::map CLKernelLibraryEx::_program_source_map #include "./cl_kernels/cast.clembed" }, { + "equal.cl", +#include "./cl_kernels/equal.clembed" + }, + { + "equal_quantized.cl", +#include "./cl_kernels/equal_quantized.clembed" + }, + { "exp.cl", #include "./cl_kernels/exp.clembed" }, diff --git a/libs/ARMComputeEx/src/core/CL/cl_kernels/equal.cl b/libs/ARMComputeEx/src/core/CL/cl_kernels/equal.cl new file mode 100644 index 0000000..9bd754f --- /dev/null +++ b/libs/ARMComputeEx/src/core/CL/cl_kernels/equal.cl @@ -0,0 +1,64 @@ +/* + * Copyright (c) 2018 Samsung Electronics Co., Ltd. All Rights Reserved + * Copyright (c) 2016-2018 ARM Limited. + * + * Licensed under the Apache License, Version 2.0 (the "License"); + * you may not use this file except in compliance with the License. + * You may obtain a copy of the License at + * + * http://www.apache.org/licenses/LICENSE-2.0 + * + * Unless required by applicable law or agreed to in writing, software + * distributed under the License is distributed on an "AS IS" BASIS, + * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. + * See the License for the specific language governing permissions and + * limitations under the License. + */ +#include "helpers.h" + +/** Checks if values in both tensors are equal. + * + * @attention Data type can be passed using the -DDATA_TYPE_IN compile flag, e.g. -DDATA_TYPE_IN=float + * @attention Vector size should be given as a preprocessor argument using -DVEC_SIZE=size. e.g. -DVEC_SIZE=16 + * @note Can only take floating point data types. + * + * @param[in] input1_ptr Pointer to the source tensor. + * @param[in] input1_stride_x Stride of the source tensor in X dimension (in bytes) + * @param[in] input1_step_x input1_stride_x * number of elements along X processed per workitem(in bytes) + * @param[in] input1_stride_y Stride of the source tensor in Y dimension (in bytes) + * @param[in] input1_step_y input1_stride_y * number of elements along Y processed per workitem(in bytes) + * @param[in] input1_stride_z Stride of the source tensor in Z dimension (in bytes) + * @param[in] input1_step_z input1_stride_z * number of elements along Z processed per workitem(in bytes) + * @param[in] input1_offset_first_element_in_bytes The offset of the first element in the source tensor + * + * @param[in] input2_ptr Pointer to the source tensor. + * @param[in] input2_stride_x Stride of the source tensor in X dimension (in bytes) + * @param[in] input2_step_x input2_stride_x * number of elements along X processed per workitem(in bytes) + * @param[in] input2_stride_y Stride of the source tensor in Y dimension (in bytes) + * @param[in] input2_step_y input2_stride_y * number of elements along Y processed per workitem(in bytes) + * @param[in] input2_stride_z Stride of the source tensor in Z dimension (in bytes) + * @param[in] input2_step_z input2_stride_z * number of elements along Z processed per workitem(in bytes) + * @param[in] input2_offset_first_element_in_bytes The offset of the first element in the source tensor + * + * @param[out] output_ptr Pointer to the destination tensor. + * @param[in] output_stride_x Stride of the destination tensor in X dimension (in bytes) + * @param[in] output_step_x output_stride_x * number of elements along X processed per workitem(in bytes) + * @param[in] output_stride_y Stride of the destination tensor in Y dimension (in bytes) + * @param[in] output_step_y output_stride_y * number of elements along Y processed per workitem(in bytes) + * @param[in] output_stride_z Stride of the destination tensor in Z dimension (in bytes) + * @param[in] output_step_z output_stride_z * number of elements along Z processed per workitem(in bytes) + */ +__kernel void equal( + TENSOR3D_DECLARATION(input1), + TENSOR3D_DECLARATION(input2), + TENSOR3D_DECLARATION(output)) +{ + Tensor3D input1 = CONVERT_TO_TENSOR3D_STRUCT(input1); + Tensor3D input2 = CONVERT_TO_TENSOR3D_STRUCT(input2); + Tensor3D output = CONVERT_TO_TENSOR3D_STRUCT(output); + + VSTORE(VEC_SIZE) + (CONVERT(VLOAD(VEC_SIZE)(0, (__global DATA_TYPE_IN *)input1.ptr) == VLOAD(VEC_SIZE)(0, (__global DATA_TYPE_IN *)input2.ptr) ? 1 : 0, + VEC_DATA_TYPE(DATA_TYPE_OUT, VEC_SIZE)), + 0, (__global DATA_TYPE_OUT *)output.ptr); +} diff --git a/libs/ARMComputeEx/src/core/CL/cl_kernels/equal_quantized.cl b/libs/ARMComputeEx/src/core/CL/cl_kernels/equal_quantized.cl new file mode 100644 index 0000000..1cffc86 --- /dev/null +++ b/libs/ARMComputeEx/src/core/CL/cl_kernels/equal_quantized.cl @@ -0,0 +1,75 @@ +/* + * Copyright (c) 2018 Samsung Electronics Co., Ltd. All Rights Reserved + * Copyright (c) 2016-2018 ARM Limited. + * + * Licensed under the Apache License, Version 2.0 (the "License"); + * you may not use this file except in compliance with the License. + * You may obtain a copy of the License at + * + * http://www.apache.org/licenses/LICENSE-2.0 + * + * Unless required by applicable law or agreed to in writing, software + * distributed under the License is distributed on an "AS IS" BASIS, + * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. + * See the License for the specific language governing permissions and + * limitations under the License. + */ +#include "helpers.h" +#define SUB(x, y) (x) - (y) + +/** Checks if values in both tensors are equal. + * + * @attention Data type can be passed using the -DDATA_TYPE_IN compile flag, e.g. -DDATA_TYPE_IN=float + * @attention Vector size should be given as a preprocessor argument using -DVEC_SIZE=size. e.g. -DVEC_SIZE=16 + * @note Can only take floating point data types. + * + * @param[in] input1_ptr Pointer to the source tensor. + * @param[in] input1_stride_x Stride of the source tensor in X dimension (in bytes) + * @param[in] input1_step_x input1_stride_x * number of elements along X processed per workitem(in bytes) + * @param[in] input1_stride_y Stride of the source tensor in Y dimension (in bytes) + * @param[in] input1_step_y input1_stride_y * number of elements along Y processed per workitem(in bytes) + * @param[in] input1_stride_z Stride of the source tensor in Z dimension (in bytes) + * @param[in] input1_step_z input1_stride_z * number of elements along Z processed per workitem(in bytes) + * @param[in] input1_offset_first_element_in_bytes The offset of the first element in the source tensor + * + * @param[in] input2_ptr Pointer to the source tensor. + * @param[in] input2_stride_x Stride of the source tensor in X dimension (in bytes) + * @param[in] input2_step_x input2_stride_x * number of elements along X processed per workitem(in bytes) + * @param[in] input2_stride_y Stride of the source tensor in Y dimension (in bytes) + * @param[in] input2_step_y input2_stride_y * number of elements along Y processed per workitem(in bytes) + * @param[in] input2_stride_z Stride of the source tensor in Z dimension (in bytes) + * @param[in] input2_step_z input2_stride_z * number of elements along Z processed per workitem(in bytes) + * @param[in] input2_offset_first_element_in_bytes The offset of the first element in the source tensor + * + * @param[out] output_ptr Pointer to the destination tensor. + * @param[in] output_stride_x Stride of the destination tensor in X dimension (in bytes) + * @param[in] output_step_x output_stride_x * number of elements along X processed per workitem(in bytes) + * @param[in] output_stride_y Stride of the destination tensor in Y dimension (in bytes) + * @param[in] output_step_y output_stride_y * number of elements along Y processed per workitem(in bytes) + * @param[in] output_stride_z Stride of the destination tensor in Z dimension (in bytes) + * @param[in] output_step_z output_stride_z * number of elements along Z processed per workitem(in bytes) + * @param[in] output_offset_first_element_in_bytes The offset of the first element in the destination tensor + */ +__kernel void equal_quantized( + TENSOR3D_DECLARATION(in1), + TENSOR3D_DECLARATION(in2), + TENSOR3D_DECLARATION(out)) +{ + // Get pixels pointer + Tensor3D in1 = CONVERT_TO_TENSOR3D_STRUCT(in1); + Tensor3D in2 = CONVERT_TO_TENSOR3D_STRUCT(in2); + Tensor3D out = CONVERT_TO_TENSOR3D_STRUCT(out); + + int16 in_a = CONVERT(vload16(0, (__global uchar *)in1.ptr), int16); + int16 in_b = CONVERT(vload16(0, (__global uchar *)in2.ptr), int16); + + in_a = SUB(in_a, (int16)((int)OFFSET_IN1)); + in_b = SUB(in_b, (int16)((int)OFFSET_IN2)); + + const float16 in1f32 = convert_float16(in_a) * (float16)((float)SCALE_IN1); + const float16 in2f32 = convert_float16(in_b) * (float16)((float)SCALE_IN2); + const float16 qresf32 = convert_float16((in1f32 == in2f32) ? 1 : 0); + const uchar16 res = convert_uchar16_sat(convert_int16_rte(qresf32)); + // Store result + vstore16(res, 0, (__global uchar *)out.ptr); +} diff --git a/libs/ARMComputeEx/src/core/CL/kernels/CLEqualKernel.cpp b/libs/ARMComputeEx/src/core/CL/kernels/CLEqualKernel.cpp new file mode 100644 index 0000000..abe4c40 --- /dev/null +++ b/libs/ARMComputeEx/src/core/CL/kernels/CLEqualKernel.cpp @@ -0,0 +1,108 @@ +/* + * Copyright (c) 2018 Samsung Electronics Co., Ltd. All Rights Reserved + * Copyright (c) 2016-2018 ARM Limited. + * + * Licensed under the Apache License, Version 2.0 (the "License"); + * you may not use this file except in compliance with the License. + * You may obtain a copy of the License at + * + * http://www.apache.org/licenses/LICENSE-2.0 + * + * Unless required by applicable law or agreed to in writing, software + * distributed under the License is distributed on an "AS IS" BASIS, + * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. + * See the License for the specific language governing permissions and + * limitations under the License. + */ +#include "arm_compute/core/CL/kernels/CLEqualKernel.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/CLKernelLibraryEx.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; + +CLEqualKernel::CLEqualKernel() : _input1(nullptr), _input2(nullptr), _output(nullptr) {} + +void CLEqualKernel::configure(const ICLTensor *input1, const ICLTensor *input2, ICLTensor *output) +{ + ARM_COMPUTE_ERROR_ON_MISMATCHING_DIMENSIONS(input1->info()->tensor_shape(), + input2->info()->tensor_shape()); + ARM_COMPUTE_ERROR_ON_MISMATCHING_DIMENSIONS(input1->info()->tensor_shape(), + output->info()->tensor_shape()); + ARM_COMPUTE_ERROR_ON_MISMATCHING_DATA_TYPES(input1, input2); + ARM_COMPUTE_ERROR_ON_MISMATCHING_DATA_TYPES(input1, output); + + _input1 = input1; + _input2 = input2; + _output = output; + + constexpr unsigned int num_elems_processed_per_iteration = 16; + + // Create kernel + std::string kernel_name = "equal"; + std::set build_opts; + build_opts.emplace(("-DDATA_TYPE_IN=" + get_cl_type_from_data_type(input1->info()->data_type()))); + build_opts.emplace( + ("-DDATA_TYPE_OUT=" + get_cl_type_from_data_type(output->info()->data_type()))); + build_opts.emplace( + ("-DVEC_SIZE=" + support::cpp11::to_string(num_elems_processed_per_iteration))); + if (is_data_type_quantized_asymmetric(input1->info()->data_type())) + { + build_opts.emplace("-DOFFSET_IN1=" + + support::cpp11::to_string(input1->info()->quantization_info().offset)); + build_opts.emplace("-DOFFSET_IN2=" + + support::cpp11::to_string(input2->info()->quantization_info().offset)); + build_opts.emplace("-DSCALE_IN1=" + + support::cpp11::to_string(input1->info()->quantization_info().scale)); + build_opts.emplace("-DSCALE_IN2=" + + support::cpp11::to_string(input2->info()->quantization_info().scale)); + kernel_name += "_quantized"; + } + + _kernel = + static_cast(CLKernelLibraryEx::get().create_kernel(kernel_name, build_opts)); + + // Configure window + Window win = calculate_max_window(*input1->info(), Steps(num_elems_processed_per_iteration)); + + AccessWindowHorizontal input1_access(input1->info(), 0, num_elems_processed_per_iteration); + AccessWindowHorizontal input2_access(input2->info(), 0, num_elems_processed_per_iteration); + AccessWindowHorizontal output_access(output->info(), 0, num_elems_processed_per_iteration); + + ValidRegion valid_region = + intersect_valid_regions(input1->info()->valid_region(), input2->info()->valid_region()); + + update_window_and_padding(win, input1_access, input2_access, output_access); + + output_access.set_valid_region(win, valid_region); + + ICLKernel::configure(win); +} + +void CLEqualKernel::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::DimZ); + Window slice = collapsed.first_slice_window_3D(); + + do + { + unsigned int idx = 0; + add_3D_tensor_argument(idx, _input1, slice); + add_3D_tensor_argument(idx, _input2, slice); + add_3D_tensor_argument(idx, _output, slice); + enqueue(queue, *this, slice); + } while (collapsed.slide_window_slice_3D(slice)); +} diff --git a/libs/ARMComputeEx/src/runtime/CL/functions/CLEqual.cpp b/libs/ARMComputeEx/src/runtime/CL/functions/CLEqual.cpp new file mode 100644 index 0000000..16553e3 --- /dev/null +++ b/libs/ARMComputeEx/src/runtime/CL/functions/CLEqual.cpp @@ -0,0 +1,30 @@ +/* + * Copyright (c) 2018 Samsung Electronics Co., Ltd. All Rights Reserved + * Copyright (c) 2016-2018 ARM Limited. + * + * Licensed under the Apache License, Version 2.0 (the "License"); + * you may not use this file except in compliance with the License. + * You may obtain a copy of the License at + * + * http://www.apache.org/licenses/LICENSE-2.0 + * + * Unless required by applicable law or agreed to in writing, software + * distributed under the License is distributed on an "AS IS" BASIS, + * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. + * See the License for the specific language governing permissions and + * limitations under the License. + */ +#include "arm_compute/runtime/CL/functions/CLEqual.h" + +#include "arm_compute/core/CL/kernels/CLEqualKernel.h" + +#include + +using namespace arm_compute; + +void CLEqual::configure(ICLTensor *input1, ICLTensor *input2, ICLTensor *output) +{ + auto k = arm_compute::support::cpp14::make_unique(); + k->configure(input1, input2, output); + _kernel = std::move(k); +} -- 2.7.4