21.02
|
OpenCL kernel to multiply two input matrices "A" and "B" and add a martix "C" if provided. More...
#include <CLGEMMMatrixMultiplyKernel.h>
Public Member Functions | |
CLGEMMMatrixMultiplyKernel () | |
Default constructor. More... | |
CLGEMMMatrixMultiplyKernel (const CLGEMMMatrixMultiplyKernel &)=delete | |
Prevent instances of this class from being copied (As this class contains pointers) More... | |
CLGEMMMatrixMultiplyKernel & | operator= (const CLGEMMMatrixMultiplyKernel &)=delete |
Prevent instances of this class from being copied (As this class contains pointers) More... | |
CLGEMMMatrixMultiplyKernel (CLGEMMMatrixMultiplyKernel &&)=default | |
Allow instances of this class to be moved. More... | |
CLGEMMMatrixMultiplyKernel & | operator= (CLGEMMMatrixMultiplyKernel &&)=default |
Allow instances of this class to be moved. More... | |
void | configure (const ICLTensor *input0, const ICLTensor *input1, const ICLTensor *input2, ICLTensor *output, float alpha, float beta=0.f, bool is_interleaved_transposed=true, const GEMMReshapeInfo &reshape_info=GEMMReshapeInfo(), bool fp_mixed_precision=false, const ActivationLayerInfo &activation_info=ActivationLayerInfo()) |
Initialise the kernel's input, output and alpha. More... | |
void | configure (const CLCompileContext &compile_context, const ICLTensor *input0, const ICLTensor *input1, const ICLTensor *input2, ICLTensor *output, float alpha, float beta=0.f, bool is_interleaved_transposed=true, const GEMMReshapeInfo &reshape_info=GEMMReshapeInfo(), bool fp_mixed_precision=false, const ActivationLayerInfo &activation_info=ActivationLayerInfo()) |
Initialise the kernel's input, output and alpha. More... | |
void | run (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... | |
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_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 Window & | window () const |
The maximum window the kernel can be executed on. More... | |
Static Public Member Functions | |
static Status | validate (const ITensorInfo *input0, const ITensorInfo *input1, const ITensorInfo *input2, const ITensorInfo *output, float alpha, float beta, bool is_interleaved_transposed, const GEMMReshapeInfo &reshape_info, GPUTarget gpu_target, bool fp_mixed_precision=false, const ActivationLayerInfo &activation_info=ActivationLayerInfo()) |
Static function to check if given info will lead to a valid configuration of CLGEMMMatrixMultiplyKernel. 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... | |
Data Fields | |
const ICLTensor * | _input0 |
const ICLTensor * | _input1 |
const ICLTensor * | _input2 |
ICLTensor * | _output |
bool | _slide_matrix_b |
bool | _reinterpret_input_as_3d |
bool | _reinterpret_output_as_3d |
bool | _add_bias |
bool | _broadcast_bias |
OpenCL kernel to multiply two input matrices "A" and "B" and add a martix "C" if provided.
All elements of the output matrix will be multiplied by alpha. In case matrix C is passed, it will be added to the previous result. For the matrix C, the broadcast addition is supported if the flag "broadcast_bias" is set in the GEMMReshapeInfo object
input0
and input1
have been reshaped respectively with CLGEMMReshapeLHSMatrixKernel" and CLGEMMReshapeRHSMatrixKernel, the flag is_interleaved_transposed
must be set to trueinput1
tensor must have at least 2 dimensions (matrix) Definition at line 42 of file CLGEMMMatrixMultiplyKernel.h.
Default constructor.
Definition at line 269 of file CLGEMMMatrixMultiplyKernel.cpp.
|
delete |
Prevent instances of this class from being copied (As this class contains pointers)
|
default |
Allow instances of this class to be moved.
void configure | ( | const ICLTensor * | input0, |
const ICLTensor * | input1, | ||
const ICLTensor * | input2, | ||
ICLTensor * | output, | ||
float | alpha, | ||
float | beta = 0.f , |
||
bool | is_interleaved_transposed = true , |
||
const GEMMReshapeInfo & | reshape_info = GEMMReshapeInfo() , |
||
bool | fp_mixed_precision = false , |
||
const ActivationLayerInfo & | activation_info = ActivationLayerInfo() |
||
) |
Initialise the kernel's input, output and alpha.
[in] | input0 | Input tensor containing the Matrix A. Data types supported: F16/F32 |
[in] | input1 | Input tensor containing the Matrix B. Data type supported: same as input0 |
[in] | input2 | Input tensor containing the Matrix C (bias). Can be nullptr. Data type supported: same as input0 |
[out] | output | Output tensor to store the result of matrix multiplication. Data type supported: same as input0 |
[in] | alpha | Weight of the matrix product |
[in] | beta | (Optional) Weight of vector C. Default value is 0. Only beta = 1 is currently supported. |
[in] | is_interleaved_transposed | (Optional) True if input0 and input1 have been reshaped respectively using CLGEMMReshapeLHSMatrixKernel and CLGEMMReshapeRHSMatrixKernel |
[in] | reshape_info | (Optional) GEMM reshape info. If is_interleaved_transposed = true, this object must contain the information to understand how the matrix A and matrix B have been reshaped |
[in] | fp_mixed_precision | (Optional) Use wider accumulators (32 bit instead of 16 for FP16) to improve accuracy |
[in] | activation_info | (Optional) Activation to apply after the matrix multiplication |
Definition at line 275 of file CLGEMMMatrixMultiplyKernel.cpp.
References CLKernelLibrary::get().
void configure | ( | const CLCompileContext & | compile_context, |
const ICLTensor * | input0, | ||
const ICLTensor * | input1, | ||
const ICLTensor * | input2, | ||
ICLTensor * | output, | ||
float | alpha, | ||
float | beta = 0.f , |
||
bool | is_interleaved_transposed = true , |
||
const GEMMReshapeInfo & | reshape_info = GEMMReshapeInfo() , |
||
bool | fp_mixed_precision = false , |
||
const ActivationLayerInfo & | activation_info = ActivationLayerInfo() |
||
) |
Initialise the kernel's input, output and alpha.
[in] | compile_context | The compile context to be used. |
[in] | input0 | Input tensor containing the Matrix A. Data types supported: F16/F32 |
[in] | input1 | Input tensor containing the Matrix B. Data type supported: same as input0 |
[in] | input2 | Input tensor containing the Matrix C (bias). Can be nullptr. Data type supported: same as input0 |
[out] | output | Output tensor to store the result of matrix multiplication. Data type supported: same as input0 |
[in] | alpha | Weight of the matrix product |
[in] | beta | (Optional) Weight of vector C. Default value is 0. Only beta = 1 is currently supported. |
[in] | is_interleaved_transposed | (Optional) True if input0 and input1 have been reshaped respectively using CLGEMMReshapeLHSMatrixKernel and CLGEMMReshapeRHSMatrixKernel |
[in] | reshape_info | (Optional) GEMM reshape info. If is_interleaved_transposed = true, this object must contain the information to understand how the matrix A and matrix B have been reshaped |
[in] | fp_mixed_precision | (Optional) Use wider accumulators (32 bit instead of 16 for FP16) to improve accuracy |
[in] | activation_info | (Optional) Activation to apply after the matrix multiplication |
Definition at line 281 of file CLGEMMMatrixMultiplyKernel.cpp.
References CLGEMMMatrixMultiplyKernel::_add_bias, CLGEMMMatrixMultiplyKernel::_broadcast_bias, CLGEMMMatrixMultiplyKernel::_input0, CLGEMMMatrixMultiplyKernel::_input1, CLGEMMMatrixMultiplyKernel::_input2, CLGEMMMatrixMultiplyKernel::_output, CLGEMMMatrixMultiplyKernel::_reinterpret_input_as_3d, CLGEMMMatrixMultiplyKernel::_reinterpret_output_as_3d, CLGEMMMatrixMultiplyKernel::_slide_matrix_b, ActivationLayerInfo::a(), ActivationLayerInfo::activation(), CLBuildOptions::add_option(), CLBuildOptions::add_option_if(), ARM_COMPUTE_ERROR_ON, ARM_COMPUTE_ERROR_ON_NULLPTR, ARM_COMPUTE_ERROR_THROW_ON, ActivationLayerInfo::b(), arm_compute::BIFROST, arm_compute::create_kernel(), ITensorInfo::data_type(), ITensorInfo::dimension(), ActivationLayerInfo::enabled(), arm_compute::F16, arm_compute::F32, arm_compute::float_to_string_with_full_precision(), arm_compute::get_arch_from_target(), arm_compute::get_cl_type_from_data_type(), arm_compute::get_padding_info(), ICLKernel::get_target(), arm_compute::has_padding_changed(), ITensor::info(), arm_compute::is_data_type_float(), arm_compute::helpers::float_ops::is_one(), arm_compute::helpers::float_ops::is_zero(), kernel_name, arm_compute::lower_string(), ITensorInfo::num_dimensions(), CLBuildOptions::options(), ICLKernel::set_lws_hint(), arm_compute::string_from_activation_func(), arm_compute::string_from_data_type(), arm_compute::support::cpp11::to_string(), and arm_compute::validate_arguments().
|
delete |
Prevent instances of this class from being copied (As this class contains pointers)
|
default |
Allow instances of this class to be moved.
|
overridevirtual |
Enqueue the OpenCL kernel to process the given window on the passed OpenCL command queue.
[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 480 of file CLGEMMMatrixMultiplyKernel.cpp.
References CLGEMMMatrixMultiplyKernel::_add_bias, CLGEMMMatrixMultiplyKernel::_input0, CLGEMMMatrixMultiplyKernel::_input1, CLGEMMMatrixMultiplyKernel::_input2, CLGEMMMatrixMultiplyKernel::_output, CLGEMMMatrixMultiplyKernel::_reinterpret_input_as_3d, CLGEMMMatrixMultiplyKernel::_reinterpret_output_as_3d, CLGEMMMatrixMultiplyKernel::_slide_matrix_b, ICLKernel::add_2D_tensor_argument(), ARM_COMPUTE_ERROR_ON, ARM_COMPUTE_ERROR_ON_INVALID_SUBWINDOW, ARM_COMPUTE_ERROR_ON_UNCONFIGURED_KERNEL, BorderSize::bottom, Window::DimX, Window::DimY, arm_compute::enqueue(), Window::first_slice_window_3D(), ITensor::info(), ICLKernel::lws_hint(), ICLKernel::num_arguments_per_2D_tensor(), ITensorInfo::num_dimensions(), ITensorInfo::padding(), Window::set(), arm_compute::test::validation::reference::slice(), Window::slide_window_slice_3D(), ITensorInfo::strides_in_bytes(), BorderSize::top, and IKernel::window().
|
static |
Static function to check if given info will lead to a valid configuration of CLGEMMMatrixMultiplyKernel.
[in] | input0 | Input tensor containing the Matrix A info. Data types supported: F16/F32 |
[in] | input1 | Input tensor containing the Matrix B info. Data type supported: same as input0 |
[in] | input2 | Input tensor containing the Matrix C (bias) info. Can be nullptr. Data type supported: same as input0 |
[in] | output | Output tensor to store the result of matrix multiplication. Data type supported: same as input0 |
[in] | alpha | Weight of the matrix product |
[in] | beta | Weight of vector C. Default value is 0. Only beta = 1 is currently supported. |
[in] | is_interleaved_transposed | True if input0 and input1 have been reshaped respectively using CLGEMMReshapeLHSMatrixKernel and CLGEMMReshapeRHSMatrixKernel |
[in] | reshape_info | GEMM reshape info. If is_interleaved_transposed = true, this object must contain the information to understand how the matrix A and matrix B have been reshaped |
[in] | gpu_target | GPU Target |
[in] | fp_mixed_precision | (Optional) Use wider accumulators (32 bit instead of 16 for FP16) to improve accuracy |
[in] | activation_info | (Optional) Activation to apply after the matrix multiplication |
Definition at line 458 of file CLGEMMMatrixMultiplyKernel.cpp.
References ARM_COMPUTE_RETURN_ON_ERROR, ARM_COMPUTE_UNUSED, ICloneable< T >::clone(), and arm_compute::validate_arguments().
bool _add_bias |
Definition at line 118 of file CLGEMMMatrixMultiplyKernel.h.
Referenced by CLGEMMMatrixMultiplyKernel::configure(), and CLGEMMMatrixMultiplyKernel::run().
bool _broadcast_bias |
Definition at line 119 of file CLGEMMMatrixMultiplyKernel.h.
Referenced by CLGEMMMatrixMultiplyKernel::configure().
const ICLTensor* _input0 |
Definition at line 111 of file CLGEMMMatrixMultiplyKernel.h.
Referenced by CLGEMMMatrixMultiplyKernel::configure(), and CLGEMMMatrixMultiplyKernel::run().
const ICLTensor* _input1 |
Definition at line 112 of file CLGEMMMatrixMultiplyKernel.h.
Referenced by CLGEMMMatrixMultiplyKernel::configure(), and CLGEMMMatrixMultiplyKernel::run().
const ICLTensor* _input2 |
Definition at line 113 of file CLGEMMMatrixMultiplyKernel.h.
Referenced by CLGEMMMatrixMultiplyKernel::configure(), and CLGEMMMatrixMultiplyKernel::run().
ICLTensor* _output |
Definition at line 114 of file CLGEMMMatrixMultiplyKernel.h.
Referenced by CLGEMMMatrixMultiplyKernel::configure(), and CLGEMMMatrixMultiplyKernel::run().
bool _reinterpret_input_as_3d |
Definition at line 116 of file CLGEMMMatrixMultiplyKernel.h.
Referenced by CLGEMMMatrixMultiplyKernel::configure(), and CLGEMMMatrixMultiplyKernel::run().
bool _reinterpret_output_as_3d |
Definition at line 117 of file CLGEMMMatrixMultiplyKernel.h.
Referenced by CLGEMMMatrixMultiplyKernel::configure(), and CLGEMMMatrixMultiplyKernel::run().
bool _slide_matrix_b |
Definition at line 115 of file CLGEMMMatrixMultiplyKernel.h.
Referenced by CLGEMMMatrixMultiplyKernel::configure(), and CLGEMMMatrixMultiplyKernel::run().