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/CLCompileContext.cpp | 369 ++++++++++++++++++++++++++++++++++ src/core/CL/CLHelpers.cpp | 8 + src/core/CL/CLKernelLibrary.cpp | 329 +++++------------------------- src/core/CL/kernels/CLFloorKernel.cpp | 15 +- 4 files changed, 434 insertions(+), 287 deletions(-) create mode 100644 src/core/CL/CLCompileContext.cpp (limited to 'src/core/CL') diff --git a/src/core/CL/CLCompileContext.cpp b/src/core/CL/CLCompileContext.cpp new file mode 100644 index 0000000000..48cc64c387 --- /dev/null +++ b/src/core/CL/CLCompileContext.cpp @@ -0,0 +1,369 @@ +/* + * Copyright (c) 2020 ARM Limited. + * + * SPDX-License-Identifier: MIT + * + * Permission is hereby granted, free of charge, to any person obtaining a copy + * of this software and associated documentation files (the "Software"), to + * deal in the Software without restriction, including without limitation the + * rights to use, copy, modify, merge, publish, distribute, sublicense, and/or + * sell copies of the Software, and to permit persons to whom the Software is + * furnished to do so, subject to the following conditions: + * + * The above copyright notice and this permission notice shall be included in all + * copies or substantial portions of the Software. + * + * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR + * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, + * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE + * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER + * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, + * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE + * SOFTWARE. + */ +#include "arm_compute/core/CL/CLCompileContext.h" +#include "arm_compute/core/CL/OpenCL.h" + +#include "arm_compute/core/CL/CLHelpers.h" +#include "arm_compute/core/Error.h" +#include "arm_compute/core/Utils.h" +#include "support/StringSupport.h" + +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())) +{ +} +CLCompileContext::CLCompileContext() + : _context(), _device(), _programs_map(), _built_programs_map() +{ +} + +CLCompileContext::CLCompileContext(cl::Context context, const cl::Device &device) + : _context(), _device(), _programs_map(), _built_programs_map() +{ + _context = std::move(context); + _device = CLDevice(device); +} + +Kernel CLCompileContext::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 +{ + const std::string build_options = generate_build_options(build_options_set, kernel_path); + 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 + { + Program program = load_program(program_name, program_source, is_binary); + + // Build program + cl_program = program.build(build_options); + + // Add built program to internal map + _built_programs_map.emplace(program_name, cl_program); + } + + // Create and return kernel + return Kernel(kernel_name, cl_program); +} + +const Program &CLCompileContext::load_program(const std::string &program_name, const std::string &program_source, bool is_binary) 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 + ARM_COMPUTE_UNUSED(is_binary); + program = Program(_context, program_name, program_source); +#else /* EMBEDDED_KERNELS */ + if(is_binary) + { + program = Program(_context, _device.cl_device(), program_name, std::vector(program_source.begin(), program_source.end())); + } + else + { + program = Program(_context, program_name, program_source); + } +#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 CLCompileContext::set_context(cl::Context context) +{ + _context = std::move(context); + if(_context.get() != nullptr) + { + const auto cl_devices = _context.getInfo(); + + if(!cl_devices.empty()) + { + _device = CLDevice(cl_devices[0]); + } + } +} + +std::string CLCompileContext::generate_build_options(const StringSet &build_options_set, const std::string &kernel_path) const +{ + 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(_device.target()); + concat_str += " -DGPU_ARCH=" + support::cpp11::to_string( + static_cast::type>(gpu_arch)); + + if(_device.supported("cl_khr_fp16")) + { + concat_str += " -DARM_COMPUTE_OPENCL_FP16_ENABLED=1 "; + } + + if(_device.supported("cl_arm_integer_dot_product_int8")) + { + concat_str += " -DARM_COMPUTE_OPENCL_DOT8_ENABLED=1 "; + } + + if(_device.supported("cl_arm_integer_dot_product_accumulate_int8")) + { + concat_str += " -DARM_COMPUTE_OPENCL_DOT8_ACC_ENABLED=1 "; + } + + if(_device.version() == CLVersion::CL20) + { + concat_str += " -cl-std=CL2.0 "; + } + else if(_device.supported("cl_arm_non_uniform_work_group_size")) + { + concat_str += " -cl-arm-non-uniform-work-group-size "; + } + else + { + ARM_COMPUTE_ERROR("Non uniform workgroup size is not supported!!"); + } + + std::string build_options = stringify_set(build_options_set, kernel_path) + concat_str; + + return build_options; +} + +bool CLCompileContext::fp16_supported() const +{ + return _device.supported("cl_khr_fp16"); +} + +std::string CLCompileContext::stringify_set(const StringSet &s, const std::string &kernel_path) const +{ + std::string concat_set; +#ifndef EMBEDDED_KERNELS + concat_set += "-I" + kernel_path + " "; +#else /* EMBEDDED_KERNELS */ + ARM_COMPUTE_UNUSED(kernel_path); +#endif /* EMBEDDED_KERNELS */ + + // Concatenate set + for(const auto &el : s) + { + concat_set += " " + el; + } + + return concat_set; +} + +void CLCompileContext::add_built_program(const std::string &built_program_name, const cl::Program &program) const +{ + _built_programs_map.emplace(built_program_name, program); +} + +void CLCompileContext::clear_programs_cache() +{ + _programs_map.clear(); + _built_programs_map.clear(); +} + +const std::map &CLCompileContext::get_built_programs() const +{ + return _built_programs_map; +} + +cl::Context &CLCompileContext::context() +{ + return _context; +} + +const cl::Device &CLCompileContext::get_device() const +{ + return _device.cl_device(); +} + +void CLCompileContext::set_device(cl::Device device) +{ + _device = std::move(device); +} + +cl::NDRange CLCompileContext::default_ndrange() const +{ + GPUTarget _target = get_target_from_device(_device.cl_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; +} + +bool CLCompileContext::int64_base_atomics_supported() const +{ + return _device.supported("cl_khr_int64_base_atomics"); +} + +size_t CLCompileContext::max_local_workgroup_size(const cl::Kernel &kernel) const +{ + size_t result; + + size_t err = kernel.getWorkGroupInfo(_device.cl_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; +} + +std::string CLCompileContext::get_device_version() const +{ + return _device.device_version(); +} + +cl_uint CLCompileContext::get_num_compute_units() const +{ + return _device.compute_units(); +} +} // namespace arm_compute diff --git a/src/core/CL/CLHelpers.cpp b/src/core/CL/CLHelpers.cpp index 84de380cc9..7d1b0ea6c7 100644 --- a/src/core/CL/CLHelpers.cpp +++ b/src/core/CL/CLHelpers.cpp @@ -384,6 +384,14 @@ cl::Kernel create_opencl_kernel(CLCoreRuntimeContext *ctx, const std::string &ke } } +cl::Kernel create_kernel(CLCompileContext &ctx, const std::string &kernel_name, const std::set &build_opts) +{ + const std::string program_name = CLKernelLibrary::get().get_program_name(kernel_name); + std::pair kernel_src = CLKernelLibrary::get().get_program(program_name); + const std::string kernel_path = CLKernelLibrary::get().get_kernel_path(); + return static_cast(ctx.create_kernel(kernel_name, program_name, kernel_src.first, kernel_path, build_opts, kernel_src.second)); +} + cl::NDRange create_lws_hint_parallel_implementations(unsigned int input_dimension, unsigned int vector_size) { const unsigned int width_leftover = input_dimension % vector_size; 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; } diff --git a/src/core/CL/kernels/CLFloorKernel.cpp b/src/core/CL/kernels/CLFloorKernel.cpp index 8f0043f08a..abfed8d18e 100644 --- a/src/core/CL/kernels/CLFloorKernel.cpp +++ b/src/core/CL/kernels/CLFloorKernel.cpp @@ -77,7 +77,7 @@ CLFloorKernel::CLFloorKernel() { } -void CLFloorKernel::configure(const ICLTensor *input, ICLTensor *output) +void CLFloorKernel::configure(CLCompileContext &compile_context, const ICLTensor *input, ICLTensor *output) { ARM_COMPUTE_ERROR_ON_NULLPTR(input, output); @@ -90,13 +90,13 @@ void CLFloorKernel::configure(const ICLTensor *input, ICLTensor *output) _input = input; _output = output; - const unsigned int num_elems_processed_per_iteration = 16 / input->info()->element_size(); - - // Create kernel + const unsigned int num_elems_processed_per_iteration = 16 / input->info()->element_size(); std::set build_opts; build_opts.emplace(("-DDATA_TYPE=" + get_cl_type_from_data_type(input->info()->data_type()))); build_opts.emplace(("-DVEC_SIZE=" + support::cpp11::to_string(num_elems_processed_per_iteration))); - _kernel = static_cast(CLKernelLibrary::get().create_kernel("floor_layer", build_opts)); + + // Create kernel + _kernel = create_kernel(compile_context, "floor_layer", build_opts); // Configure kernel window auto win_config = validate_and_configure_window(input->info(), output->info()); @@ -104,6 +104,11 @@ void CLFloorKernel::configure(const ICLTensor *input, ICLTensor *output) ICLKernel::configure_internal(win_config.second); } +void CLFloorKernel::configure(const ICLTensor *input, ICLTensor *output) +{ + configure(CLKernelLibrary::get().get_compile_context(), input, output); +} + Status CLFloorKernel::validate(const ITensorInfo *input, const ITensorInfo *output) { ARM_COMPUTE_RETURN_ON_ERROR(validate_arguments(input, output)); -- cgit v1.2.1