Compute Library
 19.08
CLChannelCombineKernel Class Reference

Interface for the channel combine kernel. More...

#include <CLChannelCombineKernel.h>

Collaboration diagram for CLChannelCombineKernel:
[legend]

Public Member Functions

 CLChannelCombineKernel ()
 Default constructor. More...
 
 CLChannelCombineKernel (const CLChannelCombineKernel &)=delete
 Prevent instances of this class from being copied (As this class contains pointers) More...
 
CLChannelCombineKerneloperator= (const CLChannelCombineKernel &)=delete
 Prevent instances of this class from being copied (As this class contains pointers) More...
 
 CLChannelCombineKernel (CLChannelCombineKernel &&)=default
 Allow instances of this class to be moved. More...
 
CLChannelCombineKerneloperator= (CLChannelCombineKernel &&)=default
 Allow instances of this class to be moved. More...
 
 ~CLChannelCombineKernel ()=default
 Default destructor. More...
 
void configure (const ICLTensor *plane0, const ICLTensor *plane1, const ICLTensor *plane2, const ICLTensor *plane3, ICLTensor *output)
 Configure function's inputs and outputs. More...
 
void configure (const ICLImage *plane0, const ICLImage *plane1, const ICLImage *plane2, ICLMultiImage *output)
 Configure function's inputs and outputs. 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...
 
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...
 
template<unsigned int dimension_size>
void add_tensor_argument (unsigned &idx, const ICLTensor *tensor, const Window &window)
 
- 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...
 

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...
 

Detailed Description

Interface for the channel combine kernel.

Definition at line 39 of file CLChannelCombineKernel.h.

Constructor & Destructor Documentation

◆ CLChannelCombineKernel() [1/3]

Default constructor.

Definition at line 49 of file CLChannelCombineKernel.cpp.

50  : _planes{ { nullptr } }, _output(nullptr), _output_multi(nullptr), _x_subsampling{ { 1, 1, 1 } }, _y_subsampling{ { 1, 1, 1 } }
51 {
52 }

◆ CLChannelCombineKernel() [2/3]

Prevent instances of this class from being copied (As this class contains pointers)

◆ CLChannelCombineKernel() [3/3]

Allow instances of this class to be moved.

◆ ~CLChannelCombineKernel()

~CLChannelCombineKernel ( )
default

Default destructor.

Member Function Documentation

◆ configure() [1/2]

void configure ( const ICLTensor plane0,
const ICLTensor plane1,
const ICLTensor plane2,
const ICLTensor plane3,
ICLTensor output 
)

Configure function's inputs and outputs.

Parameters
[in]plane0The 2D plane that forms channel 0. Must be of U8 format.
[in]plane1The 2D plane that forms channel 1. Must be of U8 format.
[in]plane2The 2D plane that forms channel 2. Must be of U8 format.
[in]plane3The 2D plane that forms channel 3. Must be of U8 format.
[out]outputThe single planar output tensor.

Definition at line 54 of file CLChannelCombineKernel.cpp.

