Compute Library
 21.02
CLColorConvertKernel Class Reference

Interface for the color convert kernel. More...

#include <CLColorConvertKernel.h>

Collaboration diagram for CLColorConvertKernel:
[legend]

Public Member Functions

 CLColorConvertKernel ()
 Default constructor. More...
 
 CLColorConvertKernel (const CLColorConvertKernel &)=delete
 Prevent instances of this class from being copied (As this class contains pointers) More...
 
CLColorConvertKerneloperator= (const CLColorConvertKernel &)=delete
 Prevent instances of this class from being copied (As this class contains pointers) More...
 
 CLColorConvertKernel (CLColorConvertKernel &&)=default
 Allow instances of this class to be moved. More...
 
CLColorConvertKerneloperator= (CLColorConvertKernel &&)=default
 Allow instances of this class to be moved. More...
 
 ~CLColorConvertKernel ()=default
 Default destructor. More...
 
void configure (const ICLTensor *input, ICLTensor *output)
 Set the input and output of the kernel. More...
 
void configure (const CLCompileContext &compile_context, const ICLTensor *input, ICLTensor *output)
 Set the input and output of the kernel. More...
 
void configure (const ICLMultiImage *input, ICLImage *output)
 Set the input and output of the kernel. More...
 
void configure (const CLCompileContext &compile_context, const ICLMultiImage *input, ICLImage *output)
 Set the input and output of the kernel. More...
 
void configure (const ICLImage *input, ICLMultiImage *output)
 Set the input and output of the kernel. More...
 
void configure (const CLCompileContext &compile_context, const ICLImage *input, ICLMultiImage *output)
 Set the input and output of the kernel. More...
 
void configure (const ICLMultiImage *input, ICLMultiImage *output)
 Set the input and output of the kernel. More...
 
void configure (const CLCompileContext &compile_context, const ICLMultiImage *input, ICLMultiImage *output)
 Set the input and output of the kernel. 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 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 color convert kernel.

Definition at line 38 of file CLColorConvertKernel.h.

Constructor & Destructor Documentation

◆ CLColorConvertKernel() [1/3]

Default constructor.

Definition at line 42 of file CLColorConvertKernel.cpp.

43  : _input(nullptr), _output(nullptr), _multi_input(nullptr), _multi_output(nullptr)
44 {
45 }

◆ CLColorConvertKernel() [2/3]

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

◆ CLColorConvertKernel() [3/3]

Allow instances of this class to be moved.

◆ ~CLColorConvertKernel()

~CLColorConvertKernel ( )
default

Default destructor.

Member Function Documentation

◆ configure() [1/8]

void configure ( const ICLTensor input,
ICLTensor output 
)

Set the input and output of the kernel.

Parameters
[in]inputSource tensor. Formats supported: RGBA8888/UYVY422/YUYV422/RGB888
[out]outputDestination tensor. Formats supported: RGB888 (if the formats of input are RGBA8888/UYVY422/YUYV422), RGBA8888 (if the formats of input are UYVY422/YUYV422/RGB888/), U8 (if the formats of input is RGB888)

Definition at line 47 of file CLColorConvertKernel.cpp.

References CLKernelLibrary::get().

Referenced by CLColorConvertKernel::configure().

48 {
49  configure(CLKernelLibrary::get().get_compile_context(), input, output);
50 }
static CLKernelLibrary & get()
Access the KernelLibrary singleton.
void configure(const ICLTensor *input, ICLTensor *output)
Set the input and output of the kernel.

◆ configure() [2/8]

void configure ( const CLCompileContext compile_context,
const ICLTensor input,
ICLTensor output 
)

Set the input and output of the kernel.

Parameters
[in]compile_contextThe compile context to be used.
[in]inputSource tensor. Formats supported: RGBA8888/UYVY422/YUYV422/RGB888
[out]outputDestination tensor. Formats supported: RGB888 (if the formats of input are RGBA8888/UYVY422/YUYV422), RGBA8888 (if the formats of input are UYVY422/YUYV422/RGB888/), U8 (if the formats of input is RGB888)

Definition at line 52 of file CLColorConvertKernel.cpp.

References ARM_COMPUTE_ERROR, ARM_COMPUTE_ERROR_ON, ARM_COMPUTE_ERROR_ON_MSG_VAR, arm_compute::calculate_max_window(), arm_compute::create_kernel(), ITensorInfo::data_type(), ITensorInfo::dimension(), ITensorInfo::format(), ITensor::info(), arm_compute::test::validation::input, kernel_name, arm_compute::lower_string(), num_elems_processed_per_iteration, arm_compute::RGB888, arm_compute::RGBA8888, arm_compute::string_from_data_type(), arm_compute::string_from_format(), arm_compute::support::cpp11::to_string(), arm_compute::U8, arm_compute::update_window_and_padding(), arm_compute::UYVY422, ITensorInfo::valid_region(), and arm_compute::YUYV422.

