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 --- Android.bp | 1 + arm_compute/core/CL/CLCompileContext.h | 324 ++++++++++++++++++++++++ arm_compute/core/CL/CLDevice.h | 152 ++++++++++++ arm_compute/core/CL/CLHelpers.h | 12 + arm_compute/core/CL/CLKernelLibrary.h | 213 +++------------- arm_compute/core/CL/CLTypes.h | 17 +- arm_compute/core/CL/kernels/CLFloorKernel.h | 11 +- arm_compute/core/IDevice.h | 60 +++++ arm_compute/runtime/CL/CLRuntimeContext.h | 1 - arm_compute/runtime/CL/functions/CLFloor.h | 9 +- 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 +- src/runtime/CL/CLRuntimeContext.cpp | 8 +- src/runtime/CL/functions/CLFloor.cpp | 7 +- tests/validation/CL/UNIT/CompileContext.cpp | 74 ++++++ 17 files changed, 1123 insertions(+), 487 deletions(-) create mode 100644 arm_compute/core/CL/CLCompileContext.h create mode 100644 arm_compute/core/CL/CLDevice.h create mode 100644 arm_compute/core/IDevice.h create mode 100644 src/core/CL/CLCompileContext.cpp create mode 100644 tests/validation/CL/UNIT/CompileContext.cpp diff --git a/Android.bp b/Android.bp index 6e34cb5cca..f31dabbe07 100644 --- a/Android.bp +++ b/Android.bp @@ -51,6 +51,7 @@ cc_library_static { "src/core/AccessWindowAutoPadding.cpp", "src/core/AccessWindowStatic.cpp", "src/core/AccessWindowTranspose.cpp", + "src/core/CL/CLCompileContext.cpp", "src/core/CL/CLCoreRuntimeContext.cpp", "src/core/CL/CLHelpers.cpp", "src/core/CL/CLKernelLibrary.cpp", diff --git a/arm_compute/core/CL/CLCompileContext.h b/arm_compute/core/CL/CLCompileContext.h new file mode 100644 index 0000000000..2b6d8cd2cb --- /dev/null +++ b/arm_compute/core/CL/CLCompileContext.h @@ -0,0 +1,324 @@ +/* + * 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. + */ +#ifndef ARM_COMPUTE_CLCOMPILECONTEXT_H +#define ARM_COMPUTE_CLCOMPILECONTEXT_H + +#include "arm_compute/core/CL/CLDevice.h" +#include "arm_compute/core/CL/OpenCL.h" + +#include +#include +#include +#include + +namespace arm_compute +{ +/** Build options */ +class CLBuildOptions final +{ + using StringSet = std::set; + +public: + /** Default constructor. */ + CLBuildOptions(); + /** Adds option to the existing build option list + * + * @param[in] option Option to add + */ + void add_option(std::string option); + /** Adds option if a given condition is true; + * + * @param[in] cond Condition to check + * @param[in] option Option to add if condition is true + */ + void add_option_if(bool cond, std::string option); + /** Adds first option if condition is true else the second one + * + * @param[in] cond Condition to check + * @param[in] option_true Option to add if condition is true + * @param[in] option_false Option to add if condition is false + */ + void add_option_if_else(bool cond, std::string option_true, std::string option_false); + /** Appends given build options to the current's objects options. + * + * @param[in] options Build options to append + */ + void add_options(const StringSet &options); + /** Appends given build options to the current's objects options if a given condition is true. + * + * @param[in] cond Condition to check + * @param[in] options Option to add if condition is true + */ + void add_options_if(bool cond, const StringSet &options); + /** Gets the current options list set + * + * @return Build options set + */ + const StringSet &options() const; + +private: + StringSet _build_opts; /**< Build options set */ +}; + +/** Program class */ +class Program final +{ +public: + /** Default constructor. */ + Program(); + /** Construct program from source file. + * + * @param[in] context CL context used to create the program. + * @param[in] name Program name. + * @param[in] source Program source. + */ + Program(cl::Context context, std::string name, std::string source); + /** Construct program from binary file. + * + * @param[in] context CL context used to create the program. + * @param[in] device CL device for which the programs are created. + * @param[in] name Program name. + * @param[in] binary Program binary. + */ + Program(cl::Context context, cl::Device device, std::string name, std::vector binary); + /** Default Copy Constructor. */ + Program(const Program &) = default; + /** Default Move Constructor. */ + Program(Program &&) = default; + /** Default copy assignment operator */ + Program &operator=(const Program &) = default; + /** Default move assignment operator */ + Program &operator=(Program &&) = default; + /** Returns program name. + * + * @return Program's name. + */ + std::string name() const + { + return _name; + } + /** User-defined conversion to the underlying CL program. + * + * @return The CL program object. + */ + explicit operator cl::Program() const; + /** Build the given CL program. + * + * @param[in] program The CL program to build. + * @param[in] build_options Options to build the CL program. + * + * @return True if the CL program builds successfully. + */ + static bool build(const cl::Program &program, const std::string &build_options = ""); + /** Build the underlying CL program. + * + * @param[in] build_options Options used to build the CL program. + * + * @return A reference to itself. + */ + cl::Program build(const std::string &build_options = "") const; + +private: + cl::Context _context; /**< Underlying CL context. */ + cl::Device _device; /**< CL device for which the programs are created. */ + bool _is_binary; /**< Create program from binary? */ + std::string _name; /**< Program name. */ + std::string _source; /**< Source code for the program. */ + std::vector _binary; /**< Binary from which to create the program. */ +}; + +/** Kernel class */ +class Kernel final +{ +public: + /** Default Constructor. */ + Kernel(); + /** Default Copy Constructor. */ + Kernel(const Kernel &) = default; + /** Default Move Constructor. */ + Kernel(Kernel &&) = default; + /** Default copy assignment operator */ + Kernel &operator=(const Kernel &) = default; + /** Default move assignment operator */ + Kernel &operator=(Kernel &&) = default; + /** Constructor. + * + * @param[in] name Kernel name. + * @param[in] program Built program. + */ + Kernel(std::string name, const cl::Program &program); + /** Returns kernel name. + * + * @return Kernel's name. + */ + std::string name() const + { + return _name; + } + /** Returns OpenCL kernel. + * + * @return OpenCL Kernel. + */ + explicit operator cl::Kernel() const + { + return _kernel; + } + +private: + std::string _name; /**< Kernel name */ + cl::Kernel _kernel; /**< OpenCL Kernel */ +}; + +/** CLCompileContext class */ +class CLCompileContext final +{ + using StringSet = std::set; + +public: + /** Constructor */ + CLCompileContext(); + /** Constructor + * + * @param[in] context A CL context. + * @param[in] device A CL device. + * */ + CLCompileContext(cl::Context context, const cl::Device &device); + + /** Accessor for the associated CL context. + * + * @return A CL context. + */ + cl::Context &context(); + + /** Sets the CL context used to create programs. + * + * @note Setting the context also resets the device to the + * first one available in the new context. + * + * @param[in] context A CL context. + */ + void set_context(cl::Context context); + + /** Gets the CL device for which the programs are created. */ + const cl::Device &get_device() const; + + /** Sets the CL device for which the programs are created. + * + * @param[in] device A CL device. + */ + void set_device(cl::Device device); + + /** Creates an OpenCL kernel. + * + * @param[in] kernel_name Kernel name. + * @param[in] program_name Program name. + * @param[in] program_source Program source. + * @param[in] kernel_path CL kernel path. + * @param[in] build_options_set Kernel build options as a set. + * @param[in] is_binary Flag to indicate if the program source is binary. + * + * @return The created kernel. + */ + 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; + + /** Clear the library's cache of binary programs + */ + void clear_programs_cache(); + + /** Access the cache of built OpenCL programs */ + const std::map &get_built_programs() const; + + /** Add a new built program to the cache + * + * @param[in] built_program_name Name of the program + * @param[in] program Built program to add to the cache + */ + void add_built_program(const std::string &built_program_name, const cl::Program &program) const; + + /** Returns true if FP16 is supported by the CL device + * + * @return true if the CL device supports FP16 + */ + bool fp16_supported() const; + + /** Return the maximum number of compute units in the device + * + * @return The content of CL_DEVICE_MAX_COMPUTE_UNITS + */ + cl_uint get_num_compute_units() const; + /** Find the maximum number of local work items in a workgroup can be supported for the kernel. + * + */ + size_t max_local_workgroup_size(const cl::Kernel &kernel) const; + /** Return the default NDRange for the device. + * + */ + cl::NDRange default_ndrange() const; + /** Return the device version + * + * @return The content of CL_DEVICE_VERSION + */ + std::string get_device_version() const; + + /** Returns true if int64_base_atomics extension is supported by the CL device + * + * @return true if the CL device supports int64_base_atomics extension + */ + bool int64_base_atomics_supported() const; + +private: + /** Load program and its dependencies. + * + * @param[in] program_name Name of the program to load. + * @param[in] program_source Source of the program. + * @param[in] is_binary Flag to indicate if the program source is binary. + */ + const Program &load_program(const std::string &program_name, const std::string &program_source, bool is_binary) const; + + /** Generates the build options given a string of user defined ones + * + * @param[in] build_options User defined build options + * @param[in] kernel_path Path of the CL kernels + * + * @return Generated build options + */ + std::string generate_build_options(const StringSet &build_options, const std::string &kernel_path) const; + + /** Concatenates contents of a set into a single string. + * + * @param[in] s Input set to concatenate. + * @param[in] kernel_path Path of the CL kernels + * + * @return Concatenated string. + */ + std::string stringify_set(const StringSet &s, const std::string &kernel_path) const; + + cl::Context _context; /**< Underlying CL context. */ + CLDevice _device; /**< Underlying CL device. */ + mutable std::map _programs_map; /**< Map with all already loaded program data. */ + mutable std::map _built_programs_map; /**< Map with all already built program data. */ +}; +} // namespace arm_compute +#endif /* ARM_COMPUTE_CLCOMPILECONTEXT_H */ diff --git a/arm_compute/core/CL/CLDevice.h b/arm_compute/core/CL/CLDevice.h new file mode 100644 index 0000000000..812834743d --- /dev/null +++ b/arm_compute/core/CL/CLDevice.h @@ -0,0 +1,152 @@ +/* + * 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. + */ +#ifndef ARM_COMPUTE_CLDEVICE_H +#define ARM_COMPUTE_CLDEVICE_H + +#include "arm_compute/core/CL/CLHelpers.h" +#include "arm_compute/core/CL/CLTypes.h" +#include "arm_compute/core/GPUTarget.h" +#include "arm_compute/core/IDevice.h" + +#include +#include + +namespace arm_compute +{ +/** OpenCL device type class + * + * Initializes and stores all the information about a cl device, + * working mainly as a cache mechanism. + * */ +class CLDevice : public IDevice +{ +public: + /** Default Constructor */ + CLDevice() + : _device(cl::Device()), _options() + { + } + + /** Constructor + * + * @param[in] cl_device OpenCL device + */ + CLDevice(const cl::Device &cl_device) + : _device(), _options() + { + _device = cl_device; + + // Get device target + std::string device_name = _device.getInfo(); + _options.gpu_target = get_target_from_name(device_name); + + // Fill extensions + std::string extensions = _device.getInfo(); + + std::istringstream iss(extensions); + for(std::string s; iss >> s;) + { + _options.extensions.insert(s); + } + + // SW workaround for G76 + if(_options.gpu_target == GPUTarget::G76) + { + _options.extensions.insert("cl_arm_integer_dot_product_int8"); + } + + // Get device version + _options.version = get_cl_version(_device); + + // Get compute units + _options.compute_units = _device.getInfo(); + + // Get device version + _options.device_version = _device.getInfo(); + } + + /** Returns the GPU target of the cl device + * + * @return The GPU target + */ + const GPUTarget &target() const + { + return _options.gpu_target; + } + + /** Returns the number of compute units available + * + * @return Number of compute units + */ + size_t compute_units() const + { + return _options.compute_units; + } + + /** Returns the underlying cl device object + * + * @return A cl device + */ + const cl::Device &cl_device() const + { + return _device; + } + + /** Returns the device's CL version + * + * @return CLVersion of the device + */ + CLVersion version() const + { + return _options.version; + } + + /** Returns the device version as a string + * + * @return CLVersion of the device + */ + std::string device_version() const + { + return _options.device_version; + } + + // Inherrited methods + DeviceType type() const override + { + return DeviceType::CL; + } + + bool supported(const std::string &extension) const override + { + return _options.extensions.count(extension) != 0; + } + +private: + cl::Device _device; /**< OpenCL device. */ + struct CLDeviceOptions _options; /**< OpenCL device options */ +}; + +} // namespace arm_compute + +#endif /* ARM_COMPUTE_CLDEVICE_H */ diff --git a/arm_compute/core/CL/CLHelpers.h b/arm_compute/core/CL/CLHelpers.h index 453384e14e..ee6397af7a 100644 --- a/arm_compute/core/CL/CLHelpers.h +++ b/arm_compute/core/CL/CLHelpers.h @@ -27,11 +27,13 @@ #include "arm_compute/core/CL/CLTypes.h" #include "arm_compute/core/CL/OpenCL.h" +#include #include namespace arm_compute { class CLCoreRuntimeContext; +class CLCompileContext; class CLBuildOptions; enum class DataType; @@ -196,6 +198,16 @@ bool preferred_dummy_work_items_support(const cl::Device &device); */ cl::Kernel create_opencl_kernel(CLCoreRuntimeContext *ctx, const std::string &kernel_name, const CLBuildOptions &build_opts); +/** Creates an opencl kernel using a compile context + * + * @param[in] ctx A compile context to be used to create the opencl kernel. + * @param[in] kernel_name The kernel name. + * @param[in] build_opts The build options to be used for the opencl kernel compilation. + * + * @return An opencl kernel + */ +cl::Kernel create_kernel(CLCompileContext &ctx, const std::string &kernel_name, const std::set &build_opts); + /** Creates a suitable LWS hint object for parallel implementations. Sets the number of WG based on the input size. * If input width is smaller than 128 we can use fewer threads than 8. * diff --git a/arm_compute/core/CL/CLKernelLibrary.h b/arm_compute/core/CL/CLKernelLibrary.h index 2d55351c95..6c5df6cb08 100644 --- a/arm_compute/core/CL/CLKernelLibrary.h +++ b/arm_compute/core/CL/CLKernelLibrary.h @@ -1,5 +1,5 @@ /* - * Copyright (c) 2016-2019 ARM Limited. + * Copyright (c) 2016-2020 ARM Limited. * * SPDX-License-Identifier: MIT * @@ -24,6 +24,7 @@ #ifndef ARM_COMPUTE_CLKERNELLIBRARY_H #define ARM_COMPUTE_CLKERNELLIBRARY_H +#include "arm_compute/core/CL/CLCompileContext.h" #include "arm_compute/core/CL/OpenCL.h" #include @@ -33,173 +34,18 @@ namespace arm_compute { -/** Build options */ -class CLBuildOptions final -{ - using StringSet = std::set; - -public: - /** Default constructor. */ - CLBuildOptions(); - /** Adds option to the existing build option list - * - * @param[in] option Option to add - */ - void add_option(std::string option); - /** Adds option if a given condition is true; - * - * @param[in] cond Condition to check - * @param[in] option Option to add if condition is true - */ - void add_option_if(bool cond, std::string option); - /** Adds first option if condition is true else the second one - * - * @param[in] cond Condition to check - * @param[in] option_true Option to add if condition is true - * @param[in] option_false Option to add if condition is false - */ - void add_option_if_else(bool cond, std::string option_true, std::string option_false); - /** Appends given build options to the current's objects options. - * - * @param[in] options Build options to append - */ - void add_options(const StringSet &options); - /** Appends given build options to the current's objects options if a given condition is true. - * - * @param[in] cond Condition to check - * @param[in] options Option to add if condition is true - */ - void add_options_if(bool cond, const StringSet &options); - /** Gets the current options list set - * - * @return Build options set - */ - const StringSet &options() const; - -private: - StringSet _build_opts; /**< Build options set */ -}; -/** Program class */ -class Program final -{ -public: - /** Default constructor. */ - Program(); - /** Construct program from source file. - * - * @param[in] context CL context used to create the program. - * @param[in] name Program name. - * @param[in] source Program source. - */ - Program(cl::Context context, std::string name, std::string source); - /** Construct program from binary file. - * - * @param[in] context CL context used to create the program. - * @param[in] device CL device for which the programs are created. - * @param[in] name Program name. - * @param[in] binary Program binary. - */ - Program(cl::Context context, cl::Device device, std::string name, std::vector binary); - /** Default Copy Constructor. */ - Program(const Program &) = default; - /** Default Move Constructor. */ - Program(Program &&) = default; - /** Default copy assignment operator */ - Program &operator=(const Program &) = default; - /** Default move assignment operator */ - Program &operator=(Program &&) = default; - /** Returns program name. - * - * @return Program's name. - */ - std::string name() const - { - return _name; - } - /** User-defined conversion to the underlying CL program. - * - * @return The CL program object. - */ - explicit operator cl::Program() const; - /** Build the given CL program. - * - * @param[in] program The CL program to build. - * @param[in] build_options Options to build the CL program. - * - * @return True if the CL program builds successfully. - */ - static bool build(const cl::Program &program, const std::string &build_options = ""); - /** Build the underlying CL program. - * - * @param[in] build_options Options used to build the CL program. - * - * @return A reference to itself. - */ - cl::Program build(const std::string &build_options = "") const; - -private: - cl::Context _context; /**< Underlying CL context. */ - cl::Device _device; /**< CL device for which the programs are created. */ - bool _is_binary; /**< Create program from binary? */ - std::string _name; /**< Program name. */ - std::string _source; /**< Source code for the program. */ - std::vector _binary; /**< Binary from which to create the program. */ -}; - -/** Kernel class */ -class Kernel final -{ -public: - /** Default Constructor. */ - Kernel(); - /** Default Copy Constructor. */ - Kernel(const Kernel &) = default; - /** Default Move Constructor. */ - Kernel(Kernel &&) = default; - /** Default copy assignment operator */ - Kernel &operator=(const Kernel &) = default; - /** Default move assignment operator */ - Kernel &operator=(Kernel &&) = default; - /** Constructor. - * - * @param[in] name Kernel name. - * @param[in] program Built program. - */ - Kernel(std::string name, const cl::Program &program); - /** Returns kernel name. - * - * @return Kernel's name. - */ - std::string name() const - { - return _name; - } - /** Returns OpenCL kernel. - * - * @return OpenCL Kernel. - */ - explicit operator cl::Kernel() const - { - return _kernel; - } - -private: - std::string _name; /**< Kernel name */ - cl::Kernel _kernel; /**< OpenCL Kernel */ -}; - /** CLKernelLibrary class */ class CLKernelLibrary final { - using StringSet = std::set; - -public: +private: /** Default Constructor. */ CLKernelLibrary(); /** Prevent instances of this class from being copied */ CLKernelLibrary(const CLKernelLibrary &) = delete; /** Prevent instances of this class from being copied */ const CLKernelLibrary &operator=(const CLKernelLibrary &) = delete; + +public: /** Access the KernelLibrary singleton. * This method has been deprecated and will be removed in the next release. * @return The KernelLibrary instance. @@ -224,17 +70,9 @@ public: * * @param[in] program_name Program name. * - * @return Source of the selected program. + * @return A pair with the source (false) or the binary (true), of the selected program. */ - std::string get_program_source(const std::string &program_name); - /** Sets the CL context used to create programs. - * - * @note Setting the context also resets the device to the - * first one available in the new context. - * - * @param[in] context A CL context. - */ - void set_context(cl::Context context); + std::pair get_program(const std::string &program_name) const; /** Accessor for the associated CL context. * @@ -243,7 +81,7 @@ public: cl::Context &context(); /** Gets the CL device for which the programs are created. */ - cl::Device &get_device(); + const cl::Device &get_device(); /** Sets the CL device for which the programs are created. * @@ -268,7 +106,7 @@ public: * * @return The created kernel. */ - Kernel create_kernel(const std::string &kernel_name, const StringSet &build_options_set = {}) const; + Kernel create_kernel(const std::string &kernel_name, const std::set &build_options_set = {}) const; /** Find the maximum number of local work items in a workgroup can be supported for the kernel. * */ @@ -304,28 +142,33 @@ public: */ bool int64_base_atomics_supported() const; -private: - /** Load program and its dependencies. + /** Returns the program name given a kernel name * - * @param[in] program_name Name of the program to load. + * @return Program name */ - const Program &load_program(const std::string &program_name) const; - /** Concatenates contents of a set into a single string. + std::string get_program_name(const std::string &kernel_name) const; + + /** Sets the CL context used to create programs. * - * @param[in] s Input set to concatenate. + * @note Setting the context also resets the device to the + * first one available in the new context. * - * @return Concatenated string. + * @param[in] context A CL context. */ - std::string stringify_set(const StringSet &s) const; + void set_context(cl::Context context); - cl::Context _context; /**< Underlying CL context. */ - cl::Device _device; /**< Underlying CL device. */ - std::string _kernel_path; /**< Path to the kernels folder. */ - mutable std::map _programs_map; /**< Map with all already loaded program data. */ - mutable std::map _built_programs_map; /**< Map with all already built program data. */ + /** Gets the compile context used + * + * @return The used compile context + */ + CLCompileContext &get_compile_context(); + +private: + CLCompileContext _compile_context; /**< Compile Context. */ + std::string _kernel_path; /**< Path to the kernels folder. */ static const std::map _kernel_program_map; /**< Map that associates kernel names with programs. */ static const std::map _program_source_map; /**< Contains sources for all programs. - Used for compile-time kernel inclusion. >*/ + Used for compile-time kernel inclusion. >*/ }; } // namespace arm_compute #endif /* ARM_COMPUTE_CLKERNELLIBRARY_H */ diff --git a/arm_compute/core/CL/CLTypes.h b/arm_compute/core/CL/CLTypes.h index 9f6ff6a82b..3643b178d3 100644 --- a/arm_compute/core/CL/CLTypes.h +++ b/arm_compute/core/CL/CLTypes.h @@ -1,5 +1,5 @@ /* - * Copyright (c) 2017-2019 ARM Limited. + * Copyright (c) 2017-2020 ARM Limited. * * SPDX-License-Identifier: MIT * @@ -27,6 +27,7 @@ #include "arm_compute/core/CL/ICLArray.h" #include "arm_compute/core/GPUTarget.h" +#include #include namespace arm_compute @@ -47,12 +48,14 @@ enum class CLVersion /** OpenCL device options */ struct CLDeviceOptions { - std::string name; /**< Device name */ - std::string extensions; /**< List of supported extensions */ - std::string ddk_version; /**< DDK version */ - GPUTarget gpu_target; /**< GPU target architecture/instance */ - size_t num_cores; /**< Number of cores */ - size_t cache_size; /**< Cache size */ + std::string name{}; /**< Device name */ + std::string device_version{}; /**< Device version string */ + std::set extensions{}; /**< List of supported extensions */ + std::string ddk_version{}; /**< DDK version */ + GPUTarget gpu_target{}; /**< GPU target architecture/instance */ + CLVersion version{}; /**< Device OpenCL version */ + size_t compute_units{}; /**< Number of compute units */ + size_t cache_size{}; /**< Cache size */ }; /** OpenCL quantization data */ diff --git a/arm_compute/core/CL/kernels/CLFloorKernel.h b/arm_compute/core/CL/kernels/CLFloorKernel.h index 00f77dcd6b..a3ccb96c61 100644 --- a/arm_compute/core/CL/kernels/CLFloorKernel.h +++ b/arm_compute/core/CL/kernels/CLFloorKernel.h @@ -1,5 +1,5 @@ /* - * Copyright (c) 2017-2019 ARM Limited. + * Copyright (c) 2017-2020 ARM Limited. * * SPDX-License-Identifier: MIT * @@ -52,6 +52,15 @@ public: * @param[out] output Destination tensor. Same as @p input */ void configure(const ICLTensor *input, ICLTensor *output); + + /** Set the source, destination of the kernel + * + * @param[in] compile_context The compile context to be used. + * @param[in] input Source tensor. Data type supported: F16/F32. + * @param[out] output Destination tensor. Same as @p input + */ + void configure(CLCompileContext &compile_context, const ICLTensor *input, ICLTensor *output); + /** Static function to check if given info will lead to a valid configuration of @ref CLFloorKernel * * @param[in] input Source tensor info. Data type supported: F16/F32. diff --git a/arm_compute/core/IDevice.h b/arm_compute/core/IDevice.h new file mode 100644 index 0000000000..5cffe646d4 --- /dev/null +++ b/arm_compute/core/IDevice.h @@ -0,0 +1,60 @@ +/* + * 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. + */ + +#ifndef ARM_COMPUTE_IDEVICE_H +#define ARM_COMPUTE_IDEVICE_H + +#include + +namespace arm_compute +{ +/** Device types */ +enum class DeviceType +{ + NEON, + CL, + GLES +}; + +/** Interface for device object */ +class IDevice +{ +public: + /** Virtual Destructor */ + virtual ~IDevice() = default; + + /** Device type accessor */ + virtual DeviceType type() const = 0; + + /** Check if extensions on a device are supported + * + * @param[in] extension An extension to check if it's supported. + * + * @return True if the extension is supported else false + */ + virtual bool supported(const std::string &extension) const = 0; +}; +} // namespace arm_compute + +#endif /* ARM_COMPUTE_IDEVICE_H */ diff --git a/arm_compute/runtime/CL/CLRuntimeContext.h b/arm_compute/runtime/CL/CLRuntimeContext.h index 791d1deaa7..54c7d3cd23 100644 --- a/arm_compute/runtime/CL/CLRuntimeContext.h +++ b/arm_compute/runtime/CL/CLRuntimeContext.h @@ -58,7 +58,6 @@ private: std::unique_ptr _gpu_owned_scheduler{ nullptr }; CLScheduler *_gpu_scheduler{ nullptr }; CLTuner _tuner{ false }; - CLKernelLibrary _kernel_lib{}; CLSymbols _symbols{}; CLCoreRuntimeContext _core_context{}; }; diff --git a/arm_compute/runtime/CL/functions/CLFloor.h b/arm_compute/runtime/CL/functions/CLFloor.h index 33d03217e7..c4a893fdeb 100644 --- a/arm_compute/runtime/CL/functions/CLFloor.h +++ b/arm_compute/runtime/CL/functions/CLFloor.h @@ -1,5 +1,5 @@ /* - * Copyright (c) 2017-2019 ARM Limited. + * Copyright (c) 2017-2020 ARM Limited. * * SPDX-License-Identifier: MIT * @@ -42,6 +42,13 @@ public: * @param[out] output Destination tensor. Same as @p input */ void configure(const ICLTensor *input, ICLTensor *output); + /** Set the source, destination of the kernel + * + * @param[in] compile_context The compile context to be used. + * @param[in] input Source tensor. Data type supported: F16/F32. + * @param[out] output Destination tensor. Same as @p input + */ + void configure(CLCompileContext &compile_context, const ICLTensor *input, ICLTensor *output); /** Static function to check if given info will lead to a valid configuration of @ref CLFloor * * @param[in] input Source tensor info. Data type supported: F16/F32. 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)); diff --git a/src/runtime/CL/CLRuntimeContext.cpp b/src/runtime/CL/CLRuntimeContext.cpp index 49e4c10c84..4d70edac2f 100644 --- a/src/runtime/CL/CLRuntimeContext.cpp +++ b/src/runtime/CL/CLRuntimeContext.cpp @@ -1,5 +1,5 @@ /* - * Copyright (c) 2019 ARM Limited. + * Copyright (c) 2019-2020 ARM Limited. * * SPDX-License-Identifier: MIT * @@ -39,13 +39,13 @@ CLRuntimeContext::CLRuntimeContext() cl::CommandQueue queue = cl::CommandQueue(ctx, dev); _gpu_owned_scheduler->init(ctx, queue, dev, &_tuner); const std::string cl_kernels_folder("./cl_kernels"); - _kernel_lib.init(cl_kernels_folder, ctx, dev); - _core_context = CLCoreRuntimeContext(&_kernel_lib, _gpu_owned_scheduler->context(), _gpu_owned_scheduler->queue()); + CLKernelLibrary::get().init(cl_kernels_folder, ctx, dev); + _core_context = CLCoreRuntimeContext(&CLKernelLibrary::get(), _gpu_owned_scheduler->context(), _gpu_owned_scheduler->queue()); } CLKernelLibrary &CLRuntimeContext::kernel_library() { - return _kernel_lib; + return CLKernelLibrary::get(); } CLCoreRuntimeContext *CLRuntimeContext::core_runtime_context() diff --git a/src/runtime/CL/functions/CLFloor.cpp b/src/runtime/CL/functions/CLFloor.cpp index 525810d400..204ca7400c 100644 --- a/src/runtime/CL/functions/CLFloor.cpp +++ b/src/runtime/CL/functions/CLFloor.cpp @@ -29,9 +29,14 @@ namespace arm_compute { void CLFloor::configure(const ICLTensor *input, ICLTensor *output) +{ + configure(CLKernelLibrary::get().get_compile_context(), input, output); +} + +void CLFloor::configure(CLCompileContext &compile_context, const ICLTensor *input, ICLTensor *output) { auto k = arm_compute::support::cpp14::make_unique(); - k->configure(input, output); + k->configure(compile_context, input, output); _kernel = std::move(k); } diff --git a/tests/validation/CL/UNIT/CompileContext.cpp b/tests/validation/CL/UNIT/CompileContext.cpp new file mode 100644 index 0000000000..5245044323 --- /dev/null +++ b/tests/validation/CL/UNIT/CompileContext.cpp @@ -0,0 +1,74 @@ +/* + * 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 "tests/AssetsLibrary.h" +#include "tests/Globals.h" +#include "tests/Utils.h" +#include "tests/framework/Asserts.h" +#include "tests/framework/Macros.h" +#include "tests/validation/Validation.h" + +namespace arm_compute +{ +namespace test +{ +namespace validation +{ +TEST_SUITE(CL) +TEST_SUITE(UNIT) +TEST_SUITE(CompileContext) + +TEST_CASE(CompileContextCache, framework::DatasetMode::ALL) +{ + // Create compile context + CLCompileContext compile_context(CLKernelLibrary::get().context(), CLKernelLibrary::get().get_device()); + + // Check if the program cache is empty + ARM_COMPUTE_EXPECT(compile_context.get_built_programs().size() == 0, framework::LogLevel::ERRORS); + + // Create a kernel using the compile context + const std::string kernel_name = "floor_layer"; + 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(); + + std::set build_opts; + build_opts.emplace("-DDATA_TYPE=float"); + build_opts.emplace("-DVEC_SIZE=16"); + compile_context.create_kernel(kernel_name, program_name, kernel_src.first, kernel_path, build_opts, kernel_src.second); + + // Check if the program is stored in the cache + ARM_COMPUTE_EXPECT(compile_context.get_built_programs().size() == 1, framework::LogLevel::ERRORS); + + // Try to build the same program and check if the program cache stayed the same + compile_context.create_kernel(kernel_name, program_name, kernel_src.first, kernel_path, build_opts, kernel_src.second); + ARM_COMPUTE_EXPECT(compile_context.get_built_programs().size() == 1, framework::LogLevel::ERRORS); +} + +TEST_SUITE_END() // CompileContext +TEST_SUITE_END() // UNIT +TEST_SUITE_END() // CL +} // namespace validation +} // namespace test +} // namespace arm_compute -- cgit v1.2.1