24.02.1
|
OpenCL kernel to multiply matrices with QASYMM8/QASYMM8_SIGNED data types when only the input matrix RHS (src1) has been reshaped using the MMUL instruction. More...
#include <ClGemmLowpMatrixMultiplyReshapedOnlyRhsMMULKernel.h>
Public Member Functions | |
ClGemmLowpMatrixMultiplyReshapedOnlyRhsMMULKernel () | |
ARM_COMPUTE_DISALLOW_COPY_ALLOW_MOVE (ClGemmLowpMatrixMultiplyReshapedOnlyRhsMMULKernel) | |
void | configure (const CLCompileContext &compile_context, const ITensorInfo *src0, const ITensorInfo *src1, ITensorInfo *dst, const GEMMKernelInfo &gemm_info, ITensorInfo *vector_sum_col=nullptr, const ITensorInfo *vector_sum_row=nullptr, ITensorInfo *bias=nullptr, ITensorInfo *output_multipliers=nullptr, ITensorInfo *output_shifts=nullptr) |
Initialise the kernel's source and destination. More... | |
void | run_op (ITensorPack &tensors, const Window &window, cl::CommandQueue &queue) override |
Enqueue the OpenCL kernel to process the given window on the passed OpenCL command queue. More... | |
Public Member Functions inherited from ICLKernel | |
ICLKernel () | |
Constructor. More... | |
cl::Kernel & | kernel () |
Returns a reference to the OpenCL kernel of this object. More... | |
CLKernelType | type () const |
Returns the CL kernel type. 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... | |
void | add_5D_tensor_argument (unsigned int &idx, const ICLTensor *tensor, const Window &window) |
Add the passed 5D tensor's parameters to the object's kernel's arguments starting from the index idx. More... | |
void | add_3d_tensor_nhw_argument (unsigned int &idx, const ICLTensor *tensor) |
Add the passed NHW 3D tensor's parameters to the object's kernel's arguments by passing strides, dimensions and the offset to the first valid element in bytes. More... | |
void | add_4d_tensor_nhwc_argument (unsigned int &idx, const ICLTensor *tensor) |
Add the passed NHWC 4D tensor's parameters to the object's kernel's arguments by passing strides, dimensions and the offset to the first valid element in bytes. 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... | |
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... | |
cl::NDRange | get_cached_gws () const |
Get the cached gws used to enqueue this kernel. More... | |
void | cache_gws (const cl::NDRange &gws) |
Cache the latest gws used to enqueue this kernel. 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 Window & | window () 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 (const ITensorInfo *src0, const ITensorInfo *src1, const ITensorInfo *dst, const GEMMKernelInfo &gemm_info, const ITensorInfo *vector_sum_col=nullptr, const ITensorInfo *vector_sum_row=nullptr, const ITensorInfo *bias=nullptr, const ITensorInfo *output_multipliers=nullptr, const ITensorInfo *output_shifts=nullptr) |
Static function to check if given info will lead to a valid configuration. More... | |
Static Public Member Functions inherited from ICLKernel | |
constexpr static unsigned int | num_arguments_per_3d_tensor_nhw () |
Returns the number of arguments enqueued per NHW 3D Tensor object. More... | |
constexpr static unsigned int | num_arguments_per_4d_tensor_nhwc () |
Returns the number of arguments enqueued per NHWC 4D Tensor object. More... | |
constexpr static unsigned int | num_arguments_per_1D_array () |
Returns the number of arguments enqueued per 1D array object. More... | |
constexpr static unsigned int | num_arguments_per_1D_tensor () |
Returns the number of arguments enqueued per 1D tensor object. More... | |
constexpr static unsigned int | num_arguments_per_2D_tensor () |
Returns the number of arguments enqueued per 2D tensor object. More... | |
constexpr static unsigned int | num_arguments_per_3D_tensor () |
Returns the number of arguments enqueued per 3D tensor object. More... | |
constexpr static 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, bool use_dummy_work_items) |
Get the global work size given an execution window. More... | |
OpenCL kernel to multiply matrices with QASYMM8/QASYMM8_SIGNED data types when only the input matrix RHS (src1) has been reshaped using the MMUL instruction.
Definition at line 43 of file ClGemmLowpMatrixMultiplyReshapedOnlyRhsMMULKernel.h.
Definition at line 314 of file ClGemmLowpMatrixMultiplyReshapedOnlyRhsMMULKernel.cpp.
References arm_compute::GEMM.
ARM_COMPUTE_DISALLOW_COPY_ALLOW_MOVE | ( | ClGemmLowpMatrixMultiplyReshapedOnlyRhsMMULKernel | ) |
void configure | ( | const CLCompileContext & | compile_context, |
const ITensorInfo * | src0, | ||
const ITensorInfo * | src1, | ||
ITensorInfo * | dst, | ||
const GEMMKernelInfo & | gemm_info, | ||
ITensorInfo * | vector_sum_col = nullptr , |
||
const ITensorInfo * | vector_sum_row = nullptr , |
||
ITensorInfo * | bias = nullptr , |
||
ITensorInfo * | output_multipliers = nullptr , |
||
ITensorInfo * | output_shifts = nullptr |
||
) |
Initialise the kernel's source and destination.
[in] | compile_context | The compile context to be used. |
[in] | src0 | Input tensor containing the LHS matrix. Data type supported: QASYMM8/QASYMM8_SIGNED |
[in] | src1 | Input tensor containing the RHS reshaped matrix. Data type supported: same as src0 |
[out] | dst | Destination tensor. Data type supported: QASYMM8/QASYMM8_SIGNED/S32. |
[in] | gemm_info | GEMM information used to retrieve the original dimensions of the input matrices, output stage information and RHS/LHS info. lhs_info.m0: 1,2,4 Only the following values are supported for RHS info: rhs_info.n0: 1,4,8 rhs_info.k0: same as lhs_info.k0: 4 rhs_info.transpose: true |
[in] | vector_sum_col | (Optional) Input row-vector of sums of all the entries in each column of matrix B. Note: vector_sum_col can be a nullptr in case a_offset = 0. Data type supported: S32 |
[in] | vector_sum_row | (Optional) Input row-vector of sums of all the entries in each row of matrix A. Note: vector_sum_row can be a nullptr in case b_offset = 0. Data type supported: S32 |
[in] | bias | (Optional) Biases tensor. Can be a nullptr if the addition of biases is not required. Biases are 1D tensor with dimensions [OFM] or same dimensionality as dst if gemm_info.broadcast_bias is false. Data type supported: S32. |
[in] | output_multipliers | (Optional) Output multipliers tensor. Supported data types: S32. |
[in] | output_shifts | (Optional) Output shifts tensor. Supported data types: S32. |
Definition at line 319 of file ClGemmLowpMatrixMultiplyReshapedOnlyRhsMMULKernel.cpp.
References GEMMKernelInfo::a_offset, ARM_COMPUTE_ERROR_ON_NULLPTR, ARM_COMPUTE_ERROR_THROW_ON, GEMMKernelInfo::b_offset, bias, arm_compute::test::validation::dst, arm_compute::get_padding_info(), GEMMKernelInfo::lhs_info, output_stage, GEMMKernelInfo::output_stage, GEMMKernelInfo::rhs_info, and arm_compute::cpu::kernels::validate_arguments().
|
overridevirtual |
Enqueue the OpenCL kernel to process the given window on the passed OpenCL command queue.
[in] | tensors | A vector containing the tensors to operato on. |
[in] | window | Region on which to execute the kernel. (Must be a valid region of the window returned by window()). |
[in,out] | queue | Command queue on which to enqueue the kernel. |
Reimplemented from ICLKernel.
Definition at line 462 of file ClGemmLowpMatrixMultiplyReshapedOnlyRhsMMULKernel.cpp.
References arm_compute::ACL_DST, arm_compute::ACL_SRC_0, arm_compute::ACL_SRC_1, arm_compute::ACL_SRC_2, arm_compute::ACL_VEC_COL_SUM, arm_compute::ACL_VEC_ROW_SUM, ARM_COMPUTE_ERROR_ON, ARM_COMPUTE_ERROR_ON_INVALID_SUBWINDOW, ARM_COMPUTE_ERROR_ON_NULLPTR, ARM_COMPUTE_ERROR_ON_UNCONFIGURED_KERNEL, arm_compute::test::validation::dst, Window::first_slice_window_3D(), ITensorPack::get_const_tensor(), ITensorPack::get_tensor(), arm_compute::test::validation::reference::slice(), and IKernel::window().
|
static |
Static function to check if given info will lead to a valid configuration.
Similar to ClGemmLowpMatrixMultiplyReshapedOnlyRhsMMULKernel::configure()
Definition at line 436 of file ClGemmLowpMatrixMultiplyReshapedOnlyRhsMMULKernel.cpp.
References ARM_COMPUTE_RETURN_ON_ERROR, bias, ICloneable< T >::clone(), arm_compute::test::validation::dst, arm_compute::cpu::kernels::validate_and_configure_window(), and arm_compute::cpu::kernels::validate_arguments().