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 --- 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 +- 6 files changed, 536 insertions(+), 193 deletions(-) create mode 100644 arm_compute/core/CL/CLCompileContext.h create mode 100644 arm_compute/core/CL/CLDevice.h (limited to 'arm_compute/core/CL') 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. -- cgit v1.2.1