55 {
56  ARM_COMPUTE_ERROR_ON_NULLPTR(plane0, plane1, plane2, output);
61 
66 
70 
71  const Format output_format = output->info()->format();
72 
73  // Check if horizontal dimension of Y plane is even and validate horizontal sub-sampling dimensions for U and V planes
74  if(Format::YUYV422 == output_format || Format::UYVY422 == output_format)
75  {
76  // Validate Y plane of input and output
77  ARM_COMPUTE_ERROR_ON_TENSORS_NOT_EVEN(output_format, plane0, output);
78 
79  // Validate U and V plane of the input
80  ARM_COMPUTE_ERROR_ON_TENSORS_NOT_SUBSAMPLED(output_format, plane0->info()->tensor_shape(), plane1, plane2);
81  }
82 
83  _planes[0] = plane0;
84  _planes[1] = plane1;
85  _planes[2] = plane2;
86  _planes[3] = nullptr;
87 
88  // Validate the last input tensor only for RGBA format
89  if(Format::RGBA8888 == output_format)
90  {
93 
96 
97  _planes[3] = plane3;
98  }
99 
100  _output = output;
101  _output_multi = nullptr;
102 
103  // Half the processed elements for U and V channels due to horizontal sub-sampling of 2
104  if(Format::YUYV422 == output_format || Format::UYVY422 == output_format)
105  {
106  _x_subsampling[1] = 2;
107  _x_subsampling[2] = 2;
108  }
109 
110  // Create kernel
111  std::string kernel_name = "channel_combine_" + string_from_format(output_format);
112  _kernel = static_cast<cl::Kernel>(CLKernelLibrary::get().create_kernel(kernel_name));
113 
114  // Configure window
115  Window win = calculate_max_window(*output->info(), Steps(num_elems_processed_per_iteration));
116 
117  AccessWindowHorizontal plane0_access(plane0->info(), 0, num_elems_processed_per_iteration);
118  AccessWindowRectangle plane1_access(plane1->info(), 0, 0, num_elems_processed_per_iteration, 1, 1.f / _x_subsampling[1], 1.f / _y_subsampling[1]);
119  AccessWindowRectangle plane2_access(plane2->info(), 0, 0, num_elems_processed_per_iteration, 1, 1.f / _x_subsampling[2], 1.f / _y_subsampling[2]);
120  AccessWindowHorizontal plane3_access(plane3 == nullptr ? nullptr : plane3->info(), 0, num_elems_processed_per_iteration);
121  AccessWindowHorizontal output_access(output->info(), 0, num_elems_processed_per_iteration);
122 
123  update_window_and_padding(win, plane0_access, plane1_access, plane2_access, plane3_access, output_access);
124 
125  ValidRegion valid_region = intersect_valid_regions(plane0->info()->valid_region(),
126  plane1->info()->valid_region(),
127  plane2->info()->valid_region());
128  if(plane3 != nullptr)
129  {
130  valid_region = intersect_valid_regions(plane3->info()->valid_region(), valid_region);
131  }
132  output_access.set_valid_region(win, ValidRegion(valid_region.anchor, output->info()->tensor_shape()));
133 
134  ICLKernel::configure_internal(win);
135 }
A single plane of 32-bit macro pixel of U0, Y0, V0, Y1 byte.
1 channel, 1 U8 per channel
#define ARM_COMPUTE_ERROR_ON_FORMAT_NOT_IN(t,...)
Definition: Validate.h:642
static CLKernelLibrary & get()
Access the KernelLibrary singleton.
Window calculate_max_window(const ValidRegion &valid_region, const Steps &steps=Steps(), bool skip_border=false, BorderSize border_size=BorderSize())
Calculate the maximum window for a given tensor shape and border setting.
Definition: Helpers.cpp:28
3 channels, 1 U8 per channel
#define ARM_COMPUTE_ERROR_ON_TENSOR_NOT_2D(t)
Definition: Validate.h:855
bool update_window_and_padding(Window &win, Ts &&... patterns)
Update window and padding size for each of the access patterns.
Definition: Helpers.h:402
Format
Image colour formats.
Definition: Types.h:52
#define ARM_COMPUTE_ERROR_ON_TENSORS_NOT_SUBSAMPLED(...)
Definition: Validate.h:351
ValidRegion intersect_valid_regions(const Ts &... regions)
Intersect multiple valid regions.
Definition: Helpers.h:503
std::unique_ptr< Kernel > create_kernel()
Helper function to create and return a unique_ptr pointed to a CL/GLES kernel object.
Definition: Helpers.h:86
#define ARM_COMPUTE_ERROR_ON_DATA_TYPE_CHANNEL_NOT_IN(t, c,...)
Definition: Validate.h:789
4 channels, 1 U8 per channel
#define ARM_COMPUTE_ERROR_ON_NULLPTR(...)
Definition: Validate.h:161
A single plane of 32-bit macro pixel of Y0, U0, Y1, V0 bytes.
Coordinates anchor
Anchor for the start of the valid region.
Definition: Types.h:246
#define ARM_COMPUTE_ERROR_ON_TENSORS_NOT_EVEN(...)
Definition: Validate.h:318
const std::string & string_from_format(Format format)
Convert a tensor format into a string.
Definition: Utils.cpp:85

