Compute Library
 21.05
ClArithmeticKernel Class Reference

#include <ClElementwiseKernel.h>

Collaboration diagram for ClArithmeticKernel:
[legend]

Public Member Functions

 ClArithmeticKernel ()=default
 
 ARM_COMPUTE_DISALLOW_COPY_ALLOW_MOVE (ClArithmeticKernel)
 
void configure (const ClCompileContext &compile_context, ArithmeticOperation op, ITensorInfo *src1, ITensorInfo *src2, ITensorInfo *dst, const ActivationLayerInfo &act_info=ActivationLayerInfo())
 Static function to check if given info will lead to a valid configuration of ClArithmeticKernel. More...
 
- Public Member Functions inherited from ClElementwiseKernel
 ClElementwiseKernel ()=default
 Default constructor. More...
 
 ARM_COMPUTE_DISALLOW_COPY_ALLOW_MOVE (ClElementwiseKernel)
 
void run_op (ITensorPack &tensors, const Window &window, ::cl::CommandQueue &queue) override
 
- Public Member Functions inherited from ICLKernel
 ICLKernel ()
 Constructor. More...
 
cl::Kernel & kernel ()
 Returns a reference to the OpenCL kernel of this object. More...
 
template<typename T >
void add_1D_array_argument (unsigned int &idx, const ICLArray< T > *array, const Strides &strides, unsigned int num_dimensions, const Window &window)
 Add the passed 1D array's parameters to the object's kernel's arguments starting from the index idx. More...
 
void add_1D_tensor_argument (unsigned int &idx, const ICLTensor *tensor, const Window &window)
 Add the passed 1D tensor's parameters to the object's kernel's arguments starting from the index idx. More...
 
void add_1D_tensor_argument_if (bool cond, unsigned int &idx, const ICLTensor *tensor, const Window &window)
 Add the passed 1D tensor's parameters to the object's kernel's arguments starting from the index idx if the condition is true. More...
 
void add_2D_tensor_argument (unsigned int &idx, const ICLTensor *tensor, const Window &window)
 Add the passed 2D tensor's parameters to the object's kernel's arguments starting from the index idx. More...
 
void add_2D_tensor_argument_if (bool cond, unsigned int &idx, const ICLTensor *tensor, const Window &window)
 Add the passed 2D tensor's parameters to the object's kernel's arguments starting from the index idx if the condition is true. More...
 
void add_3D_tensor_argument (unsigned int &idx, const ICLTensor *tensor, const Window &window)
 Add the passed 3D tensor's parameters to the object's kernel's arguments starting from the index idx. More...
 
void add_4D_tensor_argument (unsigned int &idx, const ICLTensor *tensor, const Window &window)
 Add the passed 4D tensor's parameters to the object's kernel's arguments starting from the index idx. More...
 
virtual void run (const Window &window, cl::CommandQueue &queue)
 Enqueue the OpenCL kernel to process the given window on the passed OpenCL command queue. More...
 
virtual void run_op (ITensorPack &tensors, const Window &window, cl::CommandQueue &queue)
 Enqueue the OpenCL kernel to process the given window on the passed OpenCL command queue. More...
 
template<typename T >
void add_argument (unsigned int &idx, T value)
 Add the passed parameters to the object's kernel's arguments starting from the index idx. More...
 
void set_lws_hint (const cl::NDRange &lws_hint)
 Set the Local-Workgroup-Size hint. More...
 
cl::NDRange lws_hint () const
 Return the Local-Workgroup-Size hint. More...
 
void set_wbsm_hint (const cl_int &wbsm_hint)
 Set the workgroup batch size modifier hint. More...
 
cl_int wbsm_hint () const
 Return the workgroup batch size modifier hint. More...
 
const std::string & config_id () const
 Get the configuration ID. More...
 
void set_target (GPUTarget target)
 Set the targeted GPU architecture. More...
 
void set_target (cl::Device &device)
 Set the targeted GPU architecture according to the CL device. More...
 
GPUTarget get_target () const
 Get the targeted GPU architecture. More...
 
size_t get_max_workgroup_size ()
 Get the maximum workgroup size for the device the CLKernelLibrary uses. More...
 
template<unsigned int dimension_size>
void add_tensor_argument (unsigned &idx, const ICLTensor *tensor, const Window &window)
 
