43 _build_opts.emplace(std::move(option));
61 _build_opts.insert(options.begin(), options.end());
78 : _context(), _device(), _is_binary(false), _name(), _source(), _binary()
83 : _context(
std::move(context)), _device(), _is_binary(false), _name(
std::move(name)), _source(
std::move(source)), _binary()
88 : _context(
std::move(context)), _device(
std::move(device)), _is_binary(true), _name(
std::move(name)), _source(), _binary(
std::move(binary))
92 Program::operator cl::Program()
const 96 return cl::Program(_context, { _device }, { _binary });
100 return cl::Program(_context, _source,
false);
108 return program.build(build_options.c_str()) == CL_SUCCESS;
110 catch(
const cl::Error &e)
112 cl_int err = CL_SUCCESS;
113 const auto build_info = program.getBuildInfo<CL_PROGRAM_BUILD_LOG>(&err);
115 for(
auto &pair : build_info)
117 std::cerr << pair.second << std::endl;
126 cl::Program cl_program =
static_cast<cl::Program
>(*this);
127 build(cl_program, build_options);
137 : _name(
std::move(name)),
138 _kernel(
cl::
Kernel(program, _name.c_str()))
142 : _context(), _device(), _programs_map(), _built_programs_map(), _is_wbsm_supported()
147 : _context(), _device(), _programs_map(), _built_programs_map(), _is_wbsm_supported()
149 _context = std::move(context);
155 const std::string &kernel_path,
const StringSet &build_options_set,
bool is_binary)
const 157 const std::string
build_options = generate_build_options(build_options_set, kernel_path);
158 const std::string built_program_name = program_name +
"_" +
build_options;
159 auto built_program_it = _built_programs_map.find(built_program_name);
160 cl::Program cl_program;
162 if(_built_programs_map.end() != built_program_it)
165 cl_program = built_program_it->second;
169 Program program = load_program(program_name, program_source, is_binary);
172 cl_program = program.
build(build_options);
175 _built_programs_map.emplace(built_program_name, cl_program);
179 return Kernel(kernel_name, cl_program);
182 const Program &CLCompileContext::load_program(
const std::string &program_name,
const std::string &program_source,
bool is_binary)
const 184 const auto program_it = _programs_map.find(program_name);
186 if(program_it != _programs_map.end())
188 return program_it->second;
193 #ifdef EMBEDDED_KERNELS 195 program =
Program(_context, program_name, program_source);
199 program =
Program(_context, _device.
cl_device(), program_name, std::vector<unsigned char>(program_source.begin(), program_source.end()));
203 program =
Program(_context, program_name, program_source);
208 const auto new_program = _programs_map.emplace(program_name, std::move(program));
210 return new_program.first->second;
215 _context = std::move(context);
216 if(_context.get() !=
nullptr)
218 const auto cl_devices = _context.getInfo<CL_CONTEXT_DEVICES>();
220 if(!cl_devices.empty())
227 std::string CLCompileContext::generate_build_options(
const StringSet &build_options_set,
const std::string &kernel_path)
const 229 std::string concat_str;
231 #if defined(ARM_COMPUTE_DEBUG_ENABLED) 233 concat_str +=
" -DARM_COMPUTE_DEBUG_ENABLED";
234 #endif // defined(ARM_COMPUTE_DEBUG_ENABLED) 242 concat_str +=
" -DARM_COMPUTE_OPENCL_FP16_ENABLED=1 ";
245 if(_device.
supported(
"cl_arm_integer_dot_product_int8"))
247 concat_str +=
" -DARM_COMPUTE_OPENCL_DOT8_ENABLED=1 ";
250 if(_device.
supported(
"cl_arm_integer_dot_product_accumulate_int8"))
252 concat_str +=
" -DARM_COMPUTE_OPENCL_DOT8_ACC_ENABLED=1 ";
257 concat_str +=
" -cl-std=CL2.0 ";
259 else if(_device.
supported(
"cl_arm_non_uniform_work_group_size"))
261 concat_str +=
" -cl-arm-non-uniform-work-group-size ";
271 const std::regex ddk_regex(
"r([0-9]*)p[0-9]");
272 std::smatch ddk_match;
274 if(std::regex_search(device_vers, ddk_match, ddk_regex) &&
std::stoi(ddk_match[1]) >= 11)
276 concat_str +=
" -DUNROLL_WITH_PRAGMA ";
280 std::string
build_options = stringify_set(build_options_set, kernel_path) + concat_str;
290 std::string CLCompileContext::stringify_set(
const StringSet &s,
const std::string &kernel_path)
const 292 std::string concat_set;
293 #ifndef EMBEDDED_KERNELS 294 concat_set +=
"-I" + kernel_path +
" ";
300 for(
const auto &el : s)
302 concat_set +=
" " + el;
310 _built_programs_map.emplace(built_program_name, program);
315 _programs_map.clear();
316 _built_programs_map.clear();
321 return _built_programs_map;
336 _device = std::move(device);
343 cl::NDRange default_range;
351 default_range = cl::NDRange(128u, 1);
354 default_range = cl::NullRange;
357 return default_range;
362 return _device.
supported(
"cl_khr_int64_base_atomics");
367 return _is_wbsm_supported;
374 size_t err = kernel.getWorkGroupInfo(_device.
cl_device(), CL_KERNEL_WORK_GROUP_SIZE, &result);
375 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::set< std::string > build_options
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)
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.
int stoi(const std::string &str, std::size_t *pos=0, NumericBase base=NumericBase::BASE_10)
Convert string values to integer.
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.