41 _build_opts.emplace(std::move(option));
59 _build_opts.insert(options.begin(), options.end());
76 : _context(), _device(), _is_binary(false), _name(), _source(), _binary()
81 : _context(
std::move(context)), _device(), _is_binary(false), _name(
std::move(name)), _source(
std::move(source)), _binary()
86 : _context(
std::move(context)), _device(
std::move(device)), _is_binary(true), _name(
std::move(name)), _source(), _binary(
std::move(binary))
90 Program::operator cl::Program()
const 94 return cl::Program(_context, { _device }, { _binary });
98 return cl::Program(_context, _source,
false);
106 return program.build(build_options.c_str()) == CL_SUCCESS;
108 catch(
const cl::Error &e)
110 cl_int err = CL_SUCCESS;
111 const auto build_info = program.getBuildInfo<CL_PROGRAM_BUILD_LOG>(&err);
113 for(
auto &pair : build_info)
115 std::cerr << pair.second << std::endl;
124 cl::Program cl_program =
static_cast<cl::Program
>(*this);
125 build(cl_program, build_options);
135 : _name(
std::move(name)),
136 _kernel(
cl::
Kernel(program, _name.c_str()))
140 : _context(), _device(), _programs_map(), _built_programs_map(), _is_wbsm_supported()
145 : _context(), _device(), _programs_map(), _built_programs_map(), _is_wbsm_supported()
147 _context = std::move(context);
153 const std::string &kernel_path,
const StringSet &build_options_set,
bool is_binary)
const 155 const std::string
build_options = generate_build_options(build_options_set, kernel_path);
156 const std::string built_program_name = program_name +
"_" +
build_options;
157 auto built_program_it = _built_programs_map.find(built_program_name);
158 cl::Program cl_program;
160 if(_built_programs_map.end() != built_program_it)
163 cl_program = built_program_it->second;
167 Program program = load_program(program_name, program_source, is_binary);
170 cl_program = program.
build(build_options);
173 _built_programs_map.emplace(built_program_name, cl_program);
177 return Kernel(kernel_name, cl_program);
180 const Program &CLCompileContext::load_program(
const std::string &program_name,
const std::string &program_source,
bool is_binary)
const 182 const auto program_it = _programs_map.find(program_name);
184 if(program_it != _programs_map.end())
186 return program_it->second;
191 #ifdef EMBEDDED_KERNELS 193 program =
Program(_context, program_name, program_source);
197 program =
Program(_context, _device.
cl_device(), program_name, std::vector<unsigned char>(program_source.begin(), program_source.end()));
201 program =
Program(_context, program_name, program_source);
206 const auto new_program = _programs_map.emplace(program_name, std::move(program));
208 return new_program.first->second;
213 _context = std::move(context);
214 if(_context.get() !=
nullptr)
216 const auto cl_devices = _context.getInfo<CL_CONTEXT_DEVICES>();
218 if(!cl_devices.empty())
225 std::string CLCompileContext::generate_build_options(
const StringSet &build_options_set,
const std::string &kernel_path)
const 227 std::string concat_str;
229 #if defined(ARM_COMPUTE_DEBUG_ENABLED) 231 concat_str +=
" -DARM_COMPUTE_DEBUG_ENABLED";
232 #endif // defined(ARM_COMPUTE_DEBUG_ENABLED) 240 concat_str +=
" -DARM_COMPUTE_OPENCL_FP16_ENABLED=1 ";
243 if(_device.
supported(
"cl_arm_integer_dot_product_int8"))
245 concat_str +=
" -DARM_COMPUTE_OPENCL_DOT8_ENABLED=1 ";
248 if(_device.
supported(
"cl_arm_integer_dot_product_accumulate_int8"))
250 concat_str +=
" -DARM_COMPUTE_OPENCL_DOT8_ACC_ENABLED=1 ";
255 concat_str +=
" -cl-std=CL2.0 ";
257 else if(_device.
supported(
"cl_arm_non_uniform_work_group_size"))
259 concat_str +=
" -cl-arm-non-uniform-work-group-size ";
266 std::string
build_options = stringify_set(build_options_set, kernel_path) + concat_str;
276 std::string CLCompileContext::stringify_set(
const StringSet &s,
const std::string &kernel_path)
const 278 std::string concat_set;
279 #ifndef EMBEDDED_KERNELS 280 concat_set +=
"-I" + kernel_path +
" ";
286 for(
const auto &el : s)
288 concat_set +=
" " + el;
296 _built_programs_map.emplace(built_program_name, program);
301 _programs_map.clear();
302 _built_programs_map.clear();
307 return _built_programs_map;
322 _device = std::move(device);
329 cl::NDRange default_range;
337 default_range = cl::NDRange(128u, 1);
340 default_range = cl::NullRange;
343 return default_range;
348 return _device.
supported(
"cl_khr_int64_base_atomics");
353 return _is_wbsm_supported;
360 size_t err = kernel.getWorkGroupInfo(_device.
cl_device(), CL_KERNEL_WORK_GROUP_SIZE, &result);
361 ARM_COMPUTE_ERROR_ON_MSG(err != 0,
"clGetKernelWorkGroupInfo failed to return the maximum workgroup size for the kernel");
std::string device_version() const
Returns the device version as a string.
void set_device(cl::Device device)
Sets the CL device for which the programs are created.
const cl::Device & get_device() const
Gets the CL device for which the programs are created.
cl::Context & context()
Accessor for the associated CL context.
bool is_wbsm_supported() const
Kernel()
Default Constructor.
const StringSet & options() const
Gets the current options list set.
#define ARM_COMPUTE_ERROR(msg)
Print the given message then throw an std::runtime_error.
std::string to_string(T &&value)
Convert integer and float values to string.
bool get_wbsm_support_info(const cl::Device &device)
const std::vector< unsigned char > & binary() const
Returns program binary data.
CLBuildOptions()
Default constructor.
GPUTarget get_arch_from_target(GPUTarget target)
Helper function to get the GPU arch.
decltype(strategy::transforms) typedef type
std::string get_device_version() const
Return the device version.
Copyright (c) 2017-2021 Arm Limited.
void add_option(std::string option)
Adds option to the existing build option list.
std::string name() const
Returns program name.
void add_options(const StringSet &options)
Appends given build options to the current's objects options.
CLVersion version() const
Returns the device's CL version.
#define ARM_COMPUTE_UNUSED(...)
To avoid unused variables warnings.
cl_uint get_num_compute_units() const
Return the maximum number of compute units in the device.
const GPUTarget & target() const
Returns the GPU target of the cl device.
#define ARM_COMPUTE_ERROR_ON_MSG(cond, msg)
std::set< std::string > build_options
OpenCL device type class.
static bool build(const cl::Program &program, const std::string &build_options="")
Build the given CL program.
bool supported(const std::string &extension) const override
Check if extensions on a device are supported.
cl::NDRange default_ndrange() const
Return the default NDRange for the device.
void add_option_if(bool cond, std::string option)
Adds option if a given condition is true;.
GPUTarget get_target_from_device(const cl::Device &device)
Helper function to get the GPU target from CL device.
Kernel create_kernel(const std::string &kernel_name, const std::string &program_name, const std::string &program_source, const std::string &kernel_path, const StringSet &build_options_set, bool is_binary) const
Creates an OpenCL kernel.
void clear_programs_cache()
Clear the library's cache of binary programs.
Program()
Default constructor.
bool int64_base_atomics_supported() const
Returns true if int64_base_atomics extension is supported by the CL device.
bool fp16_supported() const
Returns true if FP16 is supported by the CL device.
size_t compute_units() const
Returns the number of compute units available.
const std::map< std::string, cl::Program > & get_built_programs() const
Access the cache of built OpenCL programs.
GPUTarget
Available GPU Targets.
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...
Wrapper to configure the Khronos OpenCL C++ header.
const cl::Device & cl_device() const
Returns the underlying cl device object.
void add_built_program(const std::string &built_program_name, const cl::Program &program) const
Add a new built program to the cache.
void set_context(cl::Context context)
Sets the CL context used to create programs.
CLCompileContext()
Constructor.
std::string name() const
Returns kernel name.
void add_options_if(bool cond, const StringSet &options)
Appends given build options to the current's objects options if a given condition is true...
void add_option_if_else(bool cond, std::string option_true, std::string option_false)
Adds first option if condition is true else the second one.