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 <manuel.bottini@arm.com>
Reviewed-on: https://review.mlplatform.org/c/ml/ComputeLibrary/+/4646
Reviewed-by: Giorgio Arena <giorgio.arena@arm.com>
Tested-by: Arm Jenkins <bsgcomp@arm.com>
Comments-Addressed: Arm Jenkins <bsgcomp@arm.com>
diff --git a/src/core/CL/CLKernels.h b/src/core/CL/CLKernels.h
index 42fe79e..5d0d326 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 b88b3bc..a600bce 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,8 +23,14 @@
*/
#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)
* @param[in] in1_step_x in1_stride_x * number of elements along X processed per workitem(in bytes)
@@ -49,18 +55,32 @@
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 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)
* @param[in] in1_step_x in1_stride_x * number of elements along X processed per workitem(in bytes)
@@ -85,18 +105,32 @@
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 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)
* @param[in] in1_step_x in1_stride_x * number of elements along X processed per workitem(in bytes)
@@ -121,17 +155,31 @@
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 @@
* @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 91a6592..0000000
--- 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/CLBitwiseKernel.cpp b/src/core/CL/kernels/CLBitwiseKernel.cpp
new file mode 100644
index 0000000..b1f7c00
--- /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/CLBitwiseAndKernel.h b/src/core/CL/kernels/CLBitwiseKernel.h
similarity index 68%
rename from src/core/CL/kernels/CLBitwiseAndKernel.h
rename to src/core/CL/kernels/CLBitwiseKernel.h
index 01018ee..c5a9996 100644
--- a/src/core/CL/kernels/CLBitwiseAndKernel.h
+++ b/src/core/CL/kernels/CLBitwiseKernel.h
@@ -21,8 +21,8 @@
* 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
+#ifndef ARM_COMPUTE_CLBITWISEKERNEL_H
+#define ARM_COMPUTE_CLBITWISEKERNEL_H
#include "src/core/CL/ICLKernel.h"
@@ -30,39 +30,36 @@
{
class ICLTensor;
-/** Interface for the bitwise AND operation kernel.
+/** Interface for the bitwise operation kernel.
*
- * Result is computed by:
- * @f[ output(x,y) = input1(x,y) \land input2(x,y) @f]
+ * 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 CLBitwiseAndKernel : public ICLKernel
+class CLBitwiseKernel : public ICLKernel
{
public:
/** Default constructor. */
- CLBitwiseAndKernel();
+ CLBitwiseKernel();
/** Prevent instances of this class from being copied (As this class contains pointers) */
- CLBitwiseAndKernel(const CLBitwiseAndKernel &) = delete;
+ CLBitwiseKernel(const CLBitwiseKernel &) = delete;
/** Prevent instances of this class from being copied (As this class contains pointers) */
- CLBitwiseAndKernel &operator=(const CLBitwiseAndKernel &) = delete;
+ CLBitwiseKernel &operator=(const CLBitwiseKernel &) = delete;
/** Allow instances of this class to be moved */
- CLBitwiseAndKernel(CLBitwiseAndKernel &&) = default;
+ CLBitwiseKernel(CLBitwiseKernel &&) = 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
+ 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);
+ 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;
@@ -73,4 +70,4 @@
ICLTensor *_output; /**< Destination tensor */
};
} // namespace arm_compute
-#endif /* ARM_COMPUTE_CLBITWISEANDKERNEL_H */
+#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 118bfe8..0000000
--- 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 bf68bc7..0000000
--- 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 8954d9a..0000000
--- 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 c27d0c2..0000000
--- 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 69eb38e..0000000
--- 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 b4861ea..0000000
--- 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 */