53 {
54  ARM_COMPUTE_ERROR_ON(input == nullptr);
55  ARM_COMPUTE_ERROR_ON(output == nullptr);
56 
57  unsigned int num_elems_processed_per_iteration = 0;
58  switch(input->info()->format())
59  {
60  case Format::RGBA8888:
61  {
62  switch(output->info()->format())
63  {
64  case Format::RGB888:
65  num_elems_processed_per_iteration = 16;
66  break;
67  default:
68  ARM_COMPUTE_ERROR("Not supported");
69  break;
70  }
71  break;
72  }
73  case Format::UYVY422:
74  case Format::YUYV422:
75  {
76  switch(output->info()->format())
77  {
78  case Format::RGB888:
79  case Format::RGBA8888:
80  num_elems_processed_per_iteration = 8;
81  break;
82  default:
83  ARM_COMPUTE_ERROR("Not supported");
84  break;
85  }
86  break;
87  }
88  case Format::RGB888:
89  {
90  switch(output->info()->format())
91  {
92  case Format::RGBA8888:
93  case Format::U8:
94  num_elems_processed_per_iteration = 16;
95  break;
96  default:
97  ARM_COMPUTE_ERROR("Not supported");
98  break;
99  }
100  break;
101  }
102  default:
103  break;
104  }
105  ARM_COMPUTE_ERROR_ON_MSG_VAR(num_elems_processed_per_iteration == 0, "Conversion from %s to %s not supported",
106  string_from_format(input->info()->format()).c_str(),
107  string_from_format(output->info()->format()).c_str());
108 
109  std::stringstream kernel_name;
110 
111  kernel_name << string_from_format(input->info()->format());
112  kernel_name << "_to_";
113  kernel_name << string_from_format(output->info()->format());
114  kernel_name << "_bt709";
115 
116  _input = input;
117  _output = output;
118 
119  // Create kernel
120  _kernel = create_kernel(compile_context, kernel_name.str());
121 
122  // Configure kernel window
123  Window win = calculate_max_window(*input->info(), Steps(num_elems_processed_per_iteration));
126 
127  update_window_and_padding(win, input_access, output_access);
128 
129  output_access.set_valid_region(win, input->info()->valid_region());
130 
131  ICLKernel::configure_internal(win);
132 
133  // Set config_id for enabling LWS tuning
134  _config_id = kernel_name.str();
135  _config_id += "_";
136  _config_id += lower_string(string_from_data_type(input->info()->data_type()));
137  _config_id += "_";
138  _config_id += support::cpp11::to_string(input->info()->dimension(0));
139  _config_id += "_";
140  _config_id += support::cpp11::to_string(input->info()->dimension(1));
141 }
A single plane of 32-bit macro pixel of U0, Y0, V0, Y1 byte.
Window calculate_max_window(const ValidRegion &valid_region, const Steps &steps, bool skip_border, BorderSize border_size)
virtual size_t dimension(size_t index) const =0
Return the size of the requested dimension.
#define ARM_COMPUTE_ERROR(msg)
Print the given message then throw an std::runtime_error.
Definition: Error.h:352
1 channel, 1 U8 per channel
std::string to_string(T &&value)
Convert integer and float values to string.
virtual DataType data_type() const =0
Data type used for each element of the tensor.
#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
#define ARM_COMPUTE_ERROR_ON_MSG_VAR(cond, msg,...)
Definition: Error.h:457
std::string lower_string(const std::string &val)
Lower a given string.
Definition: Utils.cpp:350
virtual ValidRegion valid_region() const =0
Valid region of the tensor.
3 channels, 1 U8 per channel
virtual Format format() const =0
Colour format of the image.
cl::Kernel create_kernel(const CLCompileContext &ctx, const std::string &kernel_name, const std::set< std::string > &build_opts=std::set< std::string >())
Creates an opencl kernel using a compile context.
Definition: CLHelpers.cpp:403
const std::string & string_from_data_type(DataType dt)
Convert a data type identity into a string.
Definition: Utils.cpp:135
bool update_window_and_padding(Window &win, Ts &&... patterns)
Update window and padding size for each of the access patterns.
Definition: WindowHelpers.h:46
Class to describe a number of elements in each dimension.
Definition: Steps.h:40
Implementation of a row access pattern.
std::string kernel_name
virtual ITensorInfo * info() const =0
Interface to be implemented by the child class to return the tensor&#39;s metadata.
4 channels, 1 U8 per channel
unsigned int num_elems_processed_per_iteration
A single plane of 32-bit macro pixel of Y0, U0, Y1, V0 bytes.
Describe a multidimensional execution window.
Definition: Window.h:39
const std::string & string_from_format(Format format)
Convert a tensor format into a string.
Definition: Utils.cpp:76

◆ configure() [3/8]

void configure ( const ICLMultiImage input,
ICLImage output 
)

Set the input and output of the kernel.

Parameters
[in]inputMulti-planar source image. Formats supported: NV12/NV21/IYUV
[out]outputSingle-planar destination image. Formats supported: RGB888/RGBA8888

Definition at line 143 of file CLColorConvertKernel.cpp.

References CLColorConvertKernel::configure(), and CLKernelLibrary::get().

144 {
145  configure(CLKernelLibrary::get().get_compile_context(), input, output);
146 }
static CLKernelLibrary & get()
Access the KernelLibrary singleton.
void configure(const ICLTensor *input, ICLTensor *output)
Set the input and output of the kernel.

◆ configure() [4/8]

void configure ( const CLCompileContext compile_context,
const ICLMultiImage input,
ICLImage output 
)

Set the input and output of the kernel.

Parameters
[in]compile_contextThe compile context to be used.
[in]inputMulti-planar source image. Formats supported: NV12/NV21/IYUV
[out]outputSingle-planar destination image. Formats supported: RGB888/RGBA8888

Definition at line 148 of file CLColorConvertKernel.cpp.

References ValidRegion::anchor, ARM_COMPUTE_ERROR, ARM_COMPUTE_ERROR_ON, ARM_COMPUTE_ERROR_ON_MSG_VAR, ARM_COMPUTE_ERROR_ON_TENSOR_NOT_2D, arm_compute::calculate_max_window(), arm_compute::create_kernel(), ITensorInfo::data_type(), ITensorInfo::dimension(), Window::DimY, MultiImageInfo::format(), ITensorInfo::format(), ITensor::info(), IMultiImage::info(), arm_compute::test::validation::input, arm_compute::intersect_valid_regions(), arm_compute::IYUV, kernel_name, arm_compute::lower_string(), num_elems_processed_per_iteration, arm_compute::NV12, arm_compute::NV21, ICLMultiImage::plane(), arm_compute::RGB888, arm_compute::RGBA8888, Window::set_dimension_step(), arm_compute::string_from_data_type(), arm_compute::string_from_format(), ITensorInfo::tensor_shape(), arm_compute::support::cpp11::to_string(), arm_compute::update_window_and_padding(), and ITensorInfo::valid_region().

