Class to build OpenCL kernels added from nnfw. More...

#include <CLKernelLibraryEx.h>

Public Member Functions
	CLKernelLibraryEx (const CLKernelLibraryEx &)=delete
	Prevent instances of this class from being copied.

const CLKernelLibraryEx &	operator= (const CLKernelLibraryEx &)=delete
	Prevent instances of this class from being copied.

void	init (std::string kernel_path, cl::Context context, cl::Device device)
	Initialise the kernel library.

void	set_kernel_path (std::string_view kernel_path)
	Set the path that the kernels reside in.

std::string	get_kernel_path ()
	Get the path that the kernels reside in.

std::string	get_program_source (std::string_view program_name)
	Get the source of the selected program.

void	set_context (cl::Context context)
	Set the CL context used to create programs.

cl::Context &	context ()
	Return associated CL context.

void	set_device (cl::Device device)
	Set the CL device for which the programs are created.

cl::Device &	get_device ()
	Gets the CL device for which the programs are created.

std::string	get_device_version ()
	Return the device version.

Kernel	create_kernel (const std::string &kernel_name, const StringSet &build_options_set={}) const
	Create a kernel from the kernel library.

size_t	max_local_workgroup_size (const cl::Kernel &kernel) const
	Find the maximum number of local work items in a workgroup can be supported for the kernel.

cl::NDRange	default_ndrange () const
	Return the default NDRange for the device.

void	clear_programs_cache ()
	Clear the library's cache of binary programs.

const std::map< std::string, cl::Program > &	get_built_programs () const
	Access the cache of built OpenCL programs.

void	add_built_program (const std::string &built_program_name, cl::Program program)
	Add a new built program to the cache.

bool	fp16_supported () const
	Returns true if FP16 is supported by the CL device.

bool	int64_base_atomics_supported () const
	Returns true if int64_base_atomics extension is supported by the CL device.

Static Public Member Functions
static CLKernelLibraryEx &	get ()
	Get the KernelLibrary singleton.

Detailed Description

Class to build OpenCL kernels added from nnfw.

Definition at line 64 of file CLKernelLibraryEx.h.

Constructor & Destructor Documentation

◆ CLKernelLibraryEx()

arm_compute::CLKernelLibraryEx::CLKernelLibraryEx ( const CLKernelLibraryEx & )

delete

Prevent instances of this class from being copied.

Member Function Documentation

◆ add_built_program()

void CLKernelLibraryEx::add_built_program	(	const std::string &	built_program_name,
		cl::Program	program
	)

Add a new built program to the cache.

Parameters

[in]	built_program_name	Name of the program
[in]	program	Built program to add to the cache

Returns: N/A

Definition at line 230 of file CLKernelLibrary.cpp.

{
  _built_programs_map.emplace(built_program_name, program);
}

◆ clear_programs_cache()

void arm_compute::CLKernelLibraryEx::clear_programs_cache ( )

inline

Clear the library's cache of binary programs.

Returns: N/A

Definition at line 205 of file CLKernelLibraryEx.h.

  {
    _programs_map.clear();
    _built_programs_map.clear();
  }

◆ context()

cl::Context & arm_compute::CLKernelLibraryEx::context ( )

inline

Return associated CL context.

Returns: A CL context.

Definition at line 158 of file CLKernelLibraryEx.h.

158{ return _context; }

Referenced by init(), and set_context().

◆ create_kernel()

Kernel CLKernelLibraryEx::create_kernel	(	const std::string &	kernel_name,
		const StringSet &	build_options_set = `{}`
	)		const

Create a kernel from the kernel library.

Parameters

[in]	kernel_name	Kernel name.
[in]	build_options_set	Kernel build options as a set.

Returns: The created kernel.

Definition at line 170 of file CLKernelLibrary.cpp.

{
  // Find which program contains the kernel
  auto kernel_program_it = _kernel_program_map.find(kernel_name);
 
  if (_kernel_program_map.end() == kernel_program_it)
  {
    ARM_COMPUTE_ERROR_VAR("Kernel %s not found in the CLKernelLibrary", kernel_name.c_str());
  }
  std::string concat_str;
 
  if (fp16_supported())
  {
    concat_str += " -DARM_COMPUTE_OPENCL_FP16_ENABLED=1 ";
  }
 
  if (get_cl_version(_device) == CLVersion::CL20)
  {
    concat_str += " -cl-std=CL2.0 ";
  }
  else if (arm_non_uniform_workgroup_supported(_device))
  {
    concat_str += " -cl-arm-non-uniform-work-group-size ";
  }
  else
  {
    ARM_COMPUTE_ERROR("Non uniform workgroup size is not supported!!");
  }
 
  // Check if the program has been built before with same build options.
  const std::string program_name = kernel_program_it->second;
  const std::string build_options = stringify_set(build_options_set) + concat_str;
 
  const std::string built_program_name = program_name + "_" + build_options;
  auto built_program_it = _built_programs_map.find(built_program_name);
 
  cl::Program cl_program;
 
  if (_built_programs_map.end() != built_program_it)
  {
    // If program has been built, retrieve to create kernel from it
    cl_program = built_program_it->second;
  }
  else
  {
    // Get program
    Program program = load_program(program_name);
 
    // Build program
    cl_program = program.build(build_options);
 
    // Add built program to internal map
    _built_programs_map.emplace(built_program_name, cl_program);
  }
 
  // Create and return kernel
  return Kernel(kernel_name, cl_program);
}

References fp16_supported().

◆ default_ndrange()

cl::NDRange CLKernelLibraryEx::default_ndrange ( ) const

Return the default NDRange for the device.

Returns: default NDRangeof the device

Definition at line 332 of file CLKernelLibrary.cpp.

{
  //    GPUTarget   _target = get_target_from_device(_device);
  cl::Device device = cl::Device::getDefault();
  GPUTarget _target = get_target_from_device(device);
  cl::NDRange default_range;
 
  switch (_target)
  {
    case GPUTarget::MIDGARD:
    case GPUTarget::T600:
    case GPUTarget::T700:
    case GPUTarget::T800:
      default_range = cl::NDRange(128u, 1);
      break;
    default:
      default_range = cl::NullRange;
  }
 
  return default_range;
}

◆ fp16_supported()

bool CLKernelLibraryEx::fp16_supported ( ) const

Returns true if FP16 is supported by the CL device.

Returns: true if the CL device supports FP16

Definition at line 236 of file CLKernelLibrary.cpp.

236{ return ::fp16_supported(_device); }

Referenced by create_kernel().

◆ get()

CLKernelLibraryEx & CLKernelLibraryEx::get ( )

static

Get the KernelLibrary singleton.

Returns: The KernelLibrary instance

Definition at line 164 of file CLKernelLibrary.cpp.

{
  static CLKernelLibraryEx _kernel_library;
  return _kernel_library;
}

◆ get_built_programs()

const std::map< std::string, cl::Program > & arm_compute::CLKernelLibraryEx::get_built_programs ( ) const

inline

Access the cache of built OpenCL programs.

Returns: program map data structure of which key is name of kernel and value is kerel source name. (*.cl)

Definition at line 216 of file CLKernelLibraryEx.h.

  {
    return _built_programs_map;
  }

◆ get_device()

cl::Device & arm_compute::CLKernelLibraryEx::get_device ( )

inline

Gets the CL device for which the programs are created.

Returns: A CL device.

Definition at line 171 of file CLKernelLibraryEx.h.

171{ return _device; }

◆ get_device_version()

std::string CLKernelLibraryEx::get_device_version ( )

Return the device version.

Returns: The content of CL_DEVICE_VERSION

Definition at line 354 of file CLKernelLibrary.cpp.

354{ return _device.getInfo<CL_DEVICE_VERSION>(); }

◆ get_kernel_path()

std::string arm_compute::CLKernelLibraryEx::get_kernel_path ( )

inline

Get the path that the kernels reside in.

Returns: the path of kernel files

Definition at line 116 of file CLKernelLibraryEx.h.

116{ return _kernel_path; };

◆ get_program_source()

std::string CLKernelLibraryEx::get_program_source ( std::string_view program_name )

Get the source of the selected program.

Parameters

[in] program_name Program name.

Returns: Source of the selected program.

Definition at line 307 of file CLKernelLibrary.cpp.

{
  const auto program_source_it = _program_source_map.find(program_name);
 
  if (program_source_it == _program_source_map.end())
  {
    ARM_COMPUTE_ERROR_VAR("Embedded program for %s does not exist.", program_name.data());
  }
 
  return program_source_it->second;
}

◆ init()

void arm_compute::CLKernelLibraryEx::init	(	std::string	kernel_path,
		cl::Context	context,
		cl::Device	device
	)

inline

Initialise the kernel library.

Parameters

[in]	kernel_path	Path of the directory from which kernel sources are loaded.
[in]	context	CL context used to create programs.
[in]	device	CL device for which the programs are created.

Returns: N/A

Definition at line 98 of file CLKernelLibraryEx.h.

  {
    _kernel_path = std::move(kernel_path);
    _context = std::move(context);
    _device = std::move(device);
  }

References context().

Referenced by onert::backend::acl_cl::Config::initialize().

◆ int64_base_atomics_supported()

bool CLKernelLibraryEx::int64_base_atomics_supported ( ) const

Returns true if int64_base_atomics extension is supported by the CL device.

Returns: true if the CL device supports int64_base_atomics extension

Definition at line 238 of file CLKernelLibrary.cpp.

{
  return device_supports_extension(_device, "cl_khr_int64_base_atomics");
}

◆ max_local_workgroup_size()

size_t CLKernelLibraryEx::max_local_workgroup_size ( const cl::Kernel & kernel ) const

Find the maximum number of local work items in a workgroup can be supported for the kernel.

Parameters

[in] kernel kernel object

Definition at line 319 of file CLKernelLibrary.cpp.

{
  size_t result;
 
  size_t err = kernel.getWorkGroupInfo(_device, CL_KERNEL_WORK_GROUP_SIZE, &result);
  ARM_COMPUTE_ERROR_ON_MSG(
    err != 0,
    "clGetKernelWorkGroupInfo failed to return the maximum workgroup size for the kernel");
  ARM_COMPUTE_UNUSED(err);
 
  return result;
}

◆ operator=()

const CLKernelLibraryEx & arm_compute::CLKernelLibraryEx::operator= ( const CLKernelLibraryEx & )

delete

Prevent instances of this class from being copied.

◆ set_context()

void arm_compute::CLKernelLibraryEx::set_context ( cl::Context context )

inline

Set the CL context used to create programs.

Note: Setting the context also resets the device to the first one available in the new context.

Parameters

[in] context A CL context.

Returns: N/A

Definition at line 132 of file CLKernelLibraryEx.h.

  {
    _context = std::move(context);
    if (_context.get() == nullptr)
    {
      _device = cl::Device();
    }
    else
    {
      const auto cl_devices = _context.getInfo<CL_CONTEXT_DEVICES>();
 
      if (cl_devices.empty())
      {
        _device = cl::Device();
      }
      else
      {
        _device = cl_devices[0];
      }
    }
  }

References context().

◆ set_device()

void arm_compute::CLKernelLibraryEx::set_device ( cl::Device device )

inline

Set the CL device for which the programs are created.

Parameters

[in] device A CL device.

Returns: N/A

Definition at line 165 of file CLKernelLibraryEx.h.

165{ _device = std::move(device); }

◆ set_kernel_path()

void arm_compute::CLKernelLibraryEx::set_kernel_path ( std::string_view kernel_path )

inline

Set the path that the kernels reside in.

Parameters

[in] kernel_path Path of the directory from which kernel sources are loaded.

Returns: N/A

Definition at line 110 of file CLKernelLibraryEx.h.

110{ _kernel_path = kernel_path; };

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

runtime/compute/ARMComputeEx/arm_compute/core/CL/CLKernelLibraryEx.h
runtime/compute/ARMComputeEx/src/core/CL/CLKernelLibrary.cpp

Public Member Functions

Static Public Member Functions

Detailed Description

Constructor & Destructor Documentation

◆ CLKernelLibraryEx()

Member Function Documentation

◆ add_built_program()

◆ clear_programs_cache()

◆ context()

◆ create_kernel()

◆ default_ndrange()

◆ fp16_supported()

◆ get()

◆ get_built_programs()

◆ get_device()

◆ get_device_version()

◆ get_kernel_path()

◆ get_program_source()

◆ init()

◆ int64_base_atomics_supported()

◆ max_local_workgroup_size()

◆ operator=()

◆ set_context()

◆ set_device()

◆ set_kernel_path()