aboutsummaryrefslogtreecommitdiff
diff options
context:
space:
mode:
authorAnthony Barbier <anthony.barbier@arm.com>2018-02-28 13:47:58 +0000
committerAnthony Barbier <anthony.barbier@arm.com>2018-11-02 16:48:33 +0000
commit8b81195919eaab90a803fabae57ca05136e9430e (patch)
tree325798fd319f8091b18458127d3843b88c4d40af
parentfadfc6d6737716891f543f61a529e984da9b0a8b (diff)
downloadComputeLibrary-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.h53
-rw-r--r--src/graph/Graph.cpp25
-rw-r--r--src/runtime/CL/CLTuner.cpp45
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;
-}