Compute Library
18.05
|
Kernel for the Vertical pass of a Separable Convolution. More...
#include <CLConvolutionKernel.h>
Public Member Functions | |
void | configure (const ICLTensor *input, ICLTensor *output, const int16_t *conv, uint32_t scale, bool border_undefined, DataType data_type=DataType::S32) |
Initialise the kernel's input, output and border mode. More... | |
BorderSize | border_size () const override |
The size of the border for that kernel. More... | |
Public Member Functions inherited from ICLSimple2DKernel | |
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 ICLSimpleKernel | |
ICLSimpleKernel () | |
Constructor. More... | |
ICLSimpleKernel (const ICLSimpleKernel &)=delete | |
Prevent instances of this class from being copied (As this class contains pointers) More... | |
ICLSimpleKernel & | operator= (const ICLSimpleKernel &)=delete |
Prevent instances of this class from being copied (As this class contains pointers) More... | |
ICLSimpleKernel (ICLSimpleKernel &&)=default | |
Allow instances of this class to be moved. More... | |
ICLSimpleKernel & | operator= (ICLSimpleKernel &&)=default |
Allow instances of this class to be moved. More... | |
~ICLSimpleKernel ()=default | |
Default destructor. More... | |
void | configure (const ICLTensor *input, ICLTensor *output, unsigned int num_elems_processed_per_iteration, bool border_undefined=false, const BorderSize &border_size=BorderSize()) |
Configure the kernel. 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_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_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... | |
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... | |
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<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... | |
const Window & | window () const |
The maximum window the kernel can be executed on. More... | |
Additional Inherited Members | |
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... | |
Kernel for the Vertical pass of a Separable Convolution.
Currently supports 5x5, 7x7, 9x9
Definition at line 114 of file CLConvolutionKernel.h.
|
overridevirtual |
The size of the border for that kernel.
Reimplemented from IKernel.
void configure | ( | const ICLTensor * | input, |
ICLTensor * | output, | ||
const int16_t * | conv, | ||
uint32_t | scale, | ||
bool | border_undefined, | ||
DataType | data_type = DataType::S32 |
||
) |
Initialise the kernel's input, output and border mode.
[in] | input | Source tensor. Data types supported: S16. |
[out] | output | Destination tensor, Data types supported: U8, S16. |
[in] | conv | Convolution matrix to apply to the input tensor. |
[in] | scale | Scale of the convolution matrix. |
[in] | border_undefined | True if the border mode is undefined. False if it's replicate or constant. |
[in] | data_type | Data type to use for intermeidate result. |