Compute Library
 22.08
ClCompositeKernel Class Referencefinal

#include <ClCompositeKernel.h>

Collaboration diagram for ClCompositeKernel:
[legend]

Public Member Functions

void configure (const opencl::ClCompileContext &, const ClKernelCode &)
 
virtual void run_composite_op (ITensorPack &tensors, const Window &window, cl::CommandQueue &queue, const ClExecutionDescriptor &exec_desc) override
 Run the composite 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...
 
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...
 
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...
 

Additional Inherited Members

- Static Public Member Functions inherited from ICLKernel
static constexpr unsigned int num_arguments_per_3d_tensor_nhw ()
 Returns the number of arguments enqueued per NHW 3D Tensor object. More...
 
static constexpr unsigned int num_arguments_per_4d_tensor_nhwc ()
 Returns the number of arguments enqueued per NHWC 4D Tensor object. More...
 
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 42 of file ClCompositeKernel.h.

Member Function Documentation

◆ configure()

void configure ( const opencl::ClCompileContext compile_ctx,
const ClKernelCode cl_code 
)

Definition at line 42 of file ClCompositeKernel.cpp.

References ClKernelCode::arguments, ARM_COMPUTE_ERROR, BorderSize::bottom, ClKernelCode::build_options, ICLTensor::cl_buffer(), ClKernelCode::code, ClKernelCode::config_id, arm_compute::create_image2d_from_buffer(), CLCompileContext::create_kernel(), ITensorInfo::data_type(), ITensorInfo::dimension(), CLKernelLibrary::get(), ClKernelLibrary::get(), arm_compute::experimental::dynamic_fusion::Image, arm_compute::experimental::dynamic_fusion::Image_3D, arm_compute::experimental::dynamic_fusion::Image_3D_Export_To_ClImage2D, arm_compute::experimental::dynamic_fusion::Image_Export_To_ClImage2D, arm_compute::experimental::dynamic_fusion::Image_Reinterpret_As_3D, ITensor::info(), ClKernelLibrary::kernel_path(), ClKernelCode::name, CLBuildOptions::options(), ITensorInfo::padding(), arm_compute::experimental::dynamic_fusion::Scalar, ITensorInfo::strides_in_bytes(), arm_compute::experimental::dynamic_fusion::Tensor_3D, arm_compute::experimental::dynamic_fusion::Tensor_4D, arm_compute::experimental::dynamic_fusion::Tensor_4D_t_Buffer, arm_compute::experimental::dynamic_fusion::Tensor_4D_t_Image, ClKernelArgDescriptor::tensor_arg_type, ITensorInfo::tensor_shape(), BorderSize::top, TensorShape::total_size_upper(), arm_compute::experimental::dynamic_fusion::Vector, and ClKernelCode::window.

43 {
44  // Create kernel from kernel source string
45  opencl::ClKernelLibrary &klib = opencl::ClKernelLibrary::get();
46  _kernel = static_cast<cl::Kernel>(compile_ctx.create_kernel(cl_code.name,
47  "" /* Program name: Used to as part of a unique string for built kernel cache. Not needed */,
48  cl_code.code,
49  klib.kernel_path() /* Kernel path: Used in cases of embedded kernels */,
50  cl_code.build_options.options(),
51  false /* Is source binary */));
52 
53  // Configure execution window
54  IClKernel::configure_internal(cl_code.window);
55 
56  // Set config id for lws tuning
57  _config_id = cl_code.config_id;
58 
59  // Set kernel arguments
60  _arguments = cl_code.arguments;
61 }
static ClKernelLibrary & get()
Access the KernelLibrary singleton.

◆ run_composite_op()

void run_composite_op ( ITensorPack tensors,
const Window window,
cl::CommandQueue &  queue,
const ClExecutionDescriptor exec_desc 
)
overridevirtual

Run the composite kernel.

Note
The slots / keys in ITensorPack are the argument Ids of the tensors in blueprint
Parameters
tensorsITensorPack object containing run-time tensor memories
windowExecution window
queueOpenCL Command queue
exec_descDescriptor containing execution information

Reimplemented from ICLKernel.

Definition at line 155 of file ClCompositeKernel.cpp.

References ClKernelArgDescriptor::arg_id, ARM_COMPUTE_ERROR_ON, ARM_COMPUTE_ERROR_ON_INVALID_SUBWINDOW, ARM_COMPUTE_ERROR_ON_NULLPTR, ARM_COMPUTE_ERROR_ON_UNCONFIGURED_KERNEL, ARM_COMPUTE_UNUSED, Window::DimX, Window::DimY, arm_compute::enqueue(), Window::first_slice_window_3D(), ITensorPack::get_tensor(), ITensor::info(), Window::set(), ClExecutionDescriptor::skip_sliding_window, arm_compute::test::validation::reference::slice(), ClKernelArgDescriptor::slide_along_dimz, Window::slide_window_slice_3D(), ITensorInfo::strides_in_bytes(), and IKernel::window().

156 {
157  ARM_COMPUTE_UNUSED(exec_desc);
160 
162  // Don't slice matrix along the z dimension if matrix has just 2 dimensions and matrix A more than 2
163  // This scenario can happen when the matrix multiplication is used to perform a convolution operation
164  Window slice_fixed_z = slice;
165  slice_fixed_z.set(Window::DimX, Window::Dimension(0, 1, 1));
166  slice_fixed_z.set(Window::DimY, Window::Dimension(0, 1, 1));
167 
168  unsigned int idx = 0;
169  do
170  {
171  // Set kernel arguments
172  Window arg_slice = slice;
173  // CLImages created from tensor arguments. Need to be retained until enqueue
174  std::vector<cl::Image2D> cl_images;
175  for(auto id_arg : _arguments)
176  {
177  const auto arg = id_arg.second;
178  auto tensor = utils::cast::polymorphic_downcast<ICLTensor *>(tensors.get_tensor(arg.arg_id));
180  ARM_COMPUTE_ERROR_ON_NULLPTR(tensor->info());
181  if(!arg.slide_along_dimz)
182  {
183  // The stride_z for matrix must be zero if we do not slice
184  ARM_COMPUTE_ERROR_ON(tensor->info()->strides_in_bytes()[3] != 0);
185  arg_slice = slice_fixed_z;
186  }
187  add_tensor_argument(idx, arg, tensor, arg_slice, cl_images);
188  }
189 
190  // Dispatch kernel
191  bool use_dummy_work_items = false;
192  enqueue(queue, *this, slice, lws_hint(), use_dummy_work_items);
193  }
194  while(!exec_desc.skip_sliding_window && window.slide_window_slice_3D(slice));
195 }
const Window & window() const
The maximum window the kernel can be executed on.
Definition: IKernel.cpp:28
void enqueue(cl::CommandQueue &queue, ICLKernel &kernel, const Window &window, const cl::NDRange &lws_hint=CLKernelLibrary::get().default_ndrange(), bool use_dummy_work_items=false)
Add the kernel to the command queue with the given window.
Definition: ICLKernel.cpp:32
cl::NDRange lws_hint() const
Return the Local-Workgroup-Size hint.
Definition: ICLKernel.h:384
#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
static constexpr size_t DimX
Alias for dimension 0 also known as X dimension.
Definition: Window.h:43
#define ARM_COMPUTE_UNUSED(...)
To avoid unused variables warnings.
Definition: Error.h:152
bool slide_window_slice_3D(Window &slice) const
Slide the passed 3D window slice.
Definition: Window.h:349
#define ARM_COMPUTE_ERROR_ON_UNCONFIGURED_KERNEL(k)
Definition: Validate.h:915
static constexpr size_t DimY
Alias for dimension 1 also known as Y dimension.
Definition: Window.h:45
#define ARM_COMPUTE_ERROR_ON_NULLPTR(...)
Definition: Validate.h:157
Window first_slice_window_3D() const
First 3D slice of the window.
Definition: Window.h:305
#define ARM_COMPUTE_ERROR_ON_INVALID_SUBWINDOW(f, s)
Definition: Validate.h:201
SimpleTensor< T > slice(const SimpleTensor< T > &src, Coordinates starts, Coordinates ends)

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