From: Prasanna R/System SW /SRI-Bangalore/Engineer/삼성전자 Date: Wed, 10 Oct 2018 08:04:59 +0000 (+0530) Subject: Add CL Kernels for SquaredDifference op. (#3037) X-Git-Tag: 0.3~685 X-Git-Url: http://review.tizen.org/git/?a=commitdiff_plain;h=e4adbf0822f132584bf8f55ed85d81de8e685c49;p=platform%2Fcore%2Fml%2Fnnfw.git Add CL Kernels for SquaredDifference op. (#3037) This patch adds CL Kernels for SquaredDifference op. Signed-off-by: prasannar --- diff --git a/libs/ARMComputeEx/arm_compute/core/CL/kernels/CLSquaredDifferenceKernel.h b/libs/ARMComputeEx/arm_compute/core/CL/kernels/CLSquaredDifferenceKernel.h new file mode 100644 index 0000000..9ee7a21 --- /dev/null +++ b/libs/ARMComputeEx/arm_compute/core/CL/kernels/CLSquaredDifferenceKernel.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_CLSQUARED_DIFFERENCE_KERNEL_H__ +#define __ARM_COMPUTE_CLSQUARED_DIFFERENCE_KERNEL_H__ + +#include "arm_compute/core/CL/ICLKernel.h" +#include "arm_compute/core/Types.h" + +namespace arm_compute +{ +class ICLTensor; + +/** OpenCL kernel to return squared difference value of two tensors (x-y)^2*/ +class CLSquaredDifferenceKernel : public ICLKernel +{ +public: + /** Default constructor */ + CLSquaredDifferenceKernel(); + /** Prevent instances of this class from being copied (As this class contains pointers). */ + CLSquaredDifferenceKernel(const CLSquaredDifferenceKernel &) = delete; + /** Prevent instances of this class from being copied (As this class contains pointers). */ + CLSquaredDifferenceKernel &operator=(const CLSquaredDifferenceKernel &) = delete; + /** Allow instances of this class to be moved */ + CLSquaredDifferenceKernel(CLSquaredDifferenceKernel &&) = default; + /** Allow instances of this class to be moved */ + CLSquaredDifferenceKernel &operator=(CLSquaredDifferenceKernel &&) = 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_CLSQUARED_DIFFERENCE_KERNEL_H__ */ diff --git a/libs/ARMComputeEx/arm_compute/runtime/CL/functions/CLSquaredDifference.h b/libs/ARMComputeEx/arm_compute/runtime/CL/functions/CLSquaredDifference.h new file mode 100644 index 0000000..8b4905b --- /dev/null +++ b/libs/ARMComputeEx/arm_compute/runtime/CL/functions/CLSquaredDifference.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_CLSQUARED_DIFFERENCE_H__ +#define __ARM_COMPUTE_CLSQUARED_DIFFERENCE_H__ + +#include "arm_compute/core/Types.h" +#include "arm_compute/runtime/CL/ICLSimpleFunction.h" + +#include + +namespace arm_compute +{ +class ICLTensor; + +class CLSquaredDifference : 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_CLSQUARED_DIFFERENCESQUARED_DIFFERENCE*/ diff --git a/libs/ARMComputeEx/src/core/CL/CLKernelLibrary.cpp b/libs/ARMComputeEx/src/core/CL/CLKernelLibrary.cpp index 8ba01af..083cc86 100644 --- a/libs/ARMComputeEx/src/core/CL/CLKernelLibrary.cpp +++ b/libs/ARMComputeEx/src/core/CL/CLKernelLibrary.cpp @@ -262,6 +262,7 @@ const std::map CLKernelLibraryEx::_kernel_program_map {"softmax_layer_max_shift_exp_sum_quantized_parallel", "softmax_layer_quantized.cl"}, {"softmax_layer_max_shift_exp_sum_serial", "softmax_layer.cl"}, {"softmax_layer_max_shift_exp_sum_parallel", "softmax_layer.cl"}, + {"squared_difference", "squared_difference.cl"}, {"strided_slice", "strided_slice.cl"}, {"suppress_non_maximum", "canny.cl"}, {"tablelookup_U8", "tablelookup.cl"}, @@ -357,6 +358,10 @@ const std::map CLKernelLibraryEx::_program_source_map #include "./cl_kernels/reduction_mean.clembed" }, { + "squared_difference.cl", +#include "./cl_kernels/squared_difference.clembed" + }, + { "strided_slice.cl", #include "./cl_kernels/strided_slice.clembed" }, diff --git a/libs/ARMComputeEx/src/core/CL/cl_kernels/squared_difference.cl b/libs/ARMComputeEx/src/core/CL/cl_kernels/squared_difference.cl new file mode 100644 index 0000000..7c12626 --- /dev/null +++ b/libs/ARMComputeEx/src/core/CL/cl_kernels/squared_difference.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" + +/** Returns true value of squared_difference of two tensors. + * + * @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 image. Supported data types: F16/F32 + * @param[in] input1_stride_x Stride of the source image 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 image 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 image + * + * @param[in] input2_ptr Pointer to the source image. Supported data types: F16/F32 + * @param[in] input2_stride_x Stride of the source image 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 image 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 image + * + * @param[out] output_ptr Pointer to the destination image. + * @param[in] output_stride_x Stride of the destination image 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 image 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 source 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 image + */ +__kernel void squared_difference( + 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) + (pow(VLOAD(VEC_SIZE)(0, (__global DATA_TYPE *)input1.ptr) - VLOAD(VEC_SIZE)(0, (__global DATA_TYPE *)input2.ptr), 2), + 0, (__global DATA_TYPE *)output.ptr); +} diff --git a/libs/ARMComputeEx/src/core/CL/kernels/CLSquaredDifferenceKernel.cpp b/libs/ARMComputeEx/src/core/CL/kernels/CLSquaredDifferenceKernel.cpp new file mode 100644 index 0000000..03ae938 --- /dev/null +++ b/libs/ARMComputeEx/src/core/CL/kernels/CLSquaredDifferenceKernel.cpp @@ -0,0 +1,96 @@ +/* + * 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/CLSquaredDifferenceKernel.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; + +CLSquaredDifferenceKernel::CLSquaredDifferenceKernel() + : _input1(nullptr), _input2(nullptr), _output(nullptr) +{ +} + +void CLSquaredDifferenceKernel::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::set build_opts; + build_opts.emplace(("-DDATA_TYPE=" + get_cl_type_from_data_type(input1->info()->data_type()))); + build_opts.emplace( + ("-DVEC_SIZE=" + support::cpp11::to_string(num_elems_processed_per_iteration))); + _kernel = static_cast( + CLKernelLibraryEx::get().create_kernel("squared_difference", 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 CLSquaredDifferenceKernel::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/CLSquaredDifference.cpp b/libs/ARMComputeEx/src/runtime/CL/functions/CLSquaredDifference.cpp new file mode 100644 index 0000000..7c60a54 --- /dev/null +++ b/libs/ARMComputeEx/src/runtime/CL/functions/CLSquaredDifference.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/CLSquaredDifference.h" + +#include "arm_compute/core/CL/kernels/CLSquaredDifferenceKernel.h" + +#include + +using namespace arm_compute; + +void CLSquaredDifference::configure(ICLTensor *input1, ICLTensor *input2, ICLTensor *output) +{ + auto k = arm_compute::support::cpp14::make_unique(); + k->configure(input1, input2, output); + _kernel = std::move(k); +} diff --git a/runtimes/pure_arm_compute/src/compilation.cc b/runtimes/pure_arm_compute/src/compilation.cc index bef95a2..0897da2 100644 --- a/runtimes/pure_arm_compute/src/compilation.cc +++ b/runtimes/pure_arm_compute/src/compilation.cc @@ -49,6 +49,7 @@ #include #include #include +#include #include #include @@ -82,7 +83,6 @@ #include "internal/layers/PadLayer.h" #include "internal/layers/SimpleSpaceToDepth.h" #include "internal/layers/SimpleEmbeddingLookup.h" -#include "internal/layers/SquaredDifferenceOperation.h" #include "internal/layers/SimpleDepthToSpace.h" #include "internal/layers/HashtableLookupLayer.h" #include "internal/layers/SimpleSpaceToBatchND.h" @@ -3445,12 +3445,18 @@ void Planner::visit(const ::internal::tflite::op::SquaredDifference::Node &node) auto lhs_alloc = ctx.at(::internal::tflite::operand::Index{param.lhs_index}); auto rhs_alloc = ctx.at(::internal::tflite::operand::Index{param.rhs_index}); - auto fn = nnfw::make_unique(); - - // TODO Decide ConvertPolicy (WARP? SATURATE?) according to NN API specification - fn->configure(lhs_alloc, rhs_alloc, ofm_alloc); + if (::internal::arm_compute::isGpuMode()) + { + auto fn = nnfw::make_unique<::arm_compute::CLSquaredDifference>(); - builder.append("SquaredDifference", std::move(fn)); + fn->configure(CAST_CL(lhs_alloc), CAST_CL(rhs_alloc), CAST_CL(ofm_alloc)); + builder.append("SquaredDifference", std::move(fn)); + } + else + { + // TODO Enable NEON Support + throw std::runtime_error("Not supported, yet"); + } }; diff --git a/runtimes/pure_arm_compute/src/internal/layers/SquaredDifferenceOperation.cc b/runtimes/pure_arm_compute/src/internal/layers/SquaredDifferenceOperation.cc deleted file mode 100644 index 726f544..0000000 --- a/runtimes/pure_arm_compute/src/internal/layers/SquaredDifferenceOperation.cc +++ /dev/null @@ -1,69 +0,0 @@ -/* - * Copyright (c) 2018 Samsung Electronics Co., Ltd. All Rights Reserved - * - * 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 "SquaredDifferenceOperation.h" -#include "internal/arm_compute.h" - -void SquaredDifferenceOperation::configure(::arm_compute::ITensor *input1, - ::arm_compute::ITensor *input2, - ::arm_compute::ITensor *output) -{ - _input1 = input1; - _input2 = input2; - _output = output; - - if (::internal::arm_compute::isGpuMode()) - { - _cl_intermediate.allocator()->init(*input1->info()); - - _cl_sub.configure(CAST_CL(input1), CAST_CL(input2), &_cl_intermediate, - ::arm_compute::ConvertPolicy::SATURATE); - - const ::arm_compute::ActivationLayerInfo act_info{ - ::arm_compute::ActivationLayerInfo::ActivationFunction::SQUARE}; - - _cl_act.configure(&_cl_intermediate, CAST_CL(output), act_info); - - _cl_intermediate.allocator()->allocate(); - } - else - { - _neon_intermediate.allocator()->init(*input1->info()); - - _neon_sub.configure(input1, input2, &_neon_intermediate, - ::arm_compute::ConvertPolicy::SATURATE); - - const ::arm_compute::ActivationLayerInfo act_info{ - ::arm_compute::ActivationLayerInfo::ActivationFunction::SQUARE}; - - _neon_act.configure(&_neon_intermediate, output, act_info); - - _neon_intermediate.allocator()->allocate(); - } -} - -void SquaredDifferenceOperation::run(void) -{ - if (::internal::arm_compute::isGpuMode()) - { - _cl_sub.run(); - _cl_act.run(); - } - else - { - _neon_sub.run(); - _neon_act.run(); - } -} diff --git a/runtimes/pure_arm_compute/src/internal/layers/SquaredDifferenceOperation.h b/runtimes/pure_arm_compute/src/internal/layers/SquaredDifferenceOperation.h deleted file mode 100644 index f5cc37e..0000000 --- a/runtimes/pure_arm_compute/src/internal/layers/SquaredDifferenceOperation.h +++ /dev/null @@ -1,52 +0,0 @@ -/* - * Copyright (c) 2018 Samsung Electronics Co., Ltd. All Rights Reserved - * - * 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 __SQUARED_DIFFERENCE_OPERATION_H__ -#define __SQUARED_DIFFERENCE_OPERATION_H__ - -#include -#include - -#include -#include -#include -#include - -class SquaredDifferenceOperation : public ::arm_compute::IFunction -{ -public: - void configure(::arm_compute::ITensor *input1, ::arm_compute::ITensor *input2, - ::arm_compute::ITensor *output); - -public: - void run(void) override; - -private: - ::arm_compute::ITensor *_input1; - ::arm_compute::ITensor *_input2; - - ::arm_compute::CLTensor _cl_intermediate; - ::arm_compute::Tensor _neon_intermediate; - - ::arm_compute::ITensor *_output; - -private: - ::arm_compute::CLArithmeticSubtraction _cl_sub; - ::arm_compute::CLActivationLayer _cl_act; - - ::arm_compute::NEArithmeticSubtraction _neon_sub; - ::arm_compute::NEActivationLayer _neon_act; -}; -#endif // __SQUARED_DIFFERENCE_OPERATION_H__