149 {
151  ARM_COMPUTE_ERROR_ON(output == nullptr);
152 
153  unsigned int num_elems_processed_per_iteration = 0;
154 
155  switch(input->info()->format())
156  {
157  case Format::NV12:
158  case Format::NV21:
159  case Format::IYUV:
160  {
161  switch(output->info()->format())
162  {
163  case Format::RGB888:
164  case Format::RGBA8888:
165  num_elems_processed_per_iteration = 4;
166  break;
167  default:
168  ARM_COMPUTE_ERROR("Not supported");
169  break;
170  }
171  break;
172  }
173  default:
174  break;
175  }
176  ARM_COMPUTE_ERROR_ON_MSG_VAR(num_elems_processed_per_iteration == 0, "Conversion from %s to %s not supported",
177  string_from_format(input->info()->format()).c_str(),
178  string_from_format(output->info()->format()).c_str());
179 
180  std::stringstream kernel_name;
181 
182  kernel_name << string_from_format(input->info()->format());
183  kernel_name << "_to_";
184  kernel_name << string_from_format(output->info()->format());
185  kernel_name << "_bt709";
186 
187  _multi_input = input;
188  _output = output;
189 
190  // Create kernel
191  _kernel = create_kernel(compile_context, kernel_name.str());
192 
193  // Configure kernel window
194  const bool has_two_planes = (input->info()->format() == Format::NV12) || (input->info()->format() == Format::NV21);
195  const float sub_sampling = (has_two_planes || (input->info()->format() == Format::IYUV)) ? 0.5f : 1;
196 
197  Window win = calculate_max_window(*output->info(), Steps(num_elems_processed_per_iteration));
199 
200  AccessWindowHorizontal plane0_access(input->plane(0)->info(), 0, num_elems_processed_per_iteration);
201  AccessWindowRectangle plane1_access(input->plane(1)->info(), 0, 0, num_elems_processed_per_iteration, 1,
202  sub_sampling, sub_sampling);
203  AccessWindowRectangle plane2_access(has_two_planes ? nullptr : input->plane(2)->info(), 0, 0, num_elems_processed_per_iteration, 1,
204  sub_sampling, sub_sampling);
206 
208  plane0_access, plane1_access, plane2_access,
209  output_access);
210 
211  ValidRegion intersect_region = intersect_valid_regions(input->plane(0)->info()->valid_region(), input->plane(1)->info()->valid_region(),
212  input->plane(2)->info()->valid_region());
213  output_access.set_valid_region(win, ValidRegion(intersect_region.anchor, output->info()->tensor_shape()));
214 
215  ICLKernel::configure_internal(win);
216 
217  // Set config_id for enabling LWS tuning
218  _config_id = kernel_name.str();
219  _config_id += "_";
220  _config_id += lower_string(string_from_data_type(input->plane(0)->info()->data_type()));
221  _config_id += "_";
222  _config_id += support::cpp11::to_string(input->plane(0)->info()->dimension(0));
223  _config_id += "_";
224  _config_id += support::cpp11::to_string(input->plane(0)->info()->dimension(1));
225  _config_id += "_";
226  _config_id += lower_string(string_from_data_type(input->plane(1)->info()->data_type()));
227  _config_id += "_";
228  _config_id += support::cpp11::to_string(input->plane(1)->info()->dimension(0));
229  _config_id += "_";
230  _config_id += support::cpp11::to_string(input->plane(1)->info()->dimension(1));
231 }
Window calculate_max_window(const ValidRegion &valid_region, const Steps &steps, bool skip_border, BorderSize border_size)
virtual size_t dimension(size_t index) const =0
Return the size of the requested dimension.
#define ARM_COMPUTE_ERROR_ON_TENSOR_NOT_2D(t)
Definition: Validate.h:856
#define ARM_COMPUTE_ERROR(msg)
Print the given message then throw an std::runtime_error.
Definition: Error.h:352
std::string to_string(T &&value)
Convert integer and float values to string.
virtual DataType data_type() const =0
Data type used for each element of the tensor.
#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
A 2 plane YUV format of Luma (Y) and interleaved UV data at 4:2:0 sampling.
#define ARM_COMPUTE_ERROR_ON_MSG_VAR(cond, msg,...)
Definition: Error.h:457
std::string lower_string(const std::string &val)
Lower a given string.
Definition: Utils.cpp:350
A 2 plane YUV format of Luma (Y) and interleaved VU data at 4:2:0 sampling.
virtual ValidRegion valid_region() const =0
Valid region of the tensor.
3 channels, 1 U8 per channel
Implementation of a rectangular access pattern.
virtual Format format() const =0
Colour format of the image.
cl::Kernel create_kernel(const CLCompileContext &ctx, const std::string &kernel_name, const std::set< std::string > &build_opts=std::set< std::string >())
Creates an opencl kernel using a compile context.
Definition: CLHelpers.cpp:403
const std::string & string_from_data_type(DataType dt)
Convert a data type identity into a string.
Definition: Utils.cpp:135
bool update_window_and_padding(Window &win, Ts &&... patterns)
Update window and padding size for each of the access patterns.
Definition: WindowHelpers.h:46
virtual const TensorShape & tensor_shape() const =0
Size for each dimension of the tensor.
Class to describe a number of elements in each dimension.
Definition: Steps.h:40
Implementation of a row access pattern.
std::string kernel_name
virtual ITensorInfo * info() const =0
Interface to be implemented by the child class to return the tensor&#39;s metadata.
ValidRegion intersect_valid_regions(const Ts &... regions)
Intersect multiple valid regions.
Definition: WindowHelpers.h:74
IImage * plane(unsigned int index) override
Return a pointer to the requested plane of the image.
virtual const MultiImageInfo * info() const =0
Interface to be implemented by the child class to return the multi-planar image&#39;s metadata...
A 3 plane of 8-bit 4:2:0 sampled Y, U, V planes.
static constexpr size_t DimY
Alias for dimension 1 also known as Y dimension.
Definition: Window.h:45
4 channels, 1 U8 per channel
void set_dimension_step(size_t dimension, int step)
Set the step of a given dimension.
Definition: Window.inl:167
unsigned int num_elems_processed_per_iteration
Container for valid region of a window.
Definition: Types.h:188
Format format() const
Colour format of the image.
Describe a multidimensional execution window.
Definition: Window.h:39
Coordinates anchor
Anchor for the start of the valid region.
Definition: Types.h:260
const std::string & string_from_format(Format format)
Convert a tensor format into a string.
Definition: Utils.cpp:76