References ValidRegion::anchor, ARM_COMPUTE_ERROR_ON_DATA_TYPE_CHANNEL_NOT_IN, ARM_COMPUTE_ERROR_ON_FORMAT_NOT_IN, ARM_COMPUTE_ERROR_ON_NULLPTR, ARM_COMPUTE_ERROR_ON_TENSOR_NOT_2D, ARM_COMPUTE_ERROR_ON_TENSORS_NOT_EVEN, ARM_COMPUTE_ERROR_ON_TENSORS_NOT_SUBSAMPLED, arm_compute::calculate_max_window(), arm_compute::create_kernel(), ITensorInfo::format(), CLKernelLibrary::get(), ITensor::info(), arm_compute::intersect_valid_regions(), arm_compute::RGB888, arm_compute::RGBA8888, arm_compute::string_from_format(), ITensorInfo::tensor_shape(), arm_compute::U8, arm_compute::update_window_and_padding(), arm_compute::UYVY422, arm_compute::test::validation::valid_region, ITensorInfo::valid_region(), and arm_compute::YUYV422.

◆ configure() [2/2]

void configure ( const ICLImage plane0,
const ICLImage plane1,
const ICLImage plane2,
ICLMultiImage output 
)

Configure function's inputs and outputs.

Parameters
[in]plane0The 2D plane that forms channel 0. Must be of U8 format.
[in]plane1The 2D plane that forms channel 1. Must be of U8 format.
[in]plane2The 2D plane that forms channel 2. Must be of U8 format.
[out]outputThe multi planar output tensor.

Definition at line 137 of file CLChannelCombineKernel.cpp.

