diff options
Diffstat (limited to 'arm_compute/runtime/CL/CLTuner.h')
| -rw-r--r-- | arm_compute/runtime/CL/CLTuner.h | 86 |
1 files changed, 53 insertions, 33 deletions
diff --git a/arm_compute/runtime/CL/CLTuner.h b/arm_compute/runtime/CL/CLTuner.h index 3b7a42fa6a..cf293d3d27 100644 --- a/arm_compute/runtime/CL/CLTuner.h +++ b/arm_compute/runtime/CL/CLTuner.h @@ -1,5 +1,5 @@ /* - * Copyright (c) 2017-2019 ARM Limited. + * Copyright (c) 2017-2022 Arm Limited. * * SPDX-License-Identifier: MIT * @@ -25,7 +25,9 @@ #define ARM_COMPUTE_CLTUNER_H #include "arm_compute/core/CL/OpenCL.h" +#include "arm_compute/core/utils/misc/Macros.h" #include "arm_compute/runtime/CL/CLTunerTypes.h" +#include "arm_compute/runtime/CL/CLTuningParams.h" #include "arm_compute/runtime/CL/ICLTuner.h" #include <unordered_map> @@ -41,9 +43,10 @@ public: /** Constructor * * @param[in] tune_new_kernels Find the optimal local workgroup size for kernels which are not present in the table ? + * @param[in] tuning_info (Optional) opencl parameters to tune * */ - CLTuner(bool tune_new_kernels = true); + CLTuner(bool tune_new_kernels = true, CLTuningInfo tuning_info = CLTuningInfo()); /** Destructor */ ~CLTuner() = default; @@ -53,42 +56,43 @@ public: * @param[in] tune_new_kernels Find the optimal local workgroup size for kernels which are not present in the table ? */ void set_tune_new_kernels(bool tune_new_kernels); - /** Tune kernels that are not in the LWS table + + /** Tune kernels that are not in the tuning parameters table * * @return True if tuning of new kernels is enabled. */ bool tune_new_kernels() const; - /** Set OpenCL tuner mode + /** Setter for tune parameters option * - * @param[in] mode Indicates how exhaustive the search for the optimal LWS should be while tuning. Default is Exhaustive mode + * @param[in] tuning_info opencl parameters to tune */ - void set_tuner_mode(CLTunerMode mode); + void set_tuning_parameters(CLTuningInfo tuning_info); - /** Get the current OpenCL tuner mode + /** Set OpenCL tuner mode * - * @return tuner_mode Indicates how exhaustive the search for the optimal LWS should be while tuning + * @param[in] mode Indicates how exhaustive the search for the optimal tuning parameters should be while tuning. Default is Exhaustive mode */ - CLTunerMode get_tuner_mode() const; + void set_tuner_mode(CLTunerMode mode); - /** Manually add a LWS for a kernel + /** Manually add tuning parameters for a kernel * - * @param[in] kernel_id Unique identifiant of the kernel - * @param[in] optimal_lws Optimal local workgroup size to use for the given kernel + * @param[in] kernel_id Unique identifiant of the kernel + * @param[in] optimal_tuning_params Optimal tuning parameters to use for the given kernel */ - void add_lws_to_table(const std::string &kernel_id, cl::NDRange optimal_lws); + void add_tuning_params(const std::string &kernel_id, CLTuningParams optimal_tuning_params); - /** Import LWS table + /** Import tuning parameters table * - * @param[in] lws_table The unordered_map container to import + * @param[in] tuning_params_table The unordered_map container to import */ - void import_lws_table(const std::unordered_map<std::string, cl::NDRange> &lws_table); + void import_tuning_params(const std::unordered_map<std::string, CLTuningParams> &tuning_params_table); - /** Give read access to the LWS table + /** Give read access to the tuning params table * - * @return The lws table as unordered_map container + * @return The tuning params table as unordered_map container */ - const std::unordered_map<std::string, cl::NDRange> &lws_table() const; + const std::unordered_map<std::string, CLTuningParams> &tuning_params_table() const; /** Set the OpenCL kernel event * @@ -101,41 +105,57 @@ public: /** clEnqueueNDRangeKernel symbol */ std::function<decltype(clEnqueueNDRangeKernel)> real_clEnqueueNDRangeKernel; - /** Load the LWS table from file + /** Load the tuning parameters table from file. It also sets up the tuning read from the file + * + * @param[in] filename Load the tuning parameters table from this file.(Must exist) * - * @param[in] filename Load the LWS table from this file.(Must exist) */ void load_from_file(const std::string &filename); - /** Save the content of the LWS table to file + /** Save the content of the tuning parameters table to file + * + * @param[in] filename Save the tuning parameters table to this file. (Content will be overwritten) * - * @param[in] filename Save the LWS table to this file. (Content will be overwritten) + * @return true if the file was created */ - void save_to_file(const std::string &filename) const; + bool save_to_file(const std::string &filename) const; // Inherited methods overridden: void tune_kernel_static(ICLKernel &kernel) override; void tune_kernel_dynamic(ICLKernel &kernel) override; - + void tune_kernel_dynamic(ICLKernel &kernel, ITensorPack &tensors) override; /** Is the kernel_event set ? * * @return true if the kernel_event is set. */ bool kernel_event_is_set() const; + /** A wrapper wrapping tensors and other objects needed for running the kernel + */ + struct IKernelData; + private: - /** Find optimal LWS using brute-force approach + /** Perform tune_kernel_dynamic + * + * @param[in] kernel OpenCL kernel to be tuned with tuning parameters + * @param[in,out] data IKernelData object wrapping tensors and other objects needed for running the kernel + * + */ + void do_tune_kernel_dynamic(ICLKernel &kernel, IKernelData *data); + /** Find optimal tuning parameters using brute-force approach * - * @param[in] kernel OpenCL kernel to be tuned with LWS + * @param[in] kernel OpenCL kernel to be tuned with tuning parameters + * @param[in,out] data IKernelData object wrapping tensors and other objects needed for running the kernel * - * @return The optimal LWS to use + * @return The optimal tuning parameters to use */ - cl::NDRange find_optimal_lws(ICLKernel &kernel); + CLTuningParams find_optimal_tuning_params(ICLKernel &kernel, IKernelData *data); - std::unordered_map<std::string, cl::NDRange> _lws_table; - cl::Event _kernel_event; - bool _tune_new_kernels; - CLTunerMode _tuner_mode; + std::unordered_map<std::string, CLTuningParams> _tuning_params_table; + std::unordered_map<std::string, cl::NDRange> _lws_table; + cl::Event _kernel_event; + bool _tune_new_kernels; + CLTuningInfo _tuning_info; }; } // namespace arm_compute #endif /*ARM_COMPUTE_CLTUNER_H */ |