diff options
author | Anthony Barbier <anthony.barbier@arm.com> | 2018-02-28 13:47:58 +0000 |
---|---|---|
committer | Anthony Barbier <anthony.barbier@arm.com> | 2018-11-02 16:48:33 +0000 |
commit | 8b81195919eaab90a803fabae57ca05136e9430e (patch) | |
tree | 325798fd319f8091b18458127d3843b88c4d40af | |
parent | fadfc6d6737716891f543f61a529e984da9b0a8b (diff) | |
download | ComputeLibrary-8b81195919eaab90a803fabae57ca05136e9430e.tar.gz |
COMPMID-978 Load/Store tuning data from file (Part2)
Change-Id: I1819f42c0e456673543b267d51f730b6e80a0ad9
Reviewed-on: https://eu-gerrit-1.euhpc.arm.com/122629
Reviewed-by: Robert Hughes <robert.hughes@arm.com>
Tested-by: Jenkins <bsgcomp@arm.com>
Reviewed-by: Georgios Pinitas <georgios.pinitas@arm.com>
-rw-r--r-- | arm_compute/runtime/CL/CLTuner.h | 53 | ||||
-rw-r--r-- | src/graph/Graph.cpp | 25 | ||||
-rw-r--r-- | src/runtime/CL/CLTuner.cpp | 45 |
3 files changed, 57 insertions, 66 deletions
diff --git a/arm_compute/runtime/CL/CLTuner.h b/arm_compute/runtime/CL/CLTuner.h index 251848814d..a321241fee 100644 --- a/arm_compute/runtime/CL/CLTuner.h +++ b/arm_compute/runtime/CL/CLTuner.h @@ -47,11 +47,16 @@ public: /** Destructor */ ~CLTuner() = default; - /* Setter for tune_new_kernels option + /** Setter for tune_new_kernels option * * @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 + * + * @return True if tuning of new kernels is enabled. + */ + bool tune_new_kernels() const; /** Manually add a LWS for a kernel * * @param[in] kernel_id Unique identifiant of the kernel @@ -64,11 +69,11 @@ public: */ void import_lws_table(const std::unordered_map<std::string, cl::NDRange> &lws_table); - /** Export LWS table + /** Give read access to the LWS table * * return The lws table as unordered_map container */ - const std::unordered_map<std::string, cl::NDRange> &export_lws_table() const; + const std::unordered_map<std::string, cl::NDRange> &lws_table() const; /** Set the OpenCL kernel event * @@ -80,6 +85,18 @@ public: std::function<decltype(clEnqueueNDRangeKernel)> real_clEnqueueNDRangeKernel; + /** Load the LWS table from file + * + * @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 + * + * @param[in] filename Save the LWS table to this file. (Content will be overwritten) + */ + void save_to_file(const std::string &filename) const; + // Inherited methods overridden: void tune_kernel(ICLKernel &kernel) override; @@ -98,35 +115,5 @@ private: cl::Event _kernel_event; bool _tune_new_kernels; }; - -class CLFileTuner : public CLTuner -{ -public: - /** Constructor - * - * @param[in] file_path File to load/store the tuning information from - * @param[in] update_file If true, save the new LWS table to the file on exit. - * @param[in] tune_new_kernels Find the optimal local workgroup size for kernels which are not present in the table ? - */ - CLFileTuner(std::string file_path = "acl_tuner.csv", bool update_file = false, bool tune_new_kernels = false); - - /** Save the content of the LWS table to file - */ - void save_to_file() const; - /* Setter for update_file option - * - * @param[in] update_file If true, save the new LWS table to the file on exit. - */ - void set_update_file(bool update_file); - /** Destructor - * - * Will save the LWS table to the file if the CLFileTuner was created with update_file enabled. - */ - ~CLFileTuner(); - const std::string filename; - -private: - bool _update_file; -}; } #endif /*__ARM_COMPUTE_CLTUNER_H__ */ diff --git a/src/graph/Graph.cpp b/src/graph/Graph.cpp index b1698e4672..2fe3a90aef 100644 --- a/src/graph/Graph.cpp +++ b/src/graph/Graph.cpp @@ -33,8 +33,19 @@ #include "arm_compute/runtime/Tensor.h" #include "support/ToolchainSupport.h" +#include <sys/stat.h> + using namespace arm_compute::graph; +namespace +{ +bool file_exists(const std::string &filename) +{ + std::ifstream file(filename); + return file.good(); +} + +} // namespace struct Stage { ITensorObject *_input; @@ -62,16 +73,20 @@ public: std::unique_ptr<INode> _current_node{ nullptr }; ITensorObject *_current_output{ nullptr }; bool _info_enabled{ false }; - CLFileTuner _tuner{}; + CLTuner _tuner{}; private: ITensorObject *_current_input{ nullptr }; GraphHints _previous_hints{}; }; +static const std::string tuner_data_filename = "acl_tuner.csv"; Graph::~Graph() //NOLINT { - //Can't use =default because the destructor must be defined after Graph::Private's definition + if(_pimpl->_tuner.tune_new_kernels() && !_pimpl->_tuner.lws_table().empty()) + { + _pimpl->_tuner.save_to_file(tuner_data_filename); + } } Graph::Graph() @@ -85,12 +100,14 @@ void Graph::graph_init(const bool use_cl_tuner) // Check if OpenCL is available and initialize the scheduler if(opencl_is_available()) { + if(_pimpl->_tuner.lws_table().empty() && file_exists(tuner_data_filename)) + { + _pimpl->_tuner.load_from_file(tuner_data_filename); + } _pimpl->_tuner.set_tune_new_kernels(use_cl_tuner); - _pimpl->_tuner.set_update_file(use_cl_tuner); arm_compute::CLScheduler::get().default_init(&_pimpl->_tuner); } } - void Graph::run() { while(true) diff --git a/src/runtime/CL/CLTuner.cpp b/src/runtime/CL/CLTuner.cpp index 29ddbea460..9fb9d1b396 100644 --- a/src/runtime/CL/CLTuner.cpp +++ b/src/runtime/CL/CLTuner.cpp @@ -99,6 +99,10 @@ void CLTuner::set_tune_new_kernels(bool tune_new_kernels) { _tune_new_kernels = tune_new_kernels; } +bool CLTuner::tune_new_kernels() const +{ + return _tune_new_kernels; +} void CLTuner::tune_kernel(ICLKernel &kernel) { @@ -246,15 +250,16 @@ void CLTuner::import_lws_table(const std::unordered_map<std::string, cl::NDRange _lws_table = lws_table; } -const std::unordered_map<std::string, cl::NDRange> &CLTuner::export_lws_table() const +const std::unordered_map<std::string, cl::NDRange> &CLTuner::lws_table() const { return _lws_table; } -CLFileTuner::CLFileTuner(std::string file_path, bool update_file, bool tune_new_kernels) - : CLTuner(tune_new_kernels), filename(std::move(file_path)), _update_file(update_file) +void CLTuner::load_from_file(const std::string &filename) { - std::ifstream fs(filename); + std::ifstream fs; + fs.exceptions(std::ifstream::failbit | std::ifstream::badbit); + fs.open(filename, std::ios::in); if(fs.is_open()) { std::string line; @@ -264,24 +269,18 @@ CLFileTuner::CLFileTuner(std::string file_path, bool update_file, bool tune_new_ std::string token; if(!std::getline(ss, token, ';')) { - continue; + ARM_COMPUTE_ERROR("Malformed row '%s' (Should be of the form 'kernel_id;lws[0];lws[1];lws[2]')", ss.str().c_str()); } std::string kernel_id = token; cl::NDRange lws(1, 1, 1); - bool lws_is_valid = true; for(int i = 0; i < 3; i++) { if(std::getline(ss, token, ';').fail()) { - lws_is_valid = false; - break; + ARM_COMPUTE_ERROR("Malformed row '%s' (Should be of the form 'kernel_id;lws[0];lws[1];lws[2]')", ss.str().c_str()); } lws.get()[i] = support::cpp11::stoi(token); } - if(!lws_is_valid) - { - continue; // Skip this kernel - } // If all dimensions are 0: reset to NullRange (i.e nullptr) if(lws[0] == 0 && lws[1] == 0 && lws[2] == 0) @@ -294,26 +293,14 @@ CLFileTuner::CLFileTuner(std::string file_path, bool update_file, bool tune_new_ } } -void CLFileTuner::save_to_file() const +void CLTuner::save_to_file(const std::string &filename) const { - const std::unordered_map<std::string, cl::NDRange> &table = export_lws_table(); - std::ofstream fs(filename); - for(auto kernel_data : table) + std::ofstream fs; + fs.exceptions(std::ifstream::failbit | std::ifstream::badbit); + fs.open(filename, std::ios::out); + for(auto kernel_data : _lws_table) { fs << kernel_data.first << ";" << kernel_data.second[0] << ";" << kernel_data.second[1] << ";" << kernel_data.second[2] << std::endl; } fs.close(); } - -CLFileTuner::~CLFileTuner() -{ - if(_update_file) - { - save_to_file(); - } -} - -void CLFileTuner::set_update_file(bool update_file) -{ - _update_file = update_file; -} |