"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",
--- /dev/null
+/*
+ * 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 <map>
+#include <set>
+#include <string>
+#include <utility>
+
+namespace arm_compute
+{
+/** Build options */
+class CLBuildOptions final
+{
+ using StringSet = std::set<std::string>;
+
+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<unsigned char> 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<unsigned char> _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<std::string>;
+
+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<std::string, cl::Program> &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<std::string, const Program> _programs_map; /**< Map with all already loaded program data. */
+ mutable std::map<std::string, cl::Program> _built_programs_map; /**< Map with all already built program data. */
+};
+} // namespace arm_compute
+#endif /* ARM_COMPUTE_CLCOMPILECONTEXT_H */
--- /dev/null
+/*
+ * 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 <set>
+#include <string>
+
+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<CL_DEVICE_NAME>();
+ _options.gpu_target = get_target_from_name(device_name);
+
+ // Fill extensions
+ std::string extensions = _device.getInfo<CL_DEVICE_EXTENSIONS>();
+
+ 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<CL_DEVICE_MAX_COMPUTE_UNITS>();
+
+ // Get device version
+ _options.device_version = _device.getInfo<CL_DEVICE_VERSION>();
+ }
+
+ /** 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 */
#include "arm_compute/core/CL/CLTypes.h"
#include "arm_compute/core/CL/OpenCL.h"
+#include <set>
#include <string>
namespace arm_compute
{
class CLCoreRuntimeContext;
+class CLCompileContext;
class CLBuildOptions;
enum class DataType;
*/
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<std::string> &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.
*
/*
- * Copyright (c) 2016-2019 ARM Limited.
+ * Copyright (c) 2016-2020 ARM Limited.
*
* SPDX-License-Identifier: MIT
*
#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 <map>
namespace arm_compute
{
-/** Build options */
-class CLBuildOptions final
-{
- using StringSet = std::set<std::string>;
-
-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<unsigned char> 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<unsigned char> _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<std::string>;
-
-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.
*
* @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<std::string, bool> get_program(const std::string &program_name) const;
/** Accessor for the associated CL context.
*
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.
*
*
* @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<std::string> &build_options_set = {}) const;
/** Find the maximum number of local work items in a workgroup can be supported for the kernel.
*
*/
*/
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<std::string, const Program> _programs_map; /**< Map with all already loaded program data. */
- mutable std::map<std::string, cl::Program> _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<std::string, std::string> _kernel_program_map; /**< Map that associates kernel names with programs. */
static const std::map<std::string, std::string> _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 */
/*
- * Copyright (c) 2017-2019 ARM Limited.
+ * Copyright (c) 2017-2020 ARM Limited.
*
* SPDX-License-Identifier: MIT
*
#include "arm_compute/core/CL/ICLArray.h"
#include "arm_compute/core/GPUTarget.h"
+#include <set>
#include <string>
namespace arm_compute
/** 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<std::string> 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 */
/*
- * Copyright (c) 2017-2019 ARM Limited.
+ * Copyright (c) 2017-2020 ARM Limited.
*
* SPDX-License-Identifier: MIT
*
* @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.
--- /dev/null
+/*
+ * 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 <string>
+
+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 */
std::unique_ptr<CLScheduler> _gpu_owned_scheduler{ nullptr };
CLScheduler *_gpu_scheduler{ nullptr };
CLTuner _tuner{ false };
- CLKernelLibrary _kernel_lib{};
CLSymbols _symbols{};
CLCoreRuntimeContext _core_context{};
};
/*
- * Copyright (c) 2017-2019 ARM Limited.
+ * Copyright (c) 2017-2020 ARM Limited.
*
* SPDX-License-Identifier: MIT
*
* @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.
--- /dev/null
+/*
+ * 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<unsigned char> 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<CL_PROGRAM_BUILD_LOG>(&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<cl::Program>(*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<unsigned char>(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<CL_CONTEXT_DEVICES>();
+
+ 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<std::underlying_type<GPUTarget>::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<std::string, cl::Program> &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
}
}
+cl::Kernel create_kernel(CLCompileContext &ctx, const std::string &kernel_name, const std::set<std::string> &build_opts)
+{
+ const std::string program_name = CLKernelLibrary::get().get_program_name(kernel_name);
+ std::pair<std::string, bool> kernel_src = CLKernelLibrary::get().get_program(program_name);
+ const std::string kernel_path = CLKernelLibrary::get().get_kernel_path();
+ return static_cast<cl::Kernel>(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;
#include <vector>
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<unsigned char> 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<CL_PROGRAM_BUILD_LOG>(&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<cl::Program>(*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<std::string, std::string> CLKernelLibrary::_kernel_program_map =
{
{ "absdiff", "absdiff.cl" },
};
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
}
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<std::string> &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);
{
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<std::underlying_type<GPUTarget>::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()
void CLKernelLibrary::clear_programs_cache()
{
- _programs_map.clear();
- _built_programs_map.clear();
+ _compile_context.clear_programs_cache();
}
const std::map<std::string, cl::Program> &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<std::string, bool> 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<unsigned char>(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<CL_CONTEXT_DEVICES>();
-
- 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<CL_DEVICE_VERSION>();
+ return _compile_context.get_device_version();
}
cl_uint CLKernelLibrary::get_num_compute_units()
{
- return _device.getInfo<CL_DEVICE_MAX_COMPUTE_UNITS>();
+ return _compile_context.get_num_compute_units();
+}
+
+CLCompileContext &CLKernelLibrary::get_compile_context()
+{
+ return _compile_context;
}
{
}
-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);
_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<std::string> 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<cl::Kernel>(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());
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));
/*
- * Copyright (c) 2019 ARM Limited.
+ * Copyright (c) 2019-2020 ARM Limited.
*
* SPDX-License-Identifier: MIT
*
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()
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<CLFloorKernel>();
- k->configure(input, output);
+ k->configure(compile_context, input, output);
_kernel = std::move(k);
}
--- /dev/null
+/*
+ * 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<std::string, bool> kernel_src = CLKernelLibrary::get().get_program(program_name);
+ const std::string kernel_path = CLKernelLibrary::get().get_kernel_path();
+
+ std::set<std::string> 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
projects / platform / upstream / armcl.git / commitdiff