template<typename T , unsigned int dimension_size>
void add_array_argument (unsigned &idx, const ICLArray< T > *array, const Strides &strides, unsigned int num_dimensions, const Window &window)
 Add the passed array's parameters to the object's kernel's arguments starting from the index idx. More...
 
- Public Member Functions inherited from IKernel
 IKernel ()
 Constructor. More...
 
virtual ~IKernel ()=default
 Destructor. More...
 
virtual bool is_parallelisable () const
 Indicates whether or not the kernel is parallelisable. More...
 
virtual BorderSize border_size () const
 The size of the border for that kernel. More...
 
const Windowwindow () const
 The maximum window the kernel can be executed on. More...
 
bool is_window_configured () const
 Function to check if the embedded window of this kernel has been configured. More...
 

Static Public Member Functions

static Status validate (ArithmeticOperation op, const ITensorInfo *src1, const ITensorInfo *src2, const ITensorInfo *dst, const ActivationLayerInfo &act_info=ActivationLayerInfo())
 Static function to check if given info will lead to a valid configuration of ClArithmeticKernel. More...
 
- Static Public Member Functions inherited from ICLKernel
static constexpr unsigned int num_arguments_per_1D_array ()
 Returns the number of arguments enqueued per 1D array object. More...
 
static constexpr unsigned int num_arguments_per_1D_tensor ()
 Returns the number of arguments enqueued per 1D tensor object. More...
 
static constexpr unsigned int num_arguments_per_2D_tensor ()
 Returns the number of arguments enqueued per 2D tensor object. More...
 
static constexpr unsigned int num_arguments_per_3D_tensor ()
 Returns the number of arguments enqueued per 3D tensor object. More...
 
static constexpr unsigned int num_arguments_per_4D_tensor ()
 Returns the number of arguments enqueued per 4D tensor object. More...
 
static cl::NDRange gws_from_window (const Window &window)
 Get the global work size given an execution window. More...
 

Detailed Description

Definition at line 176 of file ClElementwiseKernel.h.

Constructor & Destructor Documentation

◆ ClArithmeticKernel()

ClArithmeticKernel ( )
default

Member Function Documentation

◆ ARM_COMPUTE_DISALLOW_COPY_ALLOW_MOVE()

ARM_COMPUTE_DISALLOW_COPY_ALLOW_MOVE ( ClArithmeticKernel  )

◆ configure()

void configure ( const ClCompileContext compile_context,
ArithmeticOperation  op,
ITensorInfo src1,
ITensorInfo src2,
ITensorInfo dst,
const ActivationLayerInfo act_info = ActivationLayerInfo() 
)

Static function to check if given info will lead to a valid configuration of ClArithmeticKernel.

Arithmetic operations.

Parameters
[in]compile_contextThe compile context to be used.
[in]opArithmetic operation to be executed.
[in]src1First source tensor info. Data types supported: U8/QASYMM8/QASYMM8_SIGNED/S16/QSYMM16/F16/S32/F32.
[in]src2Second source tensor info. Data types supported: same as src1.
[in]dstDestination tensor info. Data types supported: same as src1.
[in]act_info(Optional) Activation layer information in case of a fused activation.

Definition at line 471 of file ClElementwiseKernel.cpp.

473 {
474  ARM_COMPUTE_ERROR_ON_NULLPTR(src1, src2, dst);
475  ARM_COMPUTE_ERROR_THROW_ON(ClArithmeticKernel::validate(op, src1, src2, dst, act_info));
476  auto padding_info = get_padding_info({ src1, src2, dst });
477 
478  _op = op;
479  _act_info = act_info;
480  configure_common(compile_context, src1, src2, dst);
482 }
static Status validate(ArithmeticOperation op, const ITensorInfo *src1, const ITensorInfo *src2, const ITensorInfo *dst, const ActivationLayerInfo &act_info=ActivationLayerInfo())
Static function to check if given info will lead to a valid configuration of ClArithmeticKernel.
std::unordered_map< const ITensorInfo *, PaddingSize > get_padding_info(std::initializer_list< const ITensorInfo * > infos)
Stores padding information before configuring a kernel.
Definition: Utils.cpp:489
#define ARM_COMPUTE_ERROR_ON(cond)
If the condition is true then an error message is printed and an exception thrown.
Definition: Error.h:466
#define ARM_COMPUTE_ERROR_THROW_ON(status)
Definition: Error.h:455
bool has_padding_changed(const std::unordered_map< const ITensorInfo *, PaddingSize > &padding_map)
Check if the previously stored padding info has changed after configuring a kernel.
Definition: Utils.cpp:504
#define ARM_COMPUTE_ERROR_ON_NULLPTR(...)
Definition: Validate.h:157