◆ configure() [5/8]

void configure ( const ICLImage input,
ICLMultiImage output 
)

Set the input and output of the kernel.

Parameters
[in]inputSingle-planar source image. Formats supported: RGB888/RGBA8888/UYVY422/YUYV422
[out]outputMulti-planar destination image. Formats supported: NV12/IYUV/YUV444 (if the formats of input are RGB888/RGB8888)

Definition at line 233 of file CLColorConvertKernel.cpp.

References CLColorConvertKernel::configure(), and CLKernelLibrary::get().

234 {
235  configure(CLKernelLibrary::get().get_compile_context(), input, output);
236 }
static CLKernelLibrary & get()
Access the KernelLibrary singleton.
void configure(const ICLTensor *input, ICLTensor *output)
Set the input and output of the kernel.

◆ configure() [6/8]

void configure ( const CLCompileContext compile_context,
const ICLImage input,
ICLMultiImage output 
)

Set the input and output of the kernel.

Parameters
[in]compile_contextThe compile context to be used.
[in]inputSingle-planar source image. Formats supported: RGB888/RGBA8888/UYVY422/YUYV422
[out]outputMulti-planar destination image. Formats supported: NV12/IYUV/YUV444 (if the formats of input are RGB888/RGB8888)

Definition at line 238 of file CLColorConvertKernel.cpp.

References ValidRegion::anchor, ARM_COMPUTE_ERROR, ARM_COMPUTE_ERROR_ON, ARM_COMPUTE_ERROR_ON_MSG_VAR, ARM_COMPUTE_ERROR_ON_TENSOR_NOT_2D, arm_compute::calculate_max_window(), arm_compute::create_kernel(), ITensorInfo::data_type(), ITensorInfo::dimension(), Window::DimY, MultiImageInfo::format(), ITensorInfo::format(), IMultiImage::info(), ITensor::info(), arm_compute::test::validation::input, arm_compute::IYUV, kernel_name, arm_compute::lower_string(), num_elems_processed_per_iteration, arm_compute::NV12, arm_compute::NV21, ICLMultiImage::plane(), arm_compute::RGB888, arm_compute::RGBA8888, Window::set_dimension_step(), arm_compute::string_from_data_type(), arm_compute::string_from_format(), ITensorInfo::tensor_shape(), arm_compute::support::cpp11::to_string(), arm_compute::update_window_and_padding(), arm_compute::UYVY422, ITensorInfo::valid_region(), arm_compute::YUV444, and arm_compute::YUYV422.

