43 _build_opts.emplace(std::move(option));
61 _build_opts.insert(options.begin(), options.end());
79 return _build_opts == other._build_opts;
83 : _context(), _device(), _is_binary(false), _name(), _source(), _binary()
88 : _context(
std::move(context)), _device(), _is_binary(false), _name(
std::move(name)), _source(
std::move(source)), _binary()
93 : _context(
std::move(context)), _device(
std::move(device)), _is_binary(true), _name(
std::move(name)), _source(), _binary(
std::move(binary))
97 Program::operator cl::Program()
const 101 return cl::Program(_context, { _device }, { _binary });
105 return cl::Program(_context, _source,
false);
113 return program.build(build_options.c_str()) == CL_SUCCESS;
115 catch(
const cl::Error &e)
117 cl_int err = CL_SUCCESS;
118 const auto build_info = program.getBuildInfo<CL_PROGRAM_BUILD_LOG>(&err);
120 for(
auto &pair : build_info)
122 std::cerr << pair.second << std::endl;
131 cl::Program cl_program =
static_cast<cl::Program
>(*this);
132 build(cl_program, build_options);
142 : _name(
std::move(name)),
143 _kernel(
cl::
Kernel(program, _name.c_str()))
147 : _context(), _device(), _programs_map(), _built_programs_map(), _is_wbsm_supported()
152 : _context(), _device(), _programs_map(), _built_programs_map(), _is_wbsm_supported()
154 _context = std::move(context);
160 const std::string &kernel_path,
const StringSet &build_options_set,
bool is_binary)
const 162 const std::string
build_options = generate_build_options(build_options_set, kernel_path);
163 const std::string built_program_name = program_name +
"_" +
build_options;
164 auto built_program_it = _built_programs_map.find(built_program_name);
165 cl::Program cl_program;
167 if(_built_programs_map.end() != built_program_it)
170 cl_program = built_program_it->second;
174 Program program = load_program(program_name, program_source, is_binary);
177 cl_program = program.
build(build_options);
180 _built_programs_map.emplace(built_program_name, cl_program);
184 return Kernel(kernel_name, cl_program);
187 const Program &CLCompileContext::load_program(
const std::string &program_name,
const std::string &program_source,
bool is_binary)
const 189 const auto program_it = _programs_map.find(program_name);
191 if(program_it != _programs_map.end())
193 return program_it->second;
198 #ifdef EMBEDDED_KERNELS 200 program =
Program(_context, program_name, program_source);
204 program =
Program(_context, _device.
cl_device(), program_name, std::vector<unsigned char>(program_source.begin(), program_source.end()));
208 program =
Program(_context, program_name, program_source);
213 const auto new_program = _programs_map.emplace(program_name, std::move(program));
215 return new_program.first->second;
220 _context = std::move(context);
221 if(_context.get() !=
nullptr)
223 const auto cl_devices = _context.getInfo<CL_CONTEXT_DEVICES>();
225 if(!cl_devices.empty())
232 std::string CLCompileContext::generate_build_options(
const StringSet &build_options_set,
const std::string &kernel_path)
const 234 std::string concat_str;
236 #if defined(ARM_COMPUTE_DEBUG_ENABLED) 238 concat_str +=
" -DARM_COMPUTE_DEBUG_ENABLED";
239 #endif // defined(ARM_COMPUTE_DEBUG_ENABLED) 247 concat_str +=
" -DARM_COMPUTE_OPENCL_FP16_ENABLED=1 ";
250 if(_device.
supported(
"cl_arm_integer_dot_product_int8"))
252 concat_str +=
" -DARM_COMPUTE_OPENCL_DOT8_ENABLED=1 ";
255 if(_device.
supported(
"cl_arm_integer_dot_product_accumulate_int8"))
257 concat_str +=
" -DARM_COMPUTE_OPENCL_DOT8_ACC_ENABLED=1 ";
262 concat_str +=
" -cl-std=CL2.0 ";
264 else if(_device.
supported(
"cl_arm_non_uniform_work_group_size"))
266 concat_str +=
" -cl-arm-non-uniform-work-group-size ";
276 const std::regex ddk_regex(
"r([0-9]*)p[0-9]");
277 std::smatch ddk_match;
279 if(std::regex_search(device_vers, ddk_match, ddk_regex) &&
std::stoi(ddk_match[1]) >= 11)
281 concat_str +=
" -DUNROLL_WITH_PRAGMA ";
285 std::string
build_options = stringify_set(build_options_set, kernel_path) + concat_str;
295 std::string CLCompileContext::stringify_set(
const StringSet &s,
const std::string &kernel_path)
const 297 std::string concat_set;
298 #ifndef EMBEDDED_KERNELS 299 concat_set +=
"-I" + kernel_path +
" ";
305 for(
const auto &el : s)
307 concat_set +=
" " + el;
315 _built_programs_map.emplace(built_program_name, program);
320 _programs_map.clear();
321 _built_programs_map.clear();
326 return _built_programs_map;
341 _device = std::move(device);
348 cl::NDRange default_range;
356 default_range = cl::NDRange(128u, 1);
359 default_range = cl::NullRange;
362 return default_range;
367 return _device.
supported(
"cl_khr_int64_base_atomics");
372 return _is_wbsm_supported;
379 size_t err = kernel.getWorkGroupInfo(_device.
cl_device(), CL_KERNEL_WORK_GROUP_SIZE, &result);
380 ARM_COMPUTE_ERROR_ON_MSG(err != 0,
"clGetKernelWorkGroupInfo failed to return the maximum workgroup size for the kernel");
bool operator==(const CLBuildOptions &other) const
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-2022 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.