From 63bb7ca40e30b2db48d7bdd1adbc8223b53ac23c Mon Sep 17 00:00:00 2001 From: Manuel Bottini Date: Wed, 2 Dec 2020 13:22:14 +0000 Subject: COMPMID-3921: Remove OpenCL Padding CLBitwiseKernel Adding BitwiseOperation enum class Generalizing CL Bitwise kernels with a single CLBitwiseKernel Removing CL padding from CLBitwiseKernel Change-Id: I79cd79c1e425b6da7d52308a420edf8cfb7a5a36 Signed-off-by: Manuel Bottini Reviewed-on: https://review.mlplatform.org/c/ml/ComputeLibrary/+/4646 Reviewed-by: Giorgio Arena Tested-by: Arm Jenkins Comments-Addressed: Arm Jenkins --- Android.bp | 5 +- arm_compute/core/Types.h | 9 +++ src/core/CL/CLKernels.h | 5 +- src/core/CL/cl_kernels/bitwise_op.cl | 109 ++++++++++++++++++++------- src/core/CL/kernels/CLBitwiseAndKernel.cpp | 92 ----------------------- src/core/CL/kernels/CLBitwiseAndKernel.h | 76 ------------------- src/core/CL/kernels/CLBitwiseKernel.cpp | 116 +++++++++++++++++++++++++++++ src/core/CL/kernels/CLBitwiseKernel.h | 73 ++++++++++++++++++ src/core/CL/kernels/CLBitwiseNotKernel.cpp | 53 ------------- src/core/CL/kernels/CLBitwiseNotKernel.h | 56 -------------- src/core/CL/kernels/CLBitwiseOrKernel.cpp | 93 ----------------------- src/core/CL/kernels/CLBitwiseOrKernel.h | 76 ------------------- src/core/CL/kernels/CLBitwiseXorKernel.cpp | 93 ----------------------- src/core/CL/kernels/CLBitwiseXorKernel.h | 76 ------------------- src/runtime/CL/functions/CLBitwiseAnd.cpp | 11 +-- src/runtime/CL/functions/CLBitwiseNot.cpp | 11 +-- src/runtime/CL/functions/CLBitwiseOr.cpp | 11 +-- src/runtime/CL/functions/CLBitwiseXor.cpp | 11 +-- 18 files changed, 308 insertions(+), 668 deletions(-) delete mode 100644 src/core/CL/kernels/CLBitwiseAndKernel.cpp delete mode 100644 src/core/CL/kernels/CLBitwiseAndKernel.h create mode 100644 src/core/CL/kernels/CLBitwiseKernel.cpp create mode 100644 src/core/CL/kernels/CLBitwiseKernel.h delete mode 100644 src/core/CL/kernels/CLBitwiseNotKernel.cpp delete mode 100644 src/core/CL/kernels/CLBitwiseNotKernel.h delete mode 100644 src/core/CL/kernels/CLBitwiseOrKernel.cpp delete mode 100644 src/core/CL/kernels/CLBitwiseOrKernel.h delete mode 100644 src/core/CL/kernels/CLBitwiseXorKernel.cpp delete mode 100644 src/core/CL/kernels/CLBitwiseXorKernel.h diff --git a/Android.bp b/Android.bp index 2b3b6a806c..25b40f7376 100644 --- a/Android.bp +++ b/Android.bp @@ -85,10 +85,7 @@ cc_library_static { "src/core/CL/kernels/CLBatchConcatenateLayerKernel.cpp", "src/core/CL/kernels/CLBatchNormalizationLayerKernel.cpp", "src/core/CL/kernels/CLBatchToSpaceLayerKernel.cpp", - "src/core/CL/kernels/CLBitwiseAndKernel.cpp", - "src/core/CL/kernels/CLBitwiseNotKernel.cpp", - "src/core/CL/kernels/CLBitwiseOrKernel.cpp", - "src/core/CL/kernels/CLBitwiseXorKernel.cpp", + "src/core/CL/kernels/CLBitwiseKernel.cpp", "src/core/CL/kernels/CLBoundingBoxTransformKernel.cpp", "src/core/CL/kernels/CLBox3x3Kernel.cpp", "src/core/CL/kernels/CLCannyEdgeKernel.cpp", diff --git a/arm_compute/core/Types.h b/arm_compute/core/Types.h index 39cc29b0da..59ba5aabb4 100644 --- a/arm_compute/core/Types.h +++ b/arm_compute/core/Types.h @@ -556,6 +556,15 @@ enum class ElementWiseUnary LOGICAL_NOT, /**< Logical Not */ }; +/** Available bitwise operations */ +enum class BitwiseOperation +{ + AND, /**< Bitwise AND operation */ + NOT, /**< Bitwise NOT operation */ + OR, /**< Bitwise OR operation */ + XOR, /**< Bitwise XOR operation */ +}; + /** The normalization type used for the normalization layer */ enum class NormType { diff --git a/src/core/CL/CLKernels.h b/src/core/CL/CLKernels.h index 42fe79e91d..5d0d326489 100644 --- a/src/core/CL/CLKernels.h +++ b/src/core/CL/CLKernels.h @@ -32,10 +32,7 @@ #include "src/core/CL/kernels/CLBatchConcatenateLayerKernel.h" #include "src/core/CL/kernels/CLBatchNormalizationLayerKernel.h" #include "src/core/CL/kernels/CLBatchToSpaceLayerKernel.h" -#include "src/core/CL/kernels/CLBitwiseAndKernel.h" -#include "src/core/CL/kernels/CLBitwiseNotKernel.h" -#include "src/core/CL/kernels/CLBitwiseOrKernel.h" -#include "src/core/CL/kernels/CLBitwiseXorKernel.h" +#include "src/core/CL/kernels/CLBitwiseKernel.h" #include "src/core/CL/kernels/CLBoundingBoxTransformKernel.h" #include "src/core/CL/kernels/CLBox3x3Kernel.h" #include "src/core/CL/kernels/CLCannyEdgeKernel.h" diff --git a/src/core/CL/cl_kernels/bitwise_op.cl b/src/core/CL/cl_kernels/bitwise_op.cl index b88b3bca22..a600bced9e 100644 --- a/src/core/CL/cl_kernels/bitwise_op.cl +++ b/src/core/CL/cl_kernels/bitwise_op.cl @@ -1,5 +1,5 @@ /* - * Copyright (c) 2016, 2017 Arm Limited. + * Copyright (c) 2016-2020 Arm Limited. * * SPDX-License-Identifier: MIT * @@ -23,7 +23,13 @@ */ #include "helpers.h" +#if defined(VEC_SIZE) && defined(VEC_SIZE_LEFTOVER) + /** This function computes the bitwise OR of two input images. + * + * @note The following variables must be passed at compile time: + * -# -DVEC_SIZE : The number of elements processed in X dimension + * -# -DVEC_SIZE_LEFTOVER: Leftover size in the X dimension; x_dimension % VEC_SIZE * * @param[in] in1_ptr Pointer to the source image. Supported data types: U8 * @param[in] in1_stride_x Stride of the source image in X dimension (in bytes) @@ -49,17 +55,31 @@ __kernel void bitwise_or( IMAGE_DECLARATION(in2), IMAGE_DECLARATION(out)) { - Image in1 = CONVERT_TO_IMAGE_STRUCT(in1); - Image in2 = CONVERT_TO_IMAGE_STRUCT(in2); - Image out = CONVERT_TO_IMAGE_STRUCT(out); + uint x_offs = max((int)(get_global_id(0) * VEC_SIZE - (VEC_SIZE - VEC_SIZE_LEFTOVER) % VEC_SIZE), 0); + + // Get pixels pointer + __global uchar *in1_addr = in1_ptr + in1_offset_first_element_in_bytes + x_offs + get_global_id(1) * in1_step_y; + __global uchar *in2_addr = in2_ptr + in2_offset_first_element_in_bytes + x_offs + get_global_id(1) * in2_step_y; + __global uchar *out_addr = out_ptr + out_offset_first_element_in_bytes + x_offs + get_global_id(1) * out_step_y; + + // Load data + VEC_DATA_TYPE(uchar, VEC_SIZE) + in_a = VLOAD(VEC_SIZE)(0, (__global uchar *)in1_addr); + VEC_DATA_TYPE(uchar, VEC_SIZE) + in_b = VLOAD(VEC_SIZE)(0, (__global uchar *)in2_addr); - uchar16 in_a = vload16(0, in1.ptr); - uchar16 in_b = vload16(0, in2.ptr); + VEC_DATA_TYPE(uchar, VEC_SIZE) + data0 = in_a | in_b; - vstore16(in_a | in_b, 0, out.ptr); + // Boundary-aware store + STORE_VECTOR_SELECT(data, uchar, (__global uchar *)out_addr, VEC_SIZE, VEC_SIZE_LEFTOVER, VEC_SIZE_LEFTOVER != 0 && get_global_id(0) == 0); } /** This function computes the bitwise AND of two input images. + * + * @note The following variables must be passed at compile time: + * -# -DVEC_SIZE : The number of elements processed in X dimension + * -# -DVEC_SIZE_LEFTOVER: Leftover size in the X dimension; x_dimension % VEC_SIZE * * @param[in] in1_ptr Pointer to the source image. Supported data types: U8 * @param[in] in1_stride_x Stride of the source image in X dimension (in bytes) @@ -85,17 +105,31 @@ __kernel void bitwise_and( IMAGE_DECLARATION(in2), IMAGE_DECLARATION(out)) { - Image in1 = CONVERT_TO_IMAGE_STRUCT(in1); - Image in2 = CONVERT_TO_IMAGE_STRUCT(in2); - Image out = CONVERT_TO_IMAGE_STRUCT(out); + uint x_offs = max((int)(get_global_id(0) * VEC_SIZE - (VEC_SIZE - VEC_SIZE_LEFTOVER) % VEC_SIZE), 0); + + // Get pixels pointer + __global uchar *in1_addr = in1_ptr + in1_offset_first_element_in_bytes + x_offs + get_global_id(1) * in1_step_y; + __global uchar *in2_addr = in2_ptr + in2_offset_first_element_in_bytes + x_offs + get_global_id(1) * in2_step_y; + __global uchar *out_addr = out_ptr + out_offset_first_element_in_bytes + x_offs + get_global_id(1) * out_step_y; - uchar16 in_a = vload16(0, in1.ptr); - uchar16 in_b = vload16(0, in2.ptr); + // Load data + VEC_DATA_TYPE(uchar, VEC_SIZE) + in_a = VLOAD(VEC_SIZE)(0, (__global uchar *)in1_addr); + VEC_DATA_TYPE(uchar, VEC_SIZE) + in_b = VLOAD(VEC_SIZE)(0, (__global uchar *)in2_addr); - vstore16(in_a & in_b, 0, out.ptr); + VEC_DATA_TYPE(uchar, VEC_SIZE) + data0 = in_a & in_b; + + // Boundary-aware store + STORE_VECTOR_SELECT(data, uchar, (__global uchar *)out_addr, VEC_SIZE, VEC_SIZE_LEFTOVER, VEC_SIZE_LEFTOVER != 0 && get_global_id(0) == 0); } /** This function computes the bitwise XOR of two input images. + * + * @note The following variables must be passed at compile time: + * -# -DVEC_SIZE : The number of elements processed in X dimension + * -# -DVEC_SIZE_LEFTOVER: Leftover size in the X dimension; x_dimension % VEC_SIZE * * @param[in] in1_ptr Pointer to the source image. Supported data types: U8 * @param[in] in1_stride_x Stride of the source image in X dimension (in bytes) @@ -121,17 +155,31 @@ __kernel void bitwise_xor( IMAGE_DECLARATION(in2), IMAGE_DECLARATION(out)) { - Image in1 = CONVERT_TO_IMAGE_STRUCT(in1); - Image in2 = CONVERT_TO_IMAGE_STRUCT(in2); - Image out = CONVERT_TO_IMAGE_STRUCT(out); + uint x_offs = max((int)(get_global_id(0) * VEC_SIZE - (VEC_SIZE - VEC_SIZE_LEFTOVER) % VEC_SIZE), 0); - uchar16 in_a = vload16(0, in1.ptr); - uchar16 in_b = vload16(0, in2.ptr); + // Get pixels pointer + __global uchar *in1_addr = in1_ptr + in1_offset_first_element_in_bytes + x_offs + get_global_id(1) * in1_step_y; + __global uchar *in2_addr = in2_ptr + in2_offset_first_element_in_bytes + x_offs + get_global_id(1) * in2_step_y; + __global uchar *out_addr = out_ptr + out_offset_first_element_in_bytes + x_offs + get_global_id(1) * out_step_y; - vstore16(in_a ^ in_b, 0, out.ptr); + // Load data + VEC_DATA_TYPE(uchar, VEC_SIZE) + in_a = VLOAD(VEC_SIZE)(0, (__global uchar *)in1_addr); + VEC_DATA_TYPE(uchar, VEC_SIZE) + in_b = VLOAD(VEC_SIZE)(0, (__global uchar *)in2_addr); + + VEC_DATA_TYPE(uchar, VEC_SIZE) + data0 = in_a ^ in_b; + + // Boundary-aware store + STORE_VECTOR_SELECT(data, uchar, (__global uchar *)out_addr, VEC_SIZE, VEC_SIZE_LEFTOVER, VEC_SIZE_LEFTOVER != 0 && get_global_id(0) == 0); } -/** This function computes the bitwise NOT of an image. +/** This function computes the bitwise NOT of an images. + * + * @note The following variables must be passed at compile time: + * -# -DVEC_SIZE : The number of elements processed in X dimension + * -# -DVEC_SIZE_LEFTOVER: Leftover size in the X dimension; x_dimension % VEC_SIZE * * @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) @@ -147,13 +195,24 @@ __kernel void bitwise_xor( * @param[in] out_offset_first_element_in_bytes The offset of the first element in the destination image */ __kernel void bitwise_not( - IMAGE_DECLARATION(in), + IMAGE_DECLARATION(in1), IMAGE_DECLARATION(out)) { - Image in = CONVERT_TO_IMAGE_STRUCT(in); - Image out = CONVERT_TO_IMAGE_STRUCT(out); + uint x_offs = max((int)(get_global_id(0) * VEC_SIZE - (VEC_SIZE - VEC_SIZE_LEFTOVER) % VEC_SIZE), 0); - uchar16 in_data = vload16(0, in.ptr); + // Get pixels pointer + __global uchar *in1_addr = in1_ptr + in1_offset_first_element_in_bytes + x_offs + get_global_id(1) * in1_step_y; + __global uchar *out_addr = out_ptr + out_offset_first_element_in_bytes + x_offs + get_global_id(1) * out_step_y; - vstore16(~in_data, 0, out.ptr); + // Load data + VEC_DATA_TYPE(uchar, VEC_SIZE) + in_a = VLOAD(VEC_SIZE)(0, (__global uchar *)in1_addr); + + VEC_DATA_TYPE(uchar, VEC_SIZE) + data0 = ~in_a; + + // Boundary-aware store + STORE_VECTOR_SELECT(data, uchar, (__global uchar *)out_addr, VEC_SIZE, VEC_SIZE_LEFTOVER, VEC_SIZE_LEFTOVER != 0 && get_global_id(0) == 0); } + +#endif // defined(VEC_SIZE) && defined(VEC_SIZE_LEFTOVER) \ No newline at end of file diff --git a/src/core/CL/kernels/CLBitwiseAndKernel.cpp b/src/core/CL/kernels/CLBitwiseAndKernel.cpp deleted file mode 100644 index 91a659284a..0000000000 --- a/src/core/CL/kernels/CLBitwiseAndKernel.cpp +++ /dev/null @@ -1,92 +0,0 @@ -/* - * Copyright (c) 2016-2020 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 "src/core/CL/kernels/CLBitwiseAndKernel.h" - -#include "arm_compute/core/CL/CLKernelLibrary.h" -#include "arm_compute/core/CL/ICLTensor.h" -#include "arm_compute/core/CL/OpenCL.h" -#include "arm_compute/core/Helpers.h" -#include "arm_compute/core/Validate.h" -#include "src/core/helpers/WindowHelpers.h" - -using namespace arm_compute; - -CLBitwiseAndKernel::CLBitwiseAndKernel() - : _input1(nullptr), _input2(nullptr), _output(nullptr) -{ -} -void CLBitwiseAndKernel::configure(const ICLTensor *input1, const ICLTensor *input2, ICLTensor *output) -{ - configure(CLKernelLibrary::get().get_compile_context(), input1, input2, output); -} - -void CLBitwiseAndKernel::configure(const CLCompileContext &compile_context, const ICLTensor *input1, const ICLTensor *input2, ICLTensor *output) -{ - ARM_COMPUTE_ERROR_ON_DATA_TYPE_CHANNEL_NOT_IN(input1, 1, DataType::U8); - ARM_COMPUTE_ERROR_ON_DATA_TYPE_CHANNEL_NOT_IN(input2, 1, DataType::U8); - ARM_COMPUTE_ERROR_ON_DATA_TYPE_CHANNEL_NOT_IN(output, 1, DataType::U8); - - _input1 = input1; - _input2 = input2; - _output = output; - - // Create kernel - _kernel = create_kernel(compile_context, "bitwise_and"); - - // Configure kernel window - constexpr unsigned int num_elems_processed_per_iteration = 16; - - 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); - - update_window_and_padding(win, input1_access, input2_access, output_access); - - ValidRegion valid_region = intersect_valid_regions(input1->info()->valid_region(), - input2->info()->valid_region()); - - output_access.set_valid_region(win, valid_region); - - ICLKernel::configure_internal(win); -} - -void CLBitwiseAndKernel::run(const Window &window, cl::CommandQueue &queue) -{ - ARM_COMPUTE_ERROR_ON_UNCONFIGURED_KERNEL(this); - ARM_COMPUTE_ERROR_ON_INVALID_SUBWINDOW(ICLKernel::window(), window); - - Window slice = window.first_slice_window_2D(); - - do - { - unsigned int idx = 0; - add_2D_tensor_argument(idx, _input1, slice); - add_2D_tensor_argument(idx, _input2, slice); - add_2D_tensor_argument(idx, _output, slice); - enqueue(queue, *this, slice, lws_hint()); - } - while(window.slide_window_slice_2D(slice)); -} diff --git a/src/core/CL/kernels/CLBitwiseAndKernel.h b/src/core/CL/kernels/CLBitwiseAndKernel.h deleted file mode 100644 index 01018ee09d..0000000000 --- a/src/core/CL/kernels/CLBitwiseAndKernel.h +++ /dev/null @@ -1,76 +0,0 @@ -/* - * Copyright (c) 2016-2020 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_CLBITWISEANDKERNEL_H -#define ARM_COMPUTE_CLBITWISEANDKERNEL_H - -#include "src/core/CL/ICLKernel.h" - -namespace arm_compute -{ -class ICLTensor; - -/** Interface for the bitwise AND operation kernel. - * - * Result is computed by: - * @f[ output(x,y) = input1(x,y) \land input2(x,y) @f] - */ -class CLBitwiseAndKernel : public ICLKernel -{ -public: - /** Default constructor. */ - CLBitwiseAndKernel(); - /** Prevent instances of this class from being copied (As this class contains pointers) */ - CLBitwiseAndKernel(const CLBitwiseAndKernel &) = delete; - /** Prevent instances of this class from being copied (As this class contains pointers) */ - CLBitwiseAndKernel &operator=(const CLBitwiseAndKernel &) = delete; - /** Allow instances of this class to be moved */ - CLBitwiseAndKernel(CLBitwiseAndKernel &&) = default; - /** Allow instances of this class to be moved */ - CLBitwiseAndKernel &operator=(CLBitwiseAndKernel &&) = default; - /** Set the inputs and output images - * - * @param[in] input1 Source tensor. Data types supported: U8. - * @param[in] input2 Source tensor. Data types supported: U8. - * @param[out] output Destination tensor. Data types supported: U8. - */ - void configure(const ICLTensor *input1, const ICLTensor *input2, ICLTensor *output); - /** Set the inputs and output images - * - * @param[in] compile_context The compile context to be used. - * @param[in] input1 Source tensor. Data types supported: U8. - * @param[in] input2 Source tensor. Data types supported: U8. - * @param[out] output Destination tensor. Data types supported: U8. - */ - void configure(const CLCompileContext &compile_context, const ICLTensor *input1, const ICLTensor *input2, ICLTensor *output); - - // Inherited methods overridden: - void run(const Window &window, cl::CommandQueue &queue) override; - -private: - const ICLTensor *_input1; /**< Source tensor 1 */ - const ICLTensor *_input2; /**< Source tensor 2 */ - ICLTensor *_output; /**< Destination tensor */ -}; -} // namespace arm_compute -#endif /* ARM_COMPUTE_CLBITWISEANDKERNEL_H */ diff --git a/src/core/CL/kernels/CLBitwiseKernel.cpp b/src/core/CL/kernels/CLBitwiseKernel.cpp new file mode 100644 index 0000000000..b1f7c00fac --- /dev/null +++ b/src/core/CL/kernels/CLBitwiseKernel.cpp @@ -0,0 +1,116 @@ +/* + * Copyright (c) 2020 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 "src/core/CL/kernels/CLBitwiseKernel.h" + +#include "arm_compute/core/CL/CLKernelLibrary.h" +#include "arm_compute/core/CL/ICLTensor.h" +#include "arm_compute/core/CL/OpenCL.h" +#include "arm_compute/core/Helpers.h" +#include "arm_compute/core/Validate.h" +#include "src/core/helpers/AutoConfiguration.h" +#include "src/core/helpers/WindowHelpers.h" +#include "support/StringSupport.h" + +namespace arm_compute +{ +CLBitwiseKernel::CLBitwiseKernel() + : _input1(nullptr), _input2(nullptr), _output(nullptr) +{ +} + +void CLBitwiseKernel::configure(const CLCompileContext &compile_context, const ICLTensor *input1, const ICLTensor *input2, ICLTensor *output, BitwiseOperation op) +{ + ARM_COMPUTE_ERROR_ON_NULLPTR(input1); + ARM_COMPUTE_ERROR_ON_DATA_TYPE_CHANNEL_NOT_IN(input1, 1, DataType::U8); + if(op != BitwiseOperation::NOT) + { + ARM_COMPUTE_ERROR_ON_NULLPTR(input2); + ARM_COMPUTE_ERROR_ON_DATA_TYPE_CHANNEL_NOT_IN(input2, 1, DataType::U8); + } + ARM_COMPUTE_ERROR_ON_NULLPTR(output); + ARM_COMPUTE_ERROR_ON_DATA_TYPE_CHANNEL_NOT_IN(output, 1, DataType::U8); + + // Output auto inizialitation if not yet initialized + auto_init_if_empty(*(output->info()), *(input1->info())); + auto padding_info = get_padding_info({ input1, input2, output }); + + // Configure kernel window + const unsigned int vec_size_x = adjust_vec_size(16 / output->info()->element_size(), output->info()->dimension(0)); + Window win = calculate_max_window(*output->info(), Steps(vec_size_x)); + + _input1 = input1; + _input2 = input2; + _output = output; + + // Create kernel + std::string kernel_name = ""; + switch(op) + { + case BitwiseOperation::AND: + kernel_name = "bitwise_and"; + break; + case BitwiseOperation::NOT: + kernel_name = "bitwise_not"; + break; + case BitwiseOperation::OR: + kernel_name = "bitwise_or"; + break; + case BitwiseOperation::XOR: + kernel_name = "bitwise_xor"; + break; + default: + ARM_COMPUTE_ERROR("Bitwise operation not supported"); + } + + CLBuildOptions build_opts; + const int vec_size_x_leftovers = output->info()->dimension(0) % vec_size_x; + build_opts.add_option("-DVEC_SIZE=" + support::cpp11::to_string(vec_size_x)); + build_opts.add_option("-DVEC_SIZE_LEFTOVER=" + support::cpp11::to_string(vec_size_x_leftovers)); + _kernel = create_kernel(compile_context, kernel_name, build_opts.options()); + + ICLKernel::configure_internal(win); + ARM_COMPUTE_ERROR_ON(has_padding_changed(padding_info)); +} + +void CLBitwiseKernel::run(const Window &window, cl::CommandQueue &queue) +{ + ARM_COMPUTE_ERROR_ON_UNCONFIGURED_KERNEL(this); + ARM_COMPUTE_ERROR_ON_INVALID_SUBWINDOW(ICLKernel::window(), window); + + Window slice = window.first_slice_window_2D(); + + do + { + unsigned int idx = 0; + add_2D_tensor_argument(idx, _input1, slice); + if(_input2 != nullptr) + { + add_2D_tensor_argument(idx, _input2, slice); + } + add_2D_tensor_argument(idx, _output, slice); + enqueue(queue, *this, slice, lws_hint()); + } + while(window.slide_window_slice_2D(slice)); +} +} // namespace arm_compute \ No newline at end of file diff --git a/src/core/CL/kernels/CLBitwiseKernel.h b/src/core/CL/kernels/CLBitwiseKernel.h new file mode 100644 index 0000000000..c5a999643d --- /dev/null +++ b/src/core/CL/kernels/CLBitwiseKernel.h @@ -0,0 +1,73 @@ +/* + * Copyright (c) 2016-2020 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_CLBITWISEKERNEL_H +#define ARM_COMPUTE_CLBITWISEKERNEL_H + +#include "src/core/CL/ICLKernel.h" + +namespace arm_compute +{ +class ICLTensor; + +/** Interface for the bitwise operation kernel. + * + * Result depends on the \ref BitwiseOperation and is computed by: + * AND operation: @f[ output(x,y) = input1(x,y) \land input2(x,y) @f] + * NOT operation: @f[ output(x,y) = \lnot input1(x,y) @f] + * OR operation: @f[ output(x,y) = input1(x,y) \lor input2(x,y) @f] + * XOR operation: @f[ output(x,y) = input1(x,y) \oplus input2(x,y) @f] + */ +class CLBitwiseKernel : public ICLKernel +{ +public: + /** Default constructor. */ + CLBitwiseKernel(); + /** Prevent instances of this class from being copied (As this class contains pointers) */ + CLBitwiseKernel(const CLBitwiseKernel &) = delete; + /** Prevent instances of this class from being copied (As this class contains pointers) */ + CLBitwiseKernel &operator=(const CLBitwiseKernel &) = delete; + /** Allow instances of this class to be moved */ + CLBitwiseKernel(CLBitwiseKernel &&) = default; + /** Allow instances of this class to be moved */ + CLBitwiseKernel &operator=(CLBitwiseKernel &&) = default; + /** Set the inputs and output tensors + * + * @param[in] compile_context The compile context to be used. + * @param[in] input1 Source tensor. Data types supported: U8. + * @param[in] input2 Source tensor. Data types supported: U8. + * @param[out] output Destination tensor. Data types supported: U8. + * @param[in] op Bitwise operation to perform. Supported: AND, OR, NOT, XOR. + */ + void configure(const CLCompileContext &compile_context, const ICLTensor *input1, const ICLTensor *input2, ICLTensor *output, BitwiseOperation op); + + // Inherited methods overridden: + void run(const Window &window, cl::CommandQueue &queue) override; + +private: + const ICLTensor *_input1; /**< Source tensor 1 */ + const ICLTensor *_input2; /**< Source tensor 2 */ + ICLTensor *_output; /**< Destination tensor */ +}; +} // namespace arm_compute +#endif /* ARM_COMPUTE_CLBITWISEKERNEL_H */ diff --git a/src/core/CL/kernels/CLBitwiseNotKernel.cpp b/src/core/CL/kernels/CLBitwiseNotKernel.cpp deleted file mode 100644 index 118bfe8139..0000000000 --- a/src/core/CL/kernels/CLBitwiseNotKernel.cpp +++ /dev/null @@ -1,53 +0,0 @@ -/* - * Copyright (c) 2016-2020 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 "src/core/CL/kernels/CLBitwiseNotKernel.h" - -#include "arm_compute/core/CL/CLKernelLibrary.h" -#include "arm_compute/core/CL/ICLTensor.h" -#include "arm_compute/core/CL/OpenCL.h" -#include "arm_compute/core/Types.h" -#include "arm_compute/core/Validate.h" - -using namespace arm_compute; - -void CLBitwiseNotKernel::configure(const ICLTensor *input, ICLTensor *output) -{ - configure(CLKernelLibrary::get().get_compile_context(), input, output); -} - -void CLBitwiseNotKernel::configure(const CLCompileContext &compile_context, const ICLTensor *input, ICLTensor *output) -{ - ARM_COMPUTE_ERROR_ON_DATA_TYPE_CHANNEL_NOT_IN(input, 1, DataType::U8); - ARM_COMPUTE_ERROR_ON_DATA_TYPE_CHANNEL_NOT_IN(output, 1, DataType::U8); - - _input = input; - _output = output; - - // Create kernel - _kernel = create_kernel(compile_context, "bitwise_not"); - - // Configure kernel window - constexpr unsigned int num_elems_processed_per_iteration = 16; - ICLSimple2DKernel::configure(input, output, num_elems_processed_per_iteration); -} diff --git a/src/core/CL/kernels/CLBitwiseNotKernel.h b/src/core/CL/kernels/CLBitwiseNotKernel.h deleted file mode 100644 index bf68bc7ae5..0000000000 --- a/src/core/CL/kernels/CLBitwiseNotKernel.h +++ /dev/null @@ -1,56 +0,0 @@ -/* - * Copyright (c) 2016-2020 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_CLBITWISENOTKERNEL_H -#define ARM_COMPUTE_CLBITWISENOTKERNEL_H - -#include "src/core/CL/ICLSimple2DKernel.h" - -namespace arm_compute -{ -class ICLTensor; - -/** Interface for the bitwise NOT operation kernel. - * - * Result is computed by: - * @f[ output(x,y) = \lnot input(x,y) @f] - */ -class CLBitwiseNotKernel : public ICLSimple2DKernel -{ -public: - /** Set the inputs and output images. - * - * @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); - /** Set the inputs and output images. - * - * @param[in] compile_context The compile context to be used. - * @param[in] input Source tensor. Data types supported: U8. - * @param[out] output Destination tensor. Data types supported: U8. - */ - void configure(const CLCompileContext &compile_context, const ICLTensor *input, ICLTensor *output); -}; -} // namespace arm_compute -#endif /* ARM_COMPUTE_CLBITWISENOTKERNEL_H */ diff --git a/src/core/CL/kernels/CLBitwiseOrKernel.cpp b/src/core/CL/kernels/CLBitwiseOrKernel.cpp deleted file mode 100644 index 8954d9aa6d..0000000000 --- a/src/core/CL/kernels/CLBitwiseOrKernel.cpp +++ /dev/null @@ -1,93 +0,0 @@ -/* - * Copyright (c) 2016-2020 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 "src/core/CL/kernels/CLBitwiseOrKernel.h" - -#include "arm_compute/core/CL/CLKernelLibrary.h" -#include "arm_compute/core/CL/ICLTensor.h" -#include "arm_compute/core/CL/OpenCL.h" -#include "arm_compute/core/Helpers.h" -#include "arm_compute/core/Validate.h" -#include "src/core/helpers/WindowHelpers.h" - -using namespace arm_compute; - -CLBitwiseOrKernel::CLBitwiseOrKernel() - : _input1(nullptr), _input2(nullptr), _output(nullptr) -{ -} - -void CLBitwiseOrKernel::configure(const ICLTensor *input1, const ICLTensor *input2, ICLTensor *output) -{ - configure(CLKernelLibrary::get().get_compile_context(), input1, input2, output); -} - -void CLBitwiseOrKernel::configure(const CLCompileContext &compile_context, const ICLTensor *input1, const ICLTensor *input2, ICLTensor *output) -{ - ARM_COMPUTE_ERROR_ON_DATA_TYPE_CHANNEL_NOT_IN(input1, 1, DataType::U8); - ARM_COMPUTE_ERROR_ON_DATA_TYPE_CHANNEL_NOT_IN(input2, 1, DataType::U8); - ARM_COMPUTE_ERROR_ON_DATA_TYPE_CHANNEL_NOT_IN(output, 1, DataType::U8); - - _input1 = input1; - _input2 = input2; - _output = output; - - // Create kernel - _kernel = create_kernel(compile_context, "bitwise_or"); - - // Configure kernel window - constexpr unsigned int num_elems_processed_per_iteration = 16; - - 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); - - update_window_and_padding(win, input1_access, input2_access, output_access); - - ValidRegion valid_region = intersect_valid_regions(input1->info()->valid_region(), - input2->info()->valid_region()); - - output_access.set_valid_region(win, valid_region); - - ICLKernel::configure_internal(win); -} - -void CLBitwiseOrKernel::run(const Window &window, cl::CommandQueue &queue) -{ - ARM_COMPUTE_ERROR_ON_UNCONFIGURED_KERNEL(this); - ARM_COMPUTE_ERROR_ON_INVALID_SUBWINDOW(ICLKernel::window(), window); - - Window slice = window.first_slice_window_2D(); - - do - { - unsigned int idx = 0; - add_2D_tensor_argument(idx, _input1, slice); - add_2D_tensor_argument(idx, _input2, slice); - add_2D_tensor_argument(idx, _output, slice); - enqueue(queue, *this, slice, lws_hint()); - } - while(window.slide_window_slice_2D(slice)); -} diff --git a/src/core/CL/kernels/CLBitwiseOrKernel.h b/src/core/CL/kernels/CLBitwiseOrKernel.h deleted file mode 100644 index c27d0c27e2..0000000000 --- a/src/core/CL/kernels/CLBitwiseOrKernel.h +++ /dev/null @@ -1,76 +0,0 @@ -/* - * Copyright (c) 2016-2020 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_CLBITWISEORKERNEL_H -#define ARM_COMPUTE_CLBITWISEORKERNEL_H - -#include "src/core/CL/ICLKernel.h" - -namespace arm_compute -{ -class ICLTensor; - -/** Interface for the bitwise OR operation kernel. - * - * Result is computed by: - * @f[ output(x,y) = input1(x,y) \lor input2(x,y) @f] - */ -class CLBitwiseOrKernel : public ICLKernel -{ -public: - /** Default constructor. */ - CLBitwiseOrKernel(); - /** Prevent instances of this class from being copied (As this class contains pointers) */ - CLBitwiseOrKernel(const CLBitwiseOrKernel &) = delete; - /** Prevent instances of this class from being copied (As this class contains pointers) */ - CLBitwiseOrKernel &operator=(const CLBitwiseOrKernel &) = delete; - /** Allow instances of this class to be moved */ - CLBitwiseOrKernel(CLBitwiseOrKernel &&) = default; - /** Allow instances of this class to be moved */ - CLBitwiseOrKernel &operator=(CLBitwiseOrKernel &&) = default; - /** Set the inputs and output images - * - * @param[in] input1 Source tensor. Data types supported: U8. - * @param[in] input2 Source tensor. Data types supported: U8. - * @param[out] output Destination tensor. Data types supported: U8. - */ - void configure(const ICLTensor *input1, const ICLTensor *input2, ICLTensor *output); - /** Set the inputs and output images - * - * @param[in] compile_context The compile context to be used. - * @param[in] input1 Source tensor. Data types supported: U8. - * @param[in] input2 Source tensor. Data types supported: U8. - * @param[out] output Destination tensor. Data types supported: U8. - */ - void configure(const CLCompileContext &compile_context, const ICLTensor *input1, const ICLTensor *input2, ICLTensor *output); - - // Inherited methods overridden: - void run(const Window &window, cl::CommandQueue &queue) override; - -private: - const ICLTensor *_input1; /**< Source tensor 1 */ - const ICLTensor *_input2; /**< Source tensor 2 */ - ICLTensor *_output; /**< Destination tensor */ -}; -} // namespace arm_compute -#endif /* ARM_COMPUTE_CLBITWISEORKERNEL_H */ diff --git a/src/core/CL/kernels/CLBitwiseXorKernel.cpp b/src/core/CL/kernels/CLBitwiseXorKernel.cpp deleted file mode 100644 index 69eb38e2e6..0000000000 --- a/src/core/CL/kernels/CLBitwiseXorKernel.cpp +++ /dev/null @@ -1,93 +0,0 @@ -/* - * Copyright (c) 2016-2020 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 "src/core/CL/kernels/CLBitwiseXorKernel.h" - -#include "arm_compute/core/CL/CLKernelLibrary.h" -#include "arm_compute/core/CL/ICLTensor.h" -#include "arm_compute/core/CL/OpenCL.h" -#include "arm_compute/core/Helpers.h" -#include "arm_compute/core/Validate.h" -#include "src/core/helpers/WindowHelpers.h" - -using namespace arm_compute; - -CLBitwiseXorKernel::CLBitwiseXorKernel() - : _input1(nullptr), _input2(nullptr), _output(nullptr) -{ -} - -void CLBitwiseXorKernel::configure(const ICLTensor *input1, const ICLTensor *input2, ICLTensor *output) -{ - configure(CLKernelLibrary::get().get_compile_context(), input1, input2, output); -} - -void CLBitwiseXorKernel::configure(const CLCompileContext &compile_context, const ICLTensor *input1, const ICLTensor *input2, ICLTensor *output) -{ - ARM_COMPUTE_ERROR_ON_DATA_TYPE_CHANNEL_NOT_IN(input1, 1, DataType::U8); - ARM_COMPUTE_ERROR_ON_DATA_TYPE_CHANNEL_NOT_IN(input2, 1, DataType::U8); - ARM_COMPUTE_ERROR_ON_DATA_TYPE_CHANNEL_NOT_IN(output, 1, DataType::U8); - - _input1 = input1; - _input2 = input2; - _output = output; - - // Create kernel - _kernel = create_kernel(compile_context, "bitwise_xor"); - - // Configure kernel window - constexpr unsigned int num_elems_processed_per_iteration = 16; - - 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); - - update_window_and_padding(win, input1_access, input2_access, output_access); - - ValidRegion valid_region = intersect_valid_regions(input1->info()->valid_region(), - input2->info()->valid_region()); - - output_access.set_valid_region(win, valid_region); - - ICLKernel::configure_internal(win); -} - -void CLBitwiseXorKernel::run(const Window &window, cl::CommandQueue &queue) -{ - ARM_COMPUTE_ERROR_ON_UNCONFIGURED_KERNEL(this); - ARM_COMPUTE_ERROR_ON_INVALID_SUBWINDOW(ICLKernel::window(), window); - - Window slice = window.first_slice_window_2D(); - - do - { - unsigned int idx = 0; - add_2D_tensor_argument(idx, _input1, slice); - add_2D_tensor_argument(idx, _input2, slice); - add_2D_tensor_argument(idx, _output, slice); - enqueue(queue, *this, slice, lws_hint()); - } - while(window.slide_window_slice_2D(slice)); -} diff --git a/src/core/CL/kernels/CLBitwiseXorKernel.h b/src/core/CL/kernels/CLBitwiseXorKernel.h deleted file mode 100644 index b4861ea757..0000000000 --- a/src/core/CL/kernels/CLBitwiseXorKernel.h +++ /dev/null @@ -1,76 +0,0 @@ -/* - * Copyright (c) 2016-2020 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_CLBITWISEXORKERNEL_H -#define ARM_COMPUTE_CLBITWISEXORKERNEL_H - -#include "src/core/CL/ICLKernel.h" - -namespace arm_compute -{ -class ICLTensor; - -/** Interface for the bitwise XOR operation kernel. - * - * Result is computed by: - * @f[ output(x,y) = input1(x,y) \oplus input2(x,y) @f] - */ -class CLBitwiseXorKernel : public ICLKernel -{ -public: - /** Default constructor. */ - CLBitwiseXorKernel(); - /** Prevent instances of this class from being copied (As this class contains pointers) */ - CLBitwiseXorKernel(const CLBitwiseXorKernel &) = delete; - /** Prevent instances of this class from being copied (As this class contains pointers) */ - CLBitwiseXorKernel &operator=(const CLBitwiseXorKernel &) = delete; - /** Allow instances of this class to be moved */ - CLBitwiseXorKernel(CLBitwiseXorKernel &&) = default; - /** Allow instances of this class to be moved */ - CLBitwiseXorKernel &operator=(CLBitwiseXorKernel &&) = default; - /** Set the inputs and output images - * - * @param[in] input1 Source tensor. Data types supported: U8. - * @param[in] input2 Source tensor. Data types supported: U8. - * @param[out] output Destination tensor. Data types supported: U8. - */ - void configure(const ICLTensor *input1, const ICLTensor *input2, ICLTensor *output); - /** Set the inputs and output images - * - * @param[in] compile_context The compile context to be used. - * @param[in] input1 Source tensor. Data types supported: U8. - * @param[in] input2 Source tensor. Data types supported: U8. - * @param[out] output Destination tensor. Data types supported: U8. - */ - void configure(const CLCompileContext &compile_context, const ICLTensor *input1, const ICLTensor *input2, ICLTensor *output); - - // Inherited methods overridden: - void run(const Window &window, cl::CommandQueue &queue) override; - -private: - const ICLTensor *_input1; /**< Source tensor 1 */ - const ICLTensor *_input2; /**< Source tensor 2 */ - ICLTensor *_output; /**< Destination tensor */ -}; -} // namespace arm_compute -#endif /* ARM_COMPUTE_CLBITWISEXORKERNEL_H */ diff --git a/src/runtime/CL/functions/CLBitwiseAnd.cpp b/src/runtime/CL/functions/CLBitwiseAnd.cpp index 0f9f68cb9c..70e27c0cca 100644 --- a/src/runtime/CL/functions/CLBitwiseAnd.cpp +++ b/src/runtime/CL/functions/CLBitwiseAnd.cpp @@ -23,12 +23,12 @@ */ #include "arm_compute/runtime/CL/functions/CLBitwiseAnd.h" -#include "src/core/CL/kernels/CLBitwiseAndKernel.h" +#include "src/core/CL/kernels/CLBitwiseKernel.h" #include -using namespace arm_compute; - +namespace arm_compute +{ void CLBitwiseAnd::configure(const ICLTensor *input1, const ICLTensor *input2, ICLTensor *output) { configure(CLKernelLibrary::get().get_compile_context(), input1, input2, output); @@ -36,7 +36,8 @@ void CLBitwiseAnd::configure(const ICLTensor *input1, const ICLTensor *input2, I void CLBitwiseAnd::configure(const CLCompileContext &compile_context, const ICLTensor *input1, const ICLTensor *input2, ICLTensor *output) { - auto k = std::make_unique(); - k->configure(compile_context, input1, input2, output); + auto k = std::make_unique(); + k->configure(compile_context, input1, input2, output, BitwiseOperation::AND); _kernel = std::move(k); } +} // namespace arm_compute \ No newline at end of file diff --git a/src/runtime/CL/functions/CLBitwiseNot.cpp b/src/runtime/CL/functions/CLBitwiseNot.cpp index cd2384590e..7970a1698b 100644 --- a/src/runtime/CL/functions/CLBitwiseNot.cpp +++ b/src/runtime/CL/functions/CLBitwiseNot.cpp @@ -23,12 +23,12 @@ */ #include "arm_compute/runtime/CL/functions/CLBitwiseNot.h" -#include "src/core/CL/kernels/CLBitwiseNotKernel.h" +#include "src/core/CL/kernels/CLBitwiseKernel.h" #include -using namespace arm_compute; - +namespace arm_compute +{ void CLBitwiseNot::configure(const ICLTensor *input, ICLTensor *output) { configure(CLKernelLibrary::get().get_compile_context(), input, output); @@ -36,7 +36,8 @@ void CLBitwiseNot::configure(const ICLTensor *input, ICLTensor *output) void CLBitwiseNot::configure(const CLCompileContext &compile_context, const ICLTensor *input, ICLTensor *output) { - auto k = std::make_unique(); - k->configure(compile_context, input, output); + auto k = std::make_unique(); + k->configure(compile_context, input, nullptr, output, BitwiseOperation::NOT); _kernel = std::move(k); } +} // namespace arm_compute \ No newline at end of file diff --git a/src/runtime/CL/functions/CLBitwiseOr.cpp b/src/runtime/CL/functions/CLBitwiseOr.cpp index 38db5f78a0..fbda9ad289 100644 --- a/src/runtime/CL/functions/CLBitwiseOr.cpp +++ b/src/runtime/CL/functions/CLBitwiseOr.cpp @@ -23,12 +23,12 @@ */ #include "arm_compute/runtime/CL/functions/CLBitwiseOr.h" -#include "src/core/CL/kernels/CLBitwiseOrKernel.h" +#include "src/core/CL/kernels/CLBitwiseKernel.h" #include -using namespace arm_compute; - +namespace arm_compute +{ void CLBitwiseOr::configure(const ICLTensor *input1, const ICLTensor *input2, ICLTensor *output) { configure(CLKernelLibrary::get().get_compile_context(), input1, input2, output); @@ -36,7 +36,8 @@ void CLBitwiseOr::configure(const ICLTensor *input1, const ICLTensor *input2, IC void CLBitwiseOr::configure(const CLCompileContext &compile_context, const ICLTensor *input1, const ICLTensor *input2, ICLTensor *output) { - auto k = std::make_unique(); - k->configure(compile_context, input1, input2, output); + auto k = std::make_unique(); + k->configure(compile_context, input1, input2, output, BitwiseOperation::OR); _kernel = std::move(k); } +} // namespace arm_compute \ No newline at end of file diff --git a/src/runtime/CL/functions/CLBitwiseXor.cpp b/src/runtime/CL/functions/CLBitwiseXor.cpp index e477c3b847..4f4b74c04c 100644 --- a/src/runtime/CL/functions/CLBitwiseXor.cpp +++ b/src/runtime/CL/functions/CLBitwiseXor.cpp @@ -23,12 +23,12 @@ */ #include "arm_compute/runtime/CL/functions/CLBitwiseXor.h" -#include "src/core/CL/kernels/CLBitwiseXorKernel.h" +#include "src/core/CL/kernels/CLBitwiseKernel.h" #include -using namespace arm_compute; - +namespace arm_compute +{ void CLBitwiseXor::configure(const ICLTensor *input1, const ICLTensor *input2, ICLTensor *output) { configure(CLKernelLibrary::get().get_compile_context(), input1, input2, output); @@ -36,7 +36,8 @@ void CLBitwiseXor::configure(const ICLTensor *input1, const ICLTensor *input2, I void CLBitwiseXor::configure(const CLCompileContext &compile_context, const ICLTensor *input1, const ICLTensor *input2, ICLTensor *output) { - auto k = std::make_unique(); - k->configure(compile_context, input1, input2, output); + auto k = std::make_unique(); + k->configure(compile_context, input1, input2, output, BitwiseOperation::XOR); _kernel = std::move(k); } +} // namespace arm_compute \ No newline at end of file -- cgit v1.2.1