References ARM_COMPUTE_ERROR_ON, ARM_COMPUTE_ERROR_ON_NULLPTR, ARM_COMPUTE_ERROR_THROW_ON, arm_compute::test::validation::dst, arm_compute::get_padding_info(), arm_compute::has_padding_changed(), and ClArithmeticKernel::validate().

◆ validate()

Status validate ( ArithmeticOperation  op,
const ITensorInfo src1,
const ITensorInfo src2,
const ITensorInfo dst,
const ActivationLayerInfo act_info = ActivationLayerInfo() 
)
static

Static function to check if given info will lead to a valid configuration of ClArithmeticKernel.

Parameters
[in]opArithmetic operation to be executed.
[in]src1First source tensor info. Data types supported: U8/QASYMM8/QASYMM8_SIGNED/S16/QSYMM16/F16/S32/F32.
[in]src2Second source tensor info. Data types supported: same as src1.
[in]dstDestination tensor info. Data types supported: same as src1.
[in]act_info(Optional) Activation layer information in case of a fused activation.
Returns
a Status

Definition at line 484 of file ClElementwiseKernel.cpp.

485 {
487  if(op == ArithmeticOperation::DIV)
488  {
489  // Partial integer support S32/F32/F16
490  ARM_COMPUTE_RETURN_ON_ERROR(validate_arguments_divide_operation(src1, src2, dst));
491  ARM_COMPUTE_RETURN_ON_ERROR(validate_and_configure_window_for_division(*src1->clone(), *src2->clone(), *dst->clone()).first);
492  }
493  else if(op == ArithmeticOperation::POWER)
494  {
495  // Power operators doesn't support integer arithmetic
496  ARM_COMPUTE_RETURN_ON_ERROR(validate_arguments_with_float_only_supported_rules(*src1, *src2, *dst));
497  ARM_COMPUTE_RETURN_ON_ERROR(validate_and_configure_window_for_division(*src1->clone(), *src2->clone(), *dst->clone()).first);
498  }
499  else
500  {
501  ARM_COMPUTE_RETURN_ON_ERROR(validate_arguments_with_arithmetic_rules(*src1, *src2, *dst));
502  ARM_COMPUTE_RETURN_ON_ERROR(validate_and_configure_window_for_arithmetic_operators(*src1->clone(), *src2->clone(), *dst->clone()).first);
503  }
504  ARM_COMPUTE_RETURN_ERROR_ON(act_info.enabled() && !is_data_type_float(dst->data_type()));
505 
506  return Status{};
507 }
#define ARM_COMPUTE_RETURN_ON_ERROR(status)
Checks if a status contains an error and returns it.
Definition: Error.h:204
#define ARM_COMPUTE_RETURN_ERROR_ON(cond)
If the condition is true, an error is returned.
Definition: Error.h:296
#define ARM_COMPUTE_RETURN_ERROR_ON_NULLPTR(...)
Definition: Validate.h:159
bool is_data_type_float(DataType dt)
Check if a given data type is of floating point type.
Definition: Utils.h:947

References ARM_COMPUTE_RETURN_ERROR_ON, ARM_COMPUTE_RETURN_ERROR_ON_NULLPTR, ARM_COMPUTE_RETURN_ON_ERROR, ICloneable< T >::clone(), arm_compute::DIV, arm_compute::test::validation::dst, ActivationLayerInfo::enabled(), arm_compute::is_data_type_float(), and arm_compute::POWER.

Referenced by ClArithmeticKernel::configure(), ClPRelu::validate(), ClElementwiseDivision::validate(), ClElementwiseMax::validate(), ClElementwiseMin::validate(), ClElementwiseSquaredDiff::validate(), and ClElementwisePower::validate().


The documentation for this class was generated from the following files: