24 #ifndef __ARM_COMPUTE_ICLKERNEL_H__ 25 #define __ARM_COMPUTE_ICLKERNEL_H__ 50 template <
unsigned int dimension_size>
51 constexpr
static unsigned int num_arguments_per_array()
53 return num_arguments_per_tensor<dimension_size>();
59 template <
unsigned int dimension_size>
60 constexpr
static unsigned int num_arguments_per_tensor()
62 return 2 + 2 * dimension_size;
90 add_array_argument<T, 1>(idx, array, strides, num_dimensions,
window);
100 add_tensor_argument<1>(idx, tensor,
window);
110 add_tensor_argument<2>(idx, tensor,
window);
120 add_tensor_argument<3>(idx, tensor,
window);
130 add_tensor_argument<4>(idx, tensor,
window);
138 return num_arguments_per_array<1>();
146 return num_arguments_per_tensor<1>();
154 return num_arguments_per_tensor<2>();
162 return num_arguments_per_tensor<3>();
170 return num_arguments_per_tensor<4>();
185 template <
typename T>
188 _kernel.setArg(idx++, value);
272 template <
typename T,
unsigned int dimension_size>
273 void add_array_argument(
unsigned int &idx,
const ICLArray<T> *array,
const Strides &strides,
unsigned int num_dimensions,
const Window &window);
280 template <
unsigned int dimension_size>
281 void add_tensor_argument(
unsigned int &idx,
const ICLTensor *tensor,
const Window &window);
285 cl::NDRange _lws_hint;
287 std::string _config_id;
288 size_t _max_workgroup_size;
314 template <
typename T,
unsigned int dimension_size>
315 void ICLKernel::add_array_argument(
unsigned &idx,
const ICLArray<T> *array,
const Strides &strides,
unsigned int num_dimensions,
const Window &window)
320 unsigned int offset_first_element = 0;
322 for(
unsigned int n = 0; n < num_dimensions; ++n)
324 offset_first_element += window[n].start() * strides[n];
327 unsigned int idx_start = idx;
328 _kernel.setArg(idx++, array->
cl_buffer());
330 for(
unsigned int dimension = 0; dimension < dimension_size; dimension++)
332 _kernel.setArg<cl_uint>(idx++, strides[dimension]);
333 _kernel.setArg<cl_uint>(idx++, strides[dimension] * window[dimension].step());
336 _kernel.setArg<cl_uint>(idx++, offset_first_element);
339 "add_%dD_array_argument() is supposed to add exactly %d arguments to the kernel", dimension_size, num_arguments_per_array<dimension_size>());
static constexpr unsigned int num_arguments_per_1D_tensor()
Returns the number of arguments enqueued per 1D tensor object.
+
Common information for all the kernels.
+
static constexpr unsigned int num_arguments_per_1D_array()
Returns the number of arguments enqueued per 1D array object.
+
cl::NDRange lws_hint() const
Return the Local-Workgroup-Size hint.
+
cl::Kernel & kernel()
Returns a reference to the OpenCL kernel of this object.
+
void set_lws_hint(const cl::NDRange &lws_hint)
Set the Local-Workgroup-Size hint.
+
void add_argument(unsigned int &idx, T value)
Add the passed parameters to the object's kernel's arguments starting from the index idx...
+
+
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...
+
#define ARM_COMPUTE_ERROR_ON(cond)
If the condition is true then an error message is printed and an exception thrown.
-
Common interface for all the OpenCL kernels.
-
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...
-
-
unsigned int num_arguments_per_2D_tensor() const
Returns the number of arguments enqueued per 2D tensor object.
-
+
Common interface for all the OpenCL kernels.
+
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...
+
This file contains all available output stages for GEMMLowp on OpenCL.
+
static CLKernelLibrary & get()
Access the KernelLibrary singleton.
+
Interface for OpenCL Array.
+
static constexpr unsigned int num_arguments_per_3D_tensor()
Returns the number of arguments enqueued per 3D tensor object.
+
#define ARM_COMPUTE_UNUSED(...)
To avoid unused variables warnings.
+
GPUTarget get_target() const
Get the targeted GPU architecture.
+
static constexpr unsigned int num_arguments_per_2D_tensor()
Returns the number of arguments enqueued per 2D tensor object.
+
static constexpr unsigned int num_arguments_per_4D_tensor()
Returns the number of arguments enqueued per 4D tensor object.
+
const std::string & config_id() const
Get the configuration ID.
+
static cl::NDRange gws_from_window(const Window &window)
Get the global work size given an execution window.
+
Strides of an item in bytes.
+
void enqueue(cl::CommandQueue &queue, ICLKernel &kernel, const Window &window, const cl::NDRange &lws_hint=CLKernelLibrary::get().default_ndrange())
Add the kernel to the command queue with the given window.
+
virtual void run(const Window &window, cl::CommandQueue &queue)=0
Enqueue the OpenCL kernel to process the given window on the passed OpenCL command queue...
-
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...
-
void enqueue(cl::CommandQueue &queue, ICLKernel &kernel, const Window &window, const cl::NDRange &lws_hint=cl::Range_128_1)
Add the kernel to the command queue with the given window.
+
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...
Interface for OpenCL tensor.
-
+
+
GPUTarget
Available GPU Targets.
+
+
-
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...
+
+
size_t get_max_workgroup_size()
Get the maximum workgroup size for the device the CLKernelLibrary uses.
+
void set_target(GPUTarget target)
Set the targeted GPU architecture.
+
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...
+
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...
+
virtual const cl::Buffer & cl_buffer() const =0
Interface to be implemented by the child class to return a reference to the OpenCL buffer containing ...
const Window & window() const
The maximum window the kernel can be executed on.
Describe a multidimensional execution window.
+
#define ARM_COMPUTE_ERROR_ON_MSG(cond,...)