138 {
139  ARM_COMPUTE_ERROR_ON_NULLPTR(plane0, plane1, plane2, output);
143 
148 
152 
153  const Format output_format = output->info()->format();
154 
155  // Validate shape of Y plane to be even and shape of sub-sampling dimensions for U and V planes
156  // Perform validation only for formats which require sub-sampling.
157  if(Format::YUV444 != output_format)
158  {
159  // Validate Y plane of input and output
160  ARM_COMPUTE_ERROR_ON_TENSORS_NOT_EVEN(output_format, plane0, output->plane(0));
161 
162  // Validate U and V plane of the input
163  ARM_COMPUTE_ERROR_ON_TENSORS_NOT_SUBSAMPLED(output_format, plane0->info()->tensor_shape(), plane1, plane2);
164 
165  // Validate second plane U (NV12 and NV21 have a UV88 combined plane while IYUV has only the U plane)
166  // MultiImage generates the correct tensor shape but also check in case the tensor shape of planes was changed to a wrong size
167  ARM_COMPUTE_ERROR_ON_TENSORS_NOT_SUBSAMPLED(output_format, plane0->info()->tensor_shape(), output->plane(1));
168 
169  // Validate the last plane V of format IYUV
170  if(Format::IYUV == output_format)
171  {
172  // Validate Y plane of the output
173  ARM_COMPUTE_ERROR_ON_TENSORS_NOT_SUBSAMPLED(output_format, plane0->info()->tensor_shape(), output->plane(2));
174  }
175  }
176 
177  // Set input tensors
178  _planes[0] = plane0;
179  _planes[1] = plane1;
180  _planes[2] = plane2;
181  _planes[3] = nullptr;
182 
183  // Set output tensor
184  _output = nullptr;
185  _output_multi = output;
186 
187  bool has_two_planars = false;
188 
189  // Set sub-sampling parameters for each plane
190  std::string kernel_name;
191  std::set<std::string> build_opts;
192 
193  if(Format::NV12 == output_format || Format::NV21 == output_format)
194  {
195  _x_subsampling = { { 1, 2, 2 } };
196  _y_subsampling = { { 1, 2, 2 } };
197  kernel_name = "channel_combine_NV";
198  build_opts.emplace(Format::NV12 == output_format ? "-DNV12" : "-DNV21");
199  has_two_planars = true;
200  }
201  else
202  {
203  if(Format::IYUV == output_format)
204  {
205  _x_subsampling = { { 1, 2, 2 } };
206  _y_subsampling = { { 1, 2, 2 } };
207  }
208 
209  kernel_name = "copy_planes_3p";
210  build_opts.emplace(Format::IYUV == output_format ? "-DIYUV" : "-DYUV444");
211  }
212 
213  // Create kernel
214  _kernel = static_cast<cl::Kernel>(CLKernelLibrary::get().create_kernel(kernel_name, build_opts));
215 
216  // Configure window
217  Window win = calculate_max_window(*plane0->info(), Steps(num_elems_processed_per_iteration));
218 
219  AccessWindowRectangle input_plane0_access(plane0->info(), 0, 0, num_elems_processed_per_iteration, 1.f);
220  AccessWindowRectangle input_plane1_access(plane1->info(), 0, 0, num_elems_processed_per_iteration, 1.f, 1.f / _x_subsampling[1], 1.f / _y_subsampling[1]);
221  AccessWindowRectangle input_plane2_access(plane2->info(), 0, 0, num_elems_processed_per_iteration, 1.f, 1.f / _x_subsampling[2], 1.f / _y_subsampling[2]);
222  AccessWindowRectangle output_plane0_access(output->plane(0)->info(), 0, 0, num_elems_processed_per_iteration, 1.f, 1.f, 1.f / _y_subsampling[1]);
223  AccessWindowRectangle output_plane1_access(output->plane(1)->info(), 0, 0, num_elems_processed_per_iteration, 1.f, 1.f / _x_subsampling[1], 1.f / _y_subsampling[1]);
224  AccessWindowRectangle output_plane2_access(has_two_planars ? nullptr : output->plane(2)->info(), 0, 0, num_elems_processed_per_iteration, 1.f, 1.f / _x_subsampling[2], 1.f / _y_subsampling[2]);
225 
227  input_plane0_access, input_plane1_access, input_plane2_access,
228  output_plane0_access, output_plane1_access, output_plane2_access);
229 
230  ValidRegion plane0_valid_region = plane0->info()->valid_region();
231  ValidRegion output_plane1_region = has_two_planars ? intersect_valid_regions(plane1->info()->valid_region(), plane2->info()->valid_region()) : plane2->info()->valid_region();
232  output_plane0_access.set_valid_region(win, ValidRegion(plane0_valid_region.anchor, output->plane(0)->info()->tensor_shape()));
233  output_plane1_access.set_valid_region(win, ValidRegion(output_plane1_region.anchor, output->plane(1)->info()->tensor_shape()));
234  output_plane2_access.set_valid_region(win, ValidRegion(plane2->info()->valid_region().anchor, output->plane(2)->info()->tensor_shape()));
235 
236  ICLKernel::configure_internal(win);
237 }
1 channel, 1 U8 per channel
#define ARM_COMPUTE_ERROR_ON_FORMAT_NOT_IN(t,...)
Definition: Validate.h:642
static CLKernelLibrary & get()
Access the KernelLibrary singleton.
A 2 plane YUV format of Luma (Y) and interleaved UV data at 4:2:0 sampling.
Window calculate_max_window(const ValidRegion &valid_region, const Steps &steps=Steps(), bool skip_border=false, BorderSize border_size=BorderSize())
Calculate the maximum window for a given tensor shape and border setting.
Definition: Helpers.cpp:28
A 2 plane YUV format of Luma (Y) and interleaved VU data at 4:2:0 sampling.
#define ARM_COMPUTE_ERROR_ON_TENSOR_NOT_2D(t)
Definition: Validate.h:855
bool update_window_and_padding(Window &win, Ts &&... patterns)
Update window and padding size for each of the access patterns.
Definition: Helpers.h:402
Format
Image colour formats.
Definition: Types.h:52
A 3 plane of 8 bit 4:4:4 sampled Y, U, V planes.
#define ARM_COMPUTE_ERROR_ON_TENSORS_NOT_SUBSAMPLED(...)
Definition: Validate.h:351
ValidRegion intersect_valid_regions(const Ts &... regions)
Intersect multiple valid regions.
Definition: Helpers.h:503
std::unique_ptr< Kernel > create_kernel()
Helper function to create and return a unique_ptr pointed to a CL/GLES kernel object.
Definition: Helpers.h:86
#define ARM_COMPUTE_ERROR_ON_DATA_TYPE_CHANNEL_NOT_IN(t, c,...)
Definition: Validate.h:789
A 3 plane of 8-bit 4:2:0 sampled Y, U, V planes.
#define ARM_COMPUTE_ERROR_ON_NULLPTR(...)
Definition: Validate.h:161
#define ARM_COMPUTE_ERROR_ON_TENSORS_NOT_EVEN(...)
Definition: Validate.h:318