239 {
241  ARM_COMPUTE_ERROR_ON(output == nullptr);
242 
243  unsigned int num_elems_processed_per_iteration = 0;
244  unsigned int num_elems_read_per_iteration_x = 0;
245 
246  bool has_two_planes = (output->info()->format() == Format::NV12) || (output->info()->format() == Format::NV21);
247  float sub_sampling = (has_two_planes || (output->info()->format() == Format::IYUV)) ? 0.5f : 1;
248 
249  switch(input->info()->format())
250  {
251  case Format::RGB888:
252  case Format::RGBA8888:
253  {
254  switch(output->info()->format())
255  {
256  case Format::NV12:
257  case Format::IYUV:
258  num_elems_processed_per_iteration = 2;
259  num_elems_read_per_iteration_x = 8;
260  break;
261  case Format::YUV444:
262  num_elems_processed_per_iteration = 4;
263  num_elems_read_per_iteration_x = 16;
264  break;
265  default:
266  ARM_COMPUTE_ERROR("Not supported");
267  break;
268  }
269  break;
270  }
271  case Format::UYVY422:
272  case Format::YUYV422:
273  {
274  switch(output->info()->format())
275  {
276  case Format::NV12:
277  case Format::IYUV:
278  num_elems_processed_per_iteration = 8;
279  num_elems_read_per_iteration_x = 8;
280  break;
281  default:
282  ARM_COMPUTE_ERROR("Not supported");
283  break;
284  }
285  break;
286  }
287  default:
288  break;
289  }
290 
291  ARM_COMPUTE_ERROR_ON_MSG_VAR(num_elems_processed_per_iteration == 0, "Conversion from %s to %s not supported",
292  string_from_format(input->info()->format()).c_str(),
293  string_from_format(output->info()->format()).c_str());
294 
295  std::stringstream kernel_name;
296 
297  kernel_name << string_from_format(input->info()->format());
298  kernel_name << "_to_";
299  kernel_name << string_from_format(output->info()->format());
300  kernel_name << "_bt709";
301  _input = input;
302  _multi_output = output;
303 
304  // Create kernel
305  _kernel = create_kernel(compile_context, kernel_name.str());
306 
307  // Configure kernel window
308  Window win = calculate_max_window(*input->info(), Steps(num_elems_processed_per_iteration));
309  if((input->info()->format() != Format::RGB888 || output->info()->format() != Format::YUV444) && (input->info()->format() != Format::RGBA8888 || output->info()->format() != Format::YUV444))
310  {
312  }
313 
314  AccessWindowHorizontal output_plane0_access(output->plane(0)->info(), 0, num_elems_processed_per_iteration);
315  AccessWindowRectangle output_plane1_access(output->plane(1)->info(), 0, 0, num_elems_processed_per_iteration, 1, sub_sampling, sub_sampling);
316  AccessWindowRectangle output_plane2_access(has_two_planes ? nullptr : output->plane(2)->info(), 0, 0,
317  num_elems_processed_per_iteration, 1, sub_sampling, sub_sampling);
318 
319  AccessWindowHorizontal input_access(input->info(), 0, num_elems_read_per_iteration_x);
320 
322  input_access,
323  output_plane0_access,
324  output_plane1_access,
325  output_plane2_access);
326 
327  ValidRegion input_region = input->info()->valid_region();
328 
329  output_plane0_access.set_valid_region(win, ValidRegion(input_region.anchor, output->plane(0)->info()->tensor_shape()));
330  output_plane1_access.set_valid_region(win, ValidRegion(input_region.anchor, output->plane(1)->info()->tensor_shape()));
331  output_plane2_access.set_valid_region(win, ValidRegion(input_region.anchor, output->plane(2)->info()->tensor_shape()));
332 
333  ICLKernel::configure_internal(win);
334 
335  // Set config_id for enabling LWS tuning
336  _config_id = kernel_name.str();
337  _config_id += "_";
338  _config_id += lower_string(string_from_data_type(input->info()->data_type()));
339  _config_id += "_";
340  _config_id += support::cpp11::to_string(input->info()->dimension(0));
341  _config_id += "_";
342  _config_id += support::cpp11::to_string(input->info()->dimension(1));
343 }
A single plane of 32-bit macro pixel of U0, Y0, V0, Y1 byte.
Window calculate_max_window(const ValidRegion &valid_region, const Steps &steps, bool skip_border, BorderSize border_size)
virtual size_t dimension(size_t index) const =0
Return the size of the requested dimension.
#define ARM_COMPUTE_ERROR_ON_TENSOR_NOT_2D(t)
Definition: Validate.h:856
#define ARM_COMPUTE_ERROR(msg)
Print the given message then throw an std::runtime_error.
Definition: Error.h:352
std::string to_string(T &&value)
Convert integer and float values to string.
virtual DataType data_type() const =0
Data type used for each element of the tensor.
#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
A 2 plane YUV format of Luma (Y) and interleaved UV data at 4:2:0 sampling.
#define ARM_COMPUTE_ERROR_ON_MSG_VAR(cond, msg,...)
Definition: Error.h:457
std::string lower_string(const std::string &val)
Lower a given string.
Definition: Utils.cpp:350
A 2 plane YUV format of Luma (Y) and interleaved VU data at 4:2:0 sampling.
virtual ValidRegion valid_region() const =0
Valid region of the tensor.
3 channels, 1 U8 per channel
Implementation of a rectangular access pattern.
virtual Format format() const =0
Colour format of the image.
cl::Kernel create_kernel(const CLCompileContext &ctx, const std::string &kernel_name, const std::set< std::string > &build_opts=std::set< std::string >())
Creates an opencl kernel using a compile context.
Definition: CLHelpers.cpp:403
const std::string & string_from_data_type(DataType dt)
Convert a data type identity into a string.
Definition: Utils.cpp:135
bool update_window_and_padding(Window &win, Ts &&... patterns)
Update window and padding size for each of the access patterns.
Definition: WindowHelpers.h:46
virtual const TensorShape & tensor_shape() const =0
Size for each dimension of the tensor.
Class to describe a number of elements in each dimension.
Definition: Steps.h:40
Implementation of a row access pattern.
std::string kernel_name
A 3 plane of 8 bit 4:4:4 sampled Y, U, V planes.
virtual ITensorInfo * info() const =0
Interface to be implemented by the child class to return the tensor&#39;s metadata.
IImage * plane(unsigned int index) override
Return a pointer to the requested plane of the image.
virtual const MultiImageInfo * info() const =0
Interface to be implemented by the child class to return the multi-planar image&#39;s metadata...
A 3 plane of 8-bit 4:2:0 sampled Y, U, V planes.
static constexpr size_t DimY
Alias for dimension 1 also known as Y dimension.
Definition: Window.h:45
4 channels, 1 U8 per channel
void set_dimension_step(size_t dimension, int step)
Set the step of a given dimension.
Definition: Window.inl:167
unsigned int num_elems_processed_per_iteration
Container for valid region of a window.
Definition: Types.h:188
A single plane of 32-bit macro pixel of Y0, U0, Y1, V0 bytes.
Format format() const
Colour format of the image.
Describe a multidimensional execution window.
Definition: Window.h:39
Coordinates anchor
Anchor for the start of the valid region.
Definition: Types.h:260
const std::string & string_from_format(Format format)
Convert a tensor format into a string.
Definition: Utils.cpp:76

◆ configure() [7/8]

void configure ( const ICLMultiImage input,
ICLMultiImage output 
)

Set the input and output of the kernel.

Parameters
[in]inputMulti-planar source image. Formats supported: NV12/NV21/IYUV
[out]outputMulti-planar destination image. Formats supported: YUV444/IYUV (if the formats of input are NV12/NV21)/NV12 (if the format of input is IYUV)

Definition at line 345 of file CLColorConvertKernel.cpp.

References CLColorConvertKernel::configure(), and CLKernelLibrary::get().

346 {
347  configure(CLKernelLibrary::get().get_compile_context(), input, output);
348 }
static CLKernelLibrary & get()
Access the KernelLibrary singleton.
void configure(const ICLTensor *input, ICLTensor *output)
Set the input and output of the kernel.

◆ configure() [8/8]

void configure ( const CLCompileContext compile_context,
const ICLMultiImage input,
ICLMultiImage output 
)

Set the input and output of the kernel.

Parameters
[in]compile_contextThe compile context to be used.
[in]inputMulti-planar source image. Formats supported: NV12/NV21/IYUV
[out]outputMulti-planar destination image. Formats supported: YUV444/IYUV (if the formats of input are NV12/NV21)/NV12 (if the format of input is IYUV)

Definition at line 350 of file CLColorConvertKernel.cpp.

References ValidRegion::anchor, ARM_COMPUTE_ERROR, ARM_COMPUTE_ERROR_ON_MSG_VAR, arm_compute::calculate_max_window(), ICLMultiImage::cl_plane(), arm_compute::create_kernel(), ITensorInfo::data_type(), ITensorInfo::dimension(), Window::DimY, MultiImageInfo::format(), ITensor::info(), IMultiImage::info(), arm_compute::test::validation::input, arm_compute::intersect_valid_regions(), arm_compute::IYUV, kernel_name, arm_compute::lower_string(), num_elems_processed_per_iteration, arm_compute::NV12, arm_compute::NV21, ICLMultiImage::plane(), Window::set_dimension_step(), arm_compute::string_from_data_type(), arm_compute::string_from_format(), ITensorInfo::tensor_shape(), arm_compute::support::cpp11::to_string(), arm_compute::update_window_and_padding(), ITensorInfo::valid_region(), and arm_compute::YUV444.

351 {
352  unsigned int num_elems_processed_per_iteration = 0;
353  switch(input->info()->format())
354  {
355  case Format::NV12:
356  case Format::NV21:
357  {
358  switch(output->info()->format())
359  {
360  case Format::IYUV:
361  case Format::YUV444:
362  num_elems_processed_per_iteration = 16;
363  break;
364  default:
365  ARM_COMPUTE_ERROR("Not supported");
366  break;
367  }
368  break;
369  }
370  case Format::IYUV:
371  {
372  switch(output->info()->format())
373  {
374  case Format::YUV444:
375  case Format::NV12:
376  num_elems_processed_per_iteration = 16;
377  break;
378  default:
379  ARM_COMPUTE_ERROR("Not supported");
380  break;
381  }
382  break;
383  }
384  default:
385  break;
386  }
387  ARM_COMPUTE_ERROR_ON_MSG_VAR(num_elems_processed_per_iteration == 0, "Conversion from %s to %s not supported",
388  string_from_format(input->info()->format()).c_str(),
389  string_from_format(output->info()->format()).c_str());
390 
391  std::stringstream kernel_name;
392 
393  kernel_name << string_from_format(input->info()->format());
394  kernel_name << "_to_";
395  kernel_name << string_from_format(output->info()->format());
396  kernel_name << "_bt709";
397 
398  _multi_input = input;
399  _multi_output = output;
400 
401  // Create kernel
402  bool has_two_input_planars = (input->info()->format() == Format::NV12) || (input->info()->format() == Format::NV21);
403  bool has_two_output_planars = (output->info()->format() == Format::NV12) || (output->info()->format() == Format::NV21);
404 
405  float sub_sampling_input = (has_two_input_planars || (input->info()->format() == Format::IYUV)) ? 0.5f : 1;
406  float sub_sampling_output = (has_two_output_planars || (output->info()->format() == Format::IYUV)) ? 0.5f : 1;
407 
408  _kernel = create_kernel(compile_context, kernel_name.str());
409 
410  Window win = calculate_max_window(*input->cl_plane(0)->info(), Steps(num_elems_processed_per_iteration));
412 
413  AccessWindowHorizontal input_plane0_access(input->plane(0)->info(), 0, num_elems_processed_per_iteration);
414  AccessWindowRectangle input_plane1_access(input->plane(1)->info(), 0, 0, num_elems_processed_per_iteration, 1,
415  sub_sampling_input, sub_sampling_input);
416  AccessWindowRectangle input_plane2_access(has_two_input_planars ? nullptr : input->plane(2)->info(), 0, 0, num_elems_processed_per_iteration, 1,
417  sub_sampling_input, sub_sampling_input);
418  AccessWindowHorizontal output_plane0_access(output->plane(0)->info(), 0, num_elems_processed_per_iteration);
419  AccessWindowRectangle output_plane1_access(output->plane(1)->info(), 0, 0, num_elems_processed_per_iteration, 1, sub_sampling_output, sub_sampling_output);
420  AccessWindowRectangle output_plane2_access(has_two_output_planars ? nullptr : output->plane(2)->info(), 0, 0,
421  num_elems_processed_per_iteration, 1, sub_sampling_output, sub_sampling_output);
422 
424  input_plane0_access, input_plane1_access, input_plane2_access,
425  output_plane0_access, output_plane1_access, output_plane2_access);
426 
427  ValidRegion intersect_region = intersect_valid_regions(input->plane(0)->info()->valid_region(), input->plane(1)->info()->valid_region(),
428  input->plane(2)->info()->valid_region());
429  output_plane0_access.set_valid_region(win, ValidRegion(intersect_region.anchor, output->plane(0)->info()->tensor_shape()));
430  output_plane1_access.set_valid_region(win, ValidRegion(intersect_region.anchor, output->plane(1)->info()->tensor_shape()));
431  output_plane2_access.set_valid_region(win, ValidRegion(intersect_region.anchor, output->plane(2)->info()->tensor_shape()));
432 
433  ICLKernel::configure_internal(win);
434 
435  // Set config_id for enabling LWS tuning
436  _config_id = kernel_name.str();
437  _config_id += "_";
438  _config_id += lower_string(string_from_data_type(input->plane(0)->info()->data_type()));
439  _config_id += "_";
440  _config_id += support::cpp11::to_string(input->plane(0)->info()->dimension(0));
441  _config_id += "_";
442  _config_id += support::cpp11::to_string(input->plane(0)->info()->dimension(1));
443  _config_id += "_";
444  _config_id += lower_string(string_from_data_type(input->plane(1)->info()->data_type()));
445  _config_id += "_";
446  _config_id += support::cpp11::to_string(input->plane(1)->info()->dimension(0));
447  _config_id += "_";
448  _config_id += support::cpp11::to_string(input->plane(1)->info()->dimension(1));
449 }
Window calculate_max_window(const ValidRegion &valid_region, const Steps &steps, bool skip_border, BorderSize border_size)
virtual size_t dimension(size_t index) const =0
Return the size of the requested dimension.
#define ARM_COMPUTE_ERROR(msg)
Print the given message then throw an std::runtime_error.
Definition: Error.h:352
std::string to_string(T &&value)
Convert integer and float values to string.
virtual DataType data_type() const =0
Data type used for each element of the tensor.
A 2 plane YUV format of Luma (Y) and interleaved UV data at 4:2:0 sampling.
#define ARM_COMPUTE_ERROR_ON_MSG_VAR(cond, msg,...)
Definition: Error.h:457
virtual ICLImage * cl_plane(unsigned int index)=0
Return a pointer to the requested OpenCL plane of the image.
std::string lower_string(const std::string &val)
Lower a given string.
Definition: Utils.cpp:350
A 2 plane YUV format of Luma (Y) and interleaved VU data at 4:2:0 sampling.
virtual ValidRegion valid_region() const =0
Valid region of the tensor.
Implementation of a rectangular access pattern.
cl::Kernel create_kernel(const CLCompileContext &ctx, const std::string &kernel_name, const std::set< std::string > &build_opts=std::set< std::string >())
Creates an opencl kernel using a compile context.
Definition: CLHelpers.cpp:403
const std::string & string_from_data_type(DataType dt)
Convert a data type identity into a string.
Definition: Utils.cpp:135
bool update_window_and_padding(Window &win, Ts &&... patterns)
Update window and padding size for each of the access patterns.
Definition: WindowHelpers.h:46
virtual const TensorShape & tensor_shape() const =0
Size for each dimension of the tensor.
Class to describe a number of elements in each dimension.
Definition: Steps.h:40
Implementation of a row access pattern.
std::string kernel_name
A 3 plane of 8 bit 4:4:4 sampled Y, U, V planes.
virtual ITensorInfo * info() const =0
Interface to be implemented by the child class to return the tensor&#39;s metadata.
ValidRegion intersect_valid_regions(const Ts &... regions)
Intersect multiple valid regions.
Definition: WindowHelpers.h:74
IImage * plane(unsigned int index) override
Return a pointer to the requested plane of the image.
virtual const MultiImageInfo * info() const =0
Interface to be implemented by the child class to return the multi-planar image&#39;s metadata...
A 3 plane of 8-bit 4:2:0 sampled Y, U, V planes.
static constexpr size_t DimY
Alias for dimension 1 also known as Y dimension.
Definition: Window.h:45
void set_dimension_step(size_t dimension, int step)
Set the step of a given dimension.
Definition: Window.inl:167
unsigned int num_elems_processed_per_iteration
Container for valid region of a window.
Definition: Types.h:188
Format format() const
Colour format of the image.
Describe a multidimensional execution window.
Definition: Window.h:39
Coordinates anchor
Anchor for the start of the valid region.
Definition: Types.h:260
const std::string & string_from_format(Format format)
Convert a tensor format into a string.
Definition: Utils.cpp:76

◆ operator=() [1/2]

CLColorConvertKernel& operator= ( const CLColorConvertKernel )
delete

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

◆ operator=() [2/2]

CLColorConvertKernel& operator= ( CLColorConvertKernel &&  )
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.

Reimplemented from ICLKernel.

Definition at line 451 of file CLColorConvertKernel.cpp.

References ICLKernel::add_2D_tensor_argument(), ARM_COMPUTE_ERROR, 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(), ITensor::info(), arm_compute::IYUV, ICLKernel::lws_hint(), ITensorInfo::num_dimensions(), arm_compute::NV12, arm_compute::NV21, 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().

452 {
455 
457 
458  if(nullptr != _input && nullptr != _output)
459  {
460  do
461  {
462  unsigned int idx = 0;
463  add_2D_tensor_argument(idx, _input, slice);
464  add_2D_tensor_argument(idx, _output, slice);
465  enqueue(queue, *this, slice, lws_hint());
466  }
467  while(window.slide_window_slice_2D(slice));
468  }
469  else if(nullptr != _input && nullptr != _multi_output)
470  {
471  Format format = _multi_output->info()->format();
472  do
473  {
474  Window win_uv(slice);
475 
476  if((Format::NV12 == format) || (Format::NV21 == format) || (Format::IYUV == format))
477  {
478  win_uv.set(Window::DimX, Window::Dimension(win_uv.x().start() / 2, win_uv.x().end() / 2, win_uv.x().step() / 2));
479  win_uv.set(Window::DimY, Window::Dimension(win_uv.y().start() / 2, win_uv.y().end() / 2, 1));
480  }
481  unsigned int idx = 0;
482  add_2D_tensor_argument(idx, _input, slice);
483  add_2D_tensor_argument(idx, _multi_output->cl_plane(0), slice);
484  for(int i = 1; i < 3 && (0 != _multi_output->cl_plane(i)->info()->num_dimensions()); ++i)
485  {
486  add_2D_tensor_argument(idx, _multi_output->cl_plane(i), win_uv);
487  }
488  enqueue(queue, *this, slice, lws_hint());
489  }
490  while(window.slide_window_slice_2D(slice));
491  }
492  else if(nullptr != _multi_input && nullptr != _output)
493  {
494  Format format = _multi_input->info()->format();
495  do
496  {
497  Window win_uv(slice);
498 
499  if((Format::NV12 == format) || (Format::NV21 == format) || (Format::IYUV == format))
500  {
501  win_uv.set(Window::DimX, Window::Dimension(win_uv.x().start() / 2, win_uv.x().end() / 2, win_uv.x().step() / 2));
502  win_uv.set(Window::DimY, Window::Dimension(win_uv.y().start() / 2, win_uv.y().end() / 2, 1));
503  }
504 
505  unsigned int idx = 0;
506  add_2D_tensor_argument(idx, _multi_input->cl_plane(0), slice);
507 
508  for(int i = 1; i < 3 && (0 != _multi_input->cl_plane(i)->info()->num_dimensions()); ++i)
509  {
510  add_2D_tensor_argument(idx, _multi_input->cl_plane(i), win_uv);
511  }
512  add_2D_tensor_argument(idx, _output, slice);
513  enqueue(queue, *this, slice, lws_hint());
514  }
515  while(window.slide_window_slice_2D(slice));
516  }
517  else if(nullptr != _multi_input && nullptr != _multi_output)
518  {
519  Format in_format = _multi_input->info()->format();
520  Format out_format = _multi_output->info()->format();
521  do
522  {
523  Window win_in_uv(slice);
524  if((Format::NV12 == in_format) || (Format::NV21 == in_format) || (Format::IYUV == in_format))
525  {
526  win_in_uv.set(Window::DimX, Window::Dimension(win_in_uv.x().start() / 2,
527  win_in_uv.x().end() / 2, win_in_uv.x().step() / 2));
528  win_in_uv.set(Window::DimY, Window::Dimension(win_in_uv.y().start() / 2, win_in_uv.y().end() / 2, 1));
529  }
530  unsigned int idx = 0;
531  add_2D_tensor_argument(idx, _multi_input->cl_plane(0), slice);
532  for(int i = 1; i < 3 && (0 != _multi_input->cl_plane(i)->info()->num_dimensions()); ++i)
533  {
534  add_2D_tensor_argument(idx, _multi_input->cl_plane(i), win_in_uv);
535  }
536 
537  Window win_out_uv(slice);
538  if((Format::NV12 == out_format) || (Format::NV21 == out_format) || (Format::IYUV == out_format))
539  {
540  win_out_uv.set(Window::DimX, Window::Dimension(win_out_uv.x().start() / 2,
541  win_out_uv.x().end() / 2, win_out_uv.x().step() / 2));
542  win_out_uv.set(Window::DimY, Window::Dimension(win_out_uv.y().start() / 2, win_out_uv.y().end() / 2, 1));
543  }
544 
545  add_2D_tensor_argument(idx, _multi_output->cl_plane(0), slice);
546  for(int i = 1; i < 3 && (0 != _multi_output->cl_plane(i)->info()->num_dimensions()); ++i)
547  {
548  add_2D_tensor_argument(idx, _multi_output->cl_plane(i), win_out_uv);
549  }
550  enqueue(queue, *this, slice, lws_hint());
551  }
552  while(window.slide_window_slice_2D(slice));
553  }
554  else
555  {
556  ARM_COMPUTE_ERROR("Not supported");
557  }
558 }
Window first_slice_window_2D() const
First 2D slice of the window.
Definition: Window.h:283
virtual size_t num_dimensions() const =0
The number of dimensions of the tensor (rank)
const Window & window() const
The maximum window the kernel can be executed on.
Definition: IKernel.cpp:28
void enqueue(IGCKernel &kernel, const Window &window, const gles::NDRange &lws=gles::NDRange(1U, 1U, 1U))
Add the kernel to the command queue with the given window.
Definition: IGCKernel.cpp:41
#define ARM_COMPUTE_ERROR(msg)
Print the given message then throw an std::runtime_error.
Definition: Error.h:352
cl::NDRange lws_hint() const
Return the Local-Workgroup-Size hint.
Definition: ICLKernel.h:276
A 2 plane YUV format of Luma (Y) and interleaved UV data at 4:2:0 sampling.
Describe one of the image&#39;s dimensions with a start, end and step.
Definition: Window.h:77
virtual ICLImage * cl_plane(unsigned int index)=0
Return a pointer to the requested OpenCL plane of the image.
A 2 plane YUV format of Luma (Y) and interleaved VU data at 4:2:0 sampling.
bool slide_window_slice_2D(Window &slice) const
Slide the passed 2D window slice.
Definition: Window.h:323
static constexpr size_t DimX
Alias for dimension 0 also known as X dimension.
Definition: Window.h:43
Format
Image colour formats.
Definition: Types.h:54
virtual ITensorInfo * info() const =0
Interface to be implemented by the child class to return the tensor&#39;s metadata.
virtual const MultiImageInfo * info() const =0
Interface to be implemented by the child class to return the multi-planar image&#39;s metadata...
#define ARM_COMPUTE_ERROR_ON_UNCONFIGURED_KERNEL(k)
Definition: Validate.h:941
A 3 plane of 8-bit 4:2:0 sampled Y, U, V planes.
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&#39;s parameters to the object&#39;s kernel&#39;s arguments starting from the index idx...
Definition: ICLKernel.h:148
Format format() const
Colour format of the image.
Describe a multidimensional execution window.
Definition: Window.h:39
#define ARM_COMPUTE_ERROR_ON_INVALID_SUBWINDOW(f, s)
Definition: Validate.h:205
SimpleTensor< T > slice(const SimpleTensor< T > &src, Coordinates starts, Coordinates ends)

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