Compute Library
 21.05
CLTuner.h
Go to the documentation of this file.
1 /*
2  * Copyright (c) 2017-2021 Arm Limited.
3  *
4  * SPDX-License-Identifier: MIT
5  *
6  * Permission is hereby granted, free of charge, to any person obtaining a copy
7  * of this software and associated documentation files (the "Software"), to
8  * deal in the Software without restriction, including without limitation the
9  * rights to use, copy, modify, merge, publish, distribute, sublicense, and/or
10  * sell copies of the Software, and to permit persons to whom the Software is
11  * furnished to do so, subject to the following conditions:
12  *
13  * The above copyright notice and this permission notice shall be included in all
14  * copies or substantial portions of the Software.
15  *
16  * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
17  * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
18  * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
19  * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
20  * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
21  * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE
22  * SOFTWARE.
23  */
24 #ifndef ARM_COMPUTE_CLTUNER_H
25 #define ARM_COMPUTE_CLTUNER_H
26 
32 
33 #include <unordered_map>
34 
35 namespace arm_compute
36 {
37 class ICLKernel;
38 
39 /** Basic implementation of the OpenCL tuner interface */
40 class CLTuner : public ICLTuner
41 {
42 public:
43  /** Constructor
44  *
45  * @param[in] tune_new_kernels Find the optimal local workgroup size for kernels which are not present in the table ?
46  * @param[in] tuning_info (Optional) opencl parameters to tune
47  *
48  */
49  CLTuner(bool tune_new_kernels = true, CLTuningInfo tuning_info = CLTuningInfo());
50 
51  /** Destructor */
52  ~CLTuner() = default;
53 
54  /** Setter for tune_new_kernels option
55  *
56  * @param[in] tune_new_kernels Find the optimal local workgroup size for kernels which are not present in the table ?
57  */
59 
60  /** Tune kernels that are not in the tuning parameters table
61  *
62  * @return True if tuning of new kernels is enabled.
63  */
64  bool tune_new_kernels() const;
65 
66  /** Setter for tune parameters option
67  *
68  * @param[in] tuning_info opencl parameters to tune
69  */
70  void set_tuning_parameters(CLTuningInfo tuning_info);
71 
72  /** Set OpenCL tuner mode
73  *
74  * @param[in] mode Indicates how exhaustive the search for the optimal tuning parameters should be while tuning. Default is Exhaustive mode
75  */
77 
78  /** Get the current OpenCL tuner mode
79  *
80  * @return tuner_mode Indicates how exhaustive the search for the optimal tuning parameters should be while tuning
81  *
82  * @deprecated This function is deprecated and is intended to be removed in 21.08 release
83  */
85 
86  /** Manually add a LWS for a kernel
87  *
88  * @param[in] kernel_id Unique identifiant of the kernel
89  * @param[in] optimal_lws Optimal local workgroup size to use for the given kernel
90  *
91  * @deprecated This function is deprecated and is intended to be removed in 21.08 release
92  */
94  void add_lws_to_table(const std::string &kernel_id, cl::NDRange optimal_lws);
95 
96  /** Manually add tuning parameters for a kernel
97  *
98  * @param[in] kernel_id Unique identifiant of the kernel
99  * @param[in] optimal_tuning_params Optimal tuning parameters to use for the given kernel
100  */
101  void add_tuning_params(const std::string &kernel_id, CLTuningParams optimal_tuning_params);
102 
103  /** Import LWS table
104  *
105  * @param[in] lws_table The unordered_map container to import
106  *
107  * @deprecated This function is deprecated and is intended to be removed in 21.08 release
108  */
110  void import_lws_table(const std::unordered_map<std::string, cl::NDRange> &lws_table);
111 
112  /** Import tuning parameters table
113  *
114  * @param[in] tuning_params_table The unordered_map container to import
115  */
116  void import_tuning_params(const std::unordered_map<std::string, CLTuningParams> &tuning_params_table);
117 
118  /** Give read access to the LWS table
119  *
120  * @return The lws table as unordered_map container
121  *
122  * @deprecated This function is deprecated and is intended to be removed in 21.08 release
123  */
125  const std::unordered_map<std::string, cl::NDRange> &lws_table();
126 
127  /** Give read access to the tuning params table
128  *
129  * @return The tuning params table as unordered_map container
130  */
131  const std::unordered_map<std::string, CLTuningParams> &tuning_params_table() const;
132 
133  /** Set the OpenCL kernel event
134  *
135  * @note The interceptor can use this function to store the event associated to the OpenCL kernel
136  *
137  * @param[in] kernel_event The OpenCL kernel event
138  */
139  void set_cl_kernel_event(cl_event kernel_event);
140 
141  /** clEnqueueNDRangeKernel symbol */
143 
144  /** Load the tuning parameters table from file. It also sets up the tuning read from the file
145  *
146  * @param[in] filename Load the tuning parameters table from this file.(Must exist)
147  *
148  */
149  void load_from_file(const std::string &filename);
150 
151  /** Save the content of the tuning parameters table to file
152  *
153  * @param[in] filename Save the tuning parameters table to this file. (Content will be overwritten)
154  *
155  * @return true if the file was created
156  */
157  bool save_to_file(const std::string &filename) const;
158 
159  // Inherited methods overridden:
160  void tune_kernel_static(ICLKernel &kernel) override;
161  void tune_kernel_dynamic(ICLKernel &kernel) override;
162  void tune_kernel_dynamic(ICLKernel &kernel, ITensorPack &tensors) override;
163 
164  /** Is the kernel_event set ?
165  *
166  * @return true if the kernel_event is set.
167  */
168  bool kernel_event_is_set() const;
169 
170 private:
171  /** Find optimal tuning parameters using brute-force approach
172  *
173  * @param[in] kernel OpenCL kernel to be tuned with tuning parameters
174  * @param[in,out] tensors Tensors for the kernel to operate on
175  *
176  * @return The optimal tuning parameters to use
177  */
178  CLTuningParams find_optimal_tuning_params(ICLKernel &kernel, ITensorPack &tensors);
179 
180  std::unordered_map<std::string, CLTuningParams> _tuning_params_table;
181  std::unordered_map<std::string, cl::NDRange> _lws_table;
182  cl::Event _kernel_event;
183  bool _tune_new_kernels;
184  CLTuningInfo _tuning_info;
185 };
186 } // namespace arm_compute
187 #endif /*ARM_COMPUTE_CLTUNER_H */
Basic implementation of the OpenCL tuner interface.
Definition: CLTuner.h:40
void set_tuner_mode(CLTunerMode mode)
Set OpenCL tuner mode.
Definition: CLTuner.cpp:61
void set_tuning_parameters(CLTuningInfo tuning_info)
Setter for tune parameters option.
void set_cl_kernel_event(cl_event kernel_event)
Set the OpenCL kernel event.
Definition: CLTuner.cpp:47
const std::unordered_map< std::string, CLTuningParams > & tuning_params_table() const
Give read access to the tuning params table.
Definition: CLTuner.cpp:269
void import_tuning_params(const std::unordered_map< std::string, CLTuningParams > &tuning_params_table)
Import tuning parameters table.
Definition: CLTuner.cpp:274
void tune_kernel_static(ICLKernel &kernel) override
Tune OpenCL kernel statically.
Definition: CLTuner.cpp:71
void tune_kernel_dynamic(ICLKernel &kernel) override
Tune OpenCL kernel dynamically.
Definition: CLTuner.cpp:76
void load_from_file(const std::string &filename)
Load the tuning parameters table from file.
Definition: CLTuner.cpp:280
const std::unordered_map< std::string, cl::NDRange > & lws_table()
Give read access to the LWS table.
Definition: CLTuner.cpp:259
< OpenCL tuner parameters
~CLTuner()=default
Destructor.
Common interface for all the OpenCL kernels.
Definition: ICLKernel.h:46
Copyright (c) 2017-2021 Arm Limited.
CLTuner(bool tune_new_kernels=true, CLTuningInfo tuning_info=CLTuningInfo())
Constructor.
Definition: CLTuner.cpp:38
std::function< decltype(clEnqueueNDRangeKernel)> real_clEnqueueNDRangeKernel
clEnqueueNDRangeKernel symbol
Definition: CLTuner.h:142
bool kernel_event_is_set() const
Is the kernel_event set ?
Definition: CLTuner.cpp:43
bool save_to_file(const std::string &filename) const
Save the content of the tuning parameters table to file.
Definition: CLTuner.cpp:332
Interface used to tune the local work-group size of OpenCL kernels.
void add_tuning_params(const std::string &kernel_id, CLTuningParams optimal_tuning_params)
Manually add tuning parameters for a kernel.
Definition: CLTuner.cpp:127
void set_tune_new_kernels(bool tune_new_kernels)
Setter for tune_new_kernels option.
Definition: CLTuner.cpp:52
void import_lws_table(const std::unordered_map< std::string, cl::NDRange > &lws_table)
Import LWS table.
Definition: CLTuner.cpp:250
cl_int clEnqueueNDRangeKernel(cl_command_queue command_queue, cl_kernel kernel, cl_uint work_dim, const size_t *global_work_offset, const size_t *global_work_size, const size_t *local_work_size, cl_uint num_events_in_wait_list, const cl_event *event_wait_list, cl_event *event)
Definition: OpenCL.cpp:372
CLTunerMode
< OpenCL tuner modes
Definition: CLTunerTypes.h:35
Wrapper to configure the Khronos OpenCL C++ header.
Tensor packing service.
Definition: ITensorPack.h:37
Basic interface for tuning the OpenCL kernels.
Definition: ICLTuner.h:34
#define ARM_COMPUTE_DEPRECATED_REL_REPLACE(rel, replace)
Definition: Macros.h:43
bool tune_new_kernels() const
Tune kernels that are not in the tuning parameters table.
Definition: CLTuner.cpp:56
void add_lws_to_table(const std::string &kernel_id, cl::NDRange optimal_lws)
Manually add a LWS for a kernel.
Definition: CLTuner.cpp:122
CLTunerMode get_tuner_mode() const
Get the current OpenCL tuner mode.
Definition: CLTuner.cpp:66