References ValidRegion::anchor, ARM_COMPUTE_ERROR_ON_DATA_TYPE_CHANNEL_NOT_IN, ARM_COMPUTE_ERROR_ON_FORMAT_NOT_IN, ARM_COMPUTE_ERROR_ON_NULLPTR, ARM_COMPUTE_ERROR_ON_TENSOR_NOT_2D, ARM_COMPUTE_ERROR_ON_TENSORS_NOT_EVEN, ARM_COMPUTE_ERROR_ON_TENSORS_NOT_SUBSAMPLED, arm_compute::calculate_max_window(), arm_compute::create_kernel(), MultiImageInfo::format(), CLKernelLibrary::get(), IMultiImage::info(), ITensor::info(), arm_compute::intersect_valid_regions(), arm_compute::IYUV, arm_compute::NV12, arm_compute::NV21, ICLMultiImage::plane(), ITensorInfo::tensor_shape(), arm_compute::U8, arm_compute::update_window_and_padding(), ITensorInfo::valid_region(), and arm_compute::YUV444.

◆ operator=() [1/2]

CLChannelCombineKernel& operator= ( const CLChannelCombineKernel )
delete

Prevent instances of this class from being copied (As this class contains pointers)

◆ operator=() [2/2]

CLChannelCombineKernel& operator= ( CLChannelCombineKernel &&  )
default

Allow instances of this class to be moved.

◆ run()

void run ( const Window window,
cl::CommandQueue &  queue 
)
overridevirtual

Enqueue the OpenCL kernel to process the given window on the passed OpenCL command queue.

Note
The queue is not flushed by this method, and therefore the kernel will not have been executed by the time this method returns.
Parameters
[in]windowRegion on which to execute the kernel. (Must be a valid region of the window returned by window()).
[in,out]queueCommand queue on which to enqueue the kernel.

Implements ICLKernel.

Definition at line 239 of file CLChannelCombineKernel.cpp.

240 {
243 
245  slice.set_dimension_step(Window::DimY, 1);
246 
247  do
248  {
249  // Subsampling in plane 1
250  Window win_sub_plane1(slice);
251  win_sub_plane1.set(Window::DimX, Window::Dimension(win_sub_plane1.x().start() / _x_subsampling[1], win_sub_plane1.x().end() / _x_subsampling[1], win_sub_plane1.x().step() / _x_subsampling[1]));
252  win_sub_plane1.set(Window::DimY, Window::Dimension(win_sub_plane1.y().start() / _y_subsampling[1], win_sub_plane1.y().end() / _y_subsampling[1], 1));
253 
254  // Subsampling in plane 2
255  Window win_sub_plane2(slice);
256  win_sub_plane2.set(Window::DimX, Window::Dimension(win_sub_plane2.x().start() / _x_subsampling[2], win_sub_plane2.x().end() / _x_subsampling[2], win_sub_plane2.x().step() / _x_subsampling[2]));
257  win_sub_plane2.set(Window::DimY, Window::Dimension(win_sub_plane2.y().start() / _y_subsampling[2], win_sub_plane2.y().end() / _y_subsampling[2], 1));
258 
259  unsigned int idx = 0;
260 
261  // Set inputs
262  add_2D_tensor_argument(idx, _planes[0], slice);
263  add_2D_tensor_argument(idx, _planes[1], win_sub_plane1);
264  add_2D_tensor_argument(idx, _planes[2], win_sub_plane2);
265  add_2D_tensor_argument_if((nullptr != _planes[3]), idx, _planes[3], slice);
266 
267  // Set outputs
268  if(nullptr != _output) // Single planar output
269  {
270  add_2D_tensor_argument(idx, _output, slice);
271  }
272  else // Multi-planar output
273  {
274  // Reduce slice in case of subsampling to avoid out-of bounds access
275  slice.set(Window::DimY, Window::Dimension(slice.y().start() / _y_subsampling[1], slice.y().end() / _y_subsampling[1], 1));
276 
277  add_2D_tensor_argument(idx, _output_multi->cl_plane(0), slice);
278  add_2D_tensor_argument(idx, _output_multi->cl_plane(1), win_sub_plane1);
279  add_2D_tensor_argument_if((3 == num_planes_from_format(_output_multi->info()->format())), idx, _output_multi->cl_plane(2), win_sub_plane2);
280 
281  _kernel.setArg(idx++, slice.y().end());
282  }
283 
284  enqueue(queue, *this, slice);
285  }
287 }
Window first_slice_window_2D() const
First 2D slice of the window.
Definition: Window.h:267
const Window & window() const
The maximum window the kernel can be executed on.
Definition: IKernel.cpp:28
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 ...
Definition: ICLKernel.h:145
size_t num_planes_from_format(Format format)
Return the number of planes for a given format.
Definition: Utils.h:444
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:39
virtual ICLImage * cl_plane(unsigned int index)=0
Return a pointer to the requested OpenCL plane of the image.
bool slide_window_slice_2D(Window &slice) const
Slide the passed 2D window slice.
Definition: Window.h:307
static constexpr size_t DimX
Alias for dimension 0 also known as X dimension.
Definition: Window.h:43
virtual const MultiImageInfo * info() const =0
Interface to be implemented by the child class to return the multi-planar image's metadata.
static constexpr size_t DimY
Alias for dimension 1 also known as Y dimension.
Definition: Window.h:45
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.
Definition: ICLKernel.h:134
Format format() const
Colour format of the image.
#define ARM_COMPUTE_ERROR_ON_INVALID_SUBWINDOW(f, s)
Definition: Validate.h:205
#define ARM_COMPUTE_ERROR_ON_UNCONFIGURED_KERNEL(k)
Definition: Validate.h:940
SimpleTensor< T > slice(const SimpleTensor< T > &src, Coordinates starts, Coordinates ends)

References ICLKernel::add_2D_tensor_argument(), ICLKernel::add_2D_tensor_argument_if(), ARM_COMPUTE_ERROR_ON_INVALID_SUBWINDOW, ARM_COMPUTE_ERROR_ON_UNCONFIGURED_KERNEL, ICLMultiImage::cl_plane(), Window::DimX, Window::DimY, Window::Dimension::end(), arm_compute::enqueue(), Window::first_slice_window_2D(), MultiImageInfo::format(), IMultiImage::info(), arm_compute::num_planes_from_format(), Window::set(), arm_compute::test::validation::reference::slice(), Window::slide_window_slice_2D(), Window::Dimension::start(), Window::Dimension::step(), IKernel::window(), Window::x(), and Window::y().


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