From 11d4918b2321d1e590124f44dd68e6cda223dbdc Mon Sep 17 00:00:00 2001 From: Michalis Spyrou Date: Thu, 26 Mar 2020 10:31:32 +0000 Subject: COMPMID-3279: Create CLCompiler interface Change-Id: Ic9dd5288d72a690651aa03d474f2bfd6e1ebe8b2 Signed-off-by: Michalis Spyrou Reviewed-on: https://review.mlplatform.org/c/ml/ComputeLibrary/+/2957 Tested-by: Arm Jenkins Comments-Addressed: Arm Jenkins Reviewed-by: Georgios Pinitas Reviewed-by: Gian Marco Iodice --- src/core/CL/CLKernelLibrary.cpp | 329 ++++++---------------------------------- 1 file changed, 47 insertions(+), 282 deletions(-) (limited to 'src/core/CL/CLKernelLibrary.cpp') diff --git a/src/core/CL/CLKernelLibrary.cpp b/src/core/CL/CLKernelLibrary.cpp index c6c88569ce..7437f1bf22 100644 --- a/src/core/CL/CLKernelLibrary.cpp +++ b/src/core/CL/CLKernelLibrary.cpp @@ -35,113 +35,6 @@ #include using namespace arm_compute; - -CLBuildOptions::CLBuildOptions() - : _build_opts() -{ -} - -void CLBuildOptions::add_option(std::string option) -{ - _build_opts.emplace(std::move(option)); -} - -void CLBuildOptions::add_option_if(bool cond, std::string option) -{ - if(cond) - { - add_option(std::move(option)); - } -} - -void CLBuildOptions::add_option_if_else(bool cond, std::string option_true, std::string option_false) -{ - (cond) ? add_option(std::move(option_true)) : add_option(std::move(option_false)); -} - -void CLBuildOptions::add_options(const StringSet &options) -{ - _build_opts.insert(options.begin(), options.end()); -} - -void CLBuildOptions::add_options_if(bool cond, const StringSet &options) -{ - if(cond) - { - add_options(options); - } -} - -const CLBuildOptions::StringSet &CLBuildOptions::options() const -{ - return _build_opts; -} - -Program::Program() - : _context(), _device(), _is_binary(false), _name(), _source(), _binary() -{ -} - -Program::Program(cl::Context context, std::string name, std::string source) - : _context(std::move(context)), _device(), _is_binary(false), _name(std::move(name)), _source(std::move(source)), _binary() -{ -} - -Program::Program(cl::Context context, cl::Device device, std::string name, std::vector binary) - : _context(std::move(context)), _device(std::move(device)), _is_binary(true), _name(std::move(name)), _source(), _binary(std::move(binary)) -{ -} - -Program::operator cl::Program() const -{ - if(_is_binary) - { - return cl::Program(_context, { _device }, { _binary }); - } - else - { - return cl::Program(_context, _source, false); - } -} - -bool Program::build(const cl::Program &program, const std::string &build_options) -{ - try - { - return program.build(build_options.c_str()) == CL_SUCCESS; - } - catch(const cl::Error &e) - { - cl_int err = CL_SUCCESS; - const auto build_info = program.getBuildInfo(&err); - - for(auto &pair : build_info) - { - std::cerr << pair.second << std::endl; - } - - return false; - } -} - -cl::Program Program::build(const std::string &build_options) const -{ - cl::Program cl_program = static_cast(*this); - build(cl_program, build_options); - return cl_program; -} - -Kernel::Kernel() - : _name(), _kernel() -{ -} - -Kernel::Kernel(std::string name, const cl::Program &program) - : _name(std::move(name)), - _kernel(cl::Kernel(program, _name.c_str())) -{ -} - const std::map CLKernelLibrary::_kernel_program_map = { { "absdiff", "absdiff.cl" }, @@ -1066,7 +959,7 @@ const std::map CLKernelLibrary::_program_source_map = }; CLKernelLibrary::CLKernelLibrary() - : _context(), _device(), _kernel_path("."), _programs_map(), _built_programs_map() + : _compile_context(), _kernel_path() { opencl_is_available(); // Make sure the OpenCL symbols are initialised *before* the CLKernelLibrary is built } @@ -1077,7 +970,15 @@ CLKernelLibrary &CLKernelLibrary::get() return _kernel_library; } -Kernel CLKernelLibrary::create_kernel(const std::string &kernel_name, const StringSet &build_options_set) const +Kernel CLKernelLibrary::create_kernel(const std::string &kernel_name, const std::set &build_options_set) const +{ + const std::string program_name = get_program_name(kernel_name); + auto program = get_program(program_name); + + return _compile_context.create_kernel(kernel_name, program_name, program.first, _kernel_path, build_options_set, program.second); +} + +std::string CLKernelLibrary::get_program_name(const std::string &kernel_name) const { // Find which program contains the kernel auto kernel_program_it = _kernel_program_map.find(kernel_name); @@ -1086,99 +987,41 @@ Kernel CLKernelLibrary::create_kernel(const std::string &kernel_name, const Stri { ARM_COMPUTE_ERROR_VAR("Kernel %s not found in the CLKernelLibrary", kernel_name.c_str()); } - std::string concat_str; - -#if defined(ARM_COMPUTE_DEBUG_ENABLED) - // Enable debug properties in CL kernels - concat_str += " -DARM_COMPUTE_DEBUG_ENABLED"; -#endif // defined(ARM_COMPUTE_DEBUG_ENABLED) - GPUTarget gpu_arch = get_arch_from_target(get_target_from_device(_device)); - concat_str += " -DGPU_ARCH=" + support::cpp11::to_string( - static_cast::type>(gpu_arch)); - if(fp16_supported()) - { - concat_str += " -DARM_COMPUTE_OPENCL_FP16_ENABLED=1 "; - } + const std::string program_name = kernel_program_it->second; - if(dot8_supported(_device)) - { - concat_str += " -DARM_COMPUTE_OPENCL_DOT8_ENABLED=1 "; - } - - if(dot8_acc_supported(_device)) - { - concat_str += " -DARM_COMPUTE_OPENCL_DOT8_ACC_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); + return program_name; } void CLKernelLibrary::init(std::string kernel_path, cl::Context context, cl::Device device) { - _kernel_path = std::move(kernel_path); - _context = std::move(context); - _device = std::move(device); + _compile_context = CLCompileContext(context, device); + _kernel_path = kernel_path; } void CLKernelLibrary::set_kernel_path(const std::string &kernel_path) { - _kernel_path = kernel_path; + _kernel_path = std::move(kernel_path); } cl::Context &CLKernelLibrary::context() { - return _context; + return _compile_context.context(); } -cl::Device &CLKernelLibrary::get_device() +const cl::Device &CLKernelLibrary::get_device() { - return _device; + return _compile_context.get_device(); } void CLKernelLibrary::set_device(cl::Device device) { - _device = std::move(device); + _compile_context.set_device(device); +} + +void CLKernelLibrary::set_context(cl::Context context) +{ + _compile_context.set_context(context); } std::string CLKernelLibrary::get_kernel_path() @@ -1188,164 +1031,86 @@ std::string CLKernelLibrary::get_kernel_path() void CLKernelLibrary::clear_programs_cache() { - _programs_map.clear(); - _built_programs_map.clear(); + _compile_context.clear_programs_cache(); } const std::map &CLKernelLibrary::get_built_programs() const { - return _built_programs_map; + return _compile_context.get_built_programs(); } void CLKernelLibrary::add_built_program(const std::string &built_program_name, const cl::Program &program) { - _built_programs_map.emplace(built_program_name, program); + _compile_context.add_built_program(built_program_name, program); } bool CLKernelLibrary::fp16_supported() const { - return ::fp16_supported(_device); + return _compile_context.fp16_supported(); } bool CLKernelLibrary::int64_base_atomics_supported() const { - return device_supports_extension(_device, "cl_khr_int64_base_atomics"); + return _compile_context.int64_base_atomics_supported(); } -const Program &CLKernelLibrary::load_program(const std::string &program_name) const +std::pair CLKernelLibrary::get_program(const std::string &program_name) const { - const auto program_it = _programs_map.find(program_name); - - if(program_it != _programs_map.end()) - { - return program_it->second; - } - - Program program; - #ifdef EMBEDDED_KERNELS const auto program_source_it = _program_source_map.find(program_name); - if(_program_source_map.end() == program_source_it) + if(program_source_it == _program_source_map.end()) { ARM_COMPUTE_ERROR_VAR("Embedded program for %s does not exist.", program_name.c_str()); } - program = Program(_context, program_name, program_source_it->second); + return std::make_pair(program_source_it->second, false); #else /* EMBEDDED_KERNELS */ // Check for binary std::string source_name = _kernel_path + program_name; std::string binary_name = source_name + "bin"; + std::string program_source{}; + bool is_binary = false; if(std::ifstream(binary_name).is_open()) { - const std::string program_binary = read_file(binary_name, true); - program = Program(_context, _device, program_name, std::vector(program_binary.begin(), program_binary.end())); + program_source = read_file(binary_name, true); + is_binary = true; } else if(std::ifstream(source_name).is_open()) { - program = Program(_context, program_name, read_file(source_name, false)); + program_source = read_file(source_name, false); } else { ARM_COMPUTE_ERROR_VAR("Kernel file %s does not exist.", source_name.c_str()); } -#endif /* EMBEDDED_KERNELS */ - // Insert program to program map - const auto new_program = _programs_map.emplace(program_name, std::move(program)); - - return new_program.first->second; -} - -void CLKernelLibrary::set_context(cl::Context context) -{ - _context = std::move(context); - if(_context.get() == nullptr) - { - _device = cl::Device(); - } - else - { - const auto cl_devices = _context.getInfo(); - - if(cl_devices.empty()) - { - _device = cl::Device(); - } - else - { - _device = cl_devices[0]; - } - } -} - -std::string CLKernelLibrary::stringify_set(const StringSet &s) const -{ - std::string concat_set; - -#ifndef EMBEDDED_KERNELS - concat_set += "-I" + _kernel_path + " "; + return std::make_pair(program_source, is_binary); #endif /* EMBEDDED_KERNELS */ - - // Concatenate set - for(const auto &el : s) - { - concat_set += " " + el; - } - - return concat_set; -} - -std::string CLKernelLibrary::get_program_source(const std::string &program_name) -{ - 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.c_str()); - } - - return program_source_it->second; } size_t CLKernelLibrary::max_local_workgroup_size(const cl::Kernel &kernel) const { - 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; + return _compile_context.max_local_workgroup_size(kernel); } cl::NDRange CLKernelLibrary::default_ndrange() const { - 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; + return _compile_context.default_ndrange(); } std::string CLKernelLibrary::get_device_version() { - return _device.getInfo(); + return _compile_context.get_device_version(); } cl_uint CLKernelLibrary::get_num_compute_units() { - return _device.getInfo(); + return _compile_context.get_num_compute_units(); +} + +CLCompileContext &CLKernelLibrary::get_compile_context() +{ + return _compile_context; } -- cgit v1.2.1