Query to perfdb

src/include/migraphx/stringutils.hpp

@@ -168,27 +168,27 @@ inline std::string interpolate_string(const std::string& input,
 }
 
 template <class Iterator>
-inline std::string to_string_range(Iterator start, Iterator last)
+inline std::string to_string_range(Iterator start, Iterator last, const char* delim = ", ")
 {
     std::stringstream ss;
     if(start != last)
     {
         ss << *start;
-        std::for_each(std::next(start), last, [&](auto&& x) { ss << ", " << x; });
+        std::for_each(std::next(start), last, [&](auto&& x) { ss << delim << x; });
     }
     return ss.str();
 }
 
 template <class Range>
-inline std::string to_string_range(const Range& r)
+inline std::string to_string_range(const Range& r, const char* delim = ", ")
 {
-    return to_string_range(r.begin(), r.end());
+    return to_string_range(r.begin(), r.end(), delim);
 }
 
 template <class T>
-inline std::string to_string_range(const std::initializer_list<T>& r)
+inline std::string to_string_range(const std::initializer_list<T>& r, const char* delim = ", ")
 {
-    return to_string_range(r.begin(), r.end());
+    return to_string_range(r.begin(), r.end(), delim);
 }
 
 template <class T>

src/targets/gpu/CMakeLists.txt

@@ -184,6 +184,7 @@ add_library(migraphx_gpu
     pack_int8_args.cpp
     prefuse_ops.cpp
     pad.cpp
+    perfdb.cpp
     pooling.cpp
     quant_convolution.cpp
     reverse.cpp

src/targets/gpu/include/migraphx/gpu/perfdb.hpp

+#ifndef MIGRAPHX_GUARD_GPU_PERFDB_HPP
+#define MIGRAPHX_GUARD_GPU_PERFDB_HPP
+
+#include <migraphx/config.hpp>
+#include <migraphx/shape.hpp>
+#include <migraphx/operation.hpp>
+#include <string>
+#include <vector>
+
+namespace migraphx {
+inline namespace MIGRAPHX_INLINE_NS {
+namespace gpu {
+
+struct problem_params
+{
+    operation op;
+    std::vector<shape> inputs;
+    shape output;
+};
+
+std::string get_mlir_perf_for_conv(const problem_params& pp);
+
+} // namespace gpu
+} // namespace MIGRAPHX_INLINE_NS
+} // namespace migraphx
+#endif // MIGRAPHX_GUARD_GPU_PERFDB_HPP

src/targets/gpu/mlir.cpp

@@ -21,6 +21,7 @@
 #include <migraphx/gpu/context.hpp>
 #include <migraphx/gpu/device_name.hpp>
 #include <migraphx/iterator_for.hpp>
+#include <migraphx/gpu/perfdb.hpp>
 #include <deque>
 #include <variant>
 
@@ -464,6 +465,8 @@ struct mlir_program
             ops.add_attribute_value(get_operator_value(ins->get_operator()));
             if(ins->name() != "@return")
                 ops.add_results({get_shape(ins)});
+            if (ins->name())
+                pp = {ins->get_operator(), ins->inputs(), ins->get_shape()};
 
             std::vector<MlirValue> inputs;
             transform(
@@ -522,9 +525,15 @@ struct mlir_program
         MIGRAPHX_THROW("Failed to compile mlir program");
     }
 
+    std::string get_tune_params()
+    {
+        return get_mlir_perf_for_conv(pp);
+    }
+
     mlir_context ctx;
     MlirLocation location;
     mlir_module mmodule;
+    problem_params pp;
     std::deque<std::string> strings{};
 };
 

src/targets/gpu/perfdb.cpp

+#include <migraphx/gpu/perfdb.hpp>
+#include <migraphx/value.hpp>
+#include <migraphx/sqlite.hpp>
+#include <migraphx/stringutils.hpp>
+#include <migraphx/permutation.hpp>
+
+
+namespace migraphx {
+inline namespace MIGRAPHX_INLINE_NS {
+namespace gpu {
+
+namespace {
+
+// C-HxW-XxY-K-OHxOW-N-padhxpadw-stride-dilation-0-xlayout-wlayout-ylayout-type-F
+
+// case miopenFloat: return "FP32";
+// case miopenHalf: return "FP16";
+// case miopenInt8: return "INT8";
+// case miopenInt8x4: return "INT8x4";
+// case miopenInt32: return "INT32";
+// case miopenBFloat16: return "BF16";
+// case miopenDouble: return "FP64";
+
+std::string get_layout(const shape& s, std::string labels)
+{
+    auto result = labels;
+    auto p      = find_permutation(s);
+    std::transform(p.begin(), p.end(), result.begin(), [&](auto i) { return labels[i]; });
+    return result;
+}
+
+std::string get_type(const shape& s)
+{
+    static const std::unordered_map<shape::type_t, std::string> m = {
+        {shape::float_type, "FP32"},
+        {shape::half_type, "FP16"},
+        {shape::double_type, "FP64"},
+        {shape::int8_type, "INT8"},
+        {shape::int32_type, "INT32"},
+    };
+    auto it = m.find(s.type());
+    if (it == m.end())
+        return "UNKNOWN";
+    return it->second;
+}
+
+std::string generate_miopen_config(const problem_params& pp)
+{
+    value v = pp.op.to_value();
+    auto input = pp.inputs[0].lens();
+    auto weights = pp.inputs[1].lens();
+    auto output = pp.output.lens();
+    auto padding = v["padding"].to_vector<std::size_t>();
+    auto stride = v["stride"].to_vector<std::size_t>();
+    auto dilation = v["dilation"].to_vector<std::size_t>();
+    if (padding.size() != stride.size())
+        padding.erase(padding.begin()+padding.size()/2, padding.end());
+
+    return to_string_range({
+        to_string(input[1]),
+        to_string_range(input.begin()+2, input.end(), "x"),
+        to_string_range(weights.begin()+2, weights.end(), "x"),
+        to_string(weights[0]),
+        to_string_range(output.begin()+2, output.end(), "x"),
+        to_string(input[0]),
+        to_string_range(padding.begin()+2, padding.end(), "x"),
+        to_string_range(stride.begin()+2, stride.end(), "x"),
+        to_string_range(dilation.begin()+2, dilation.end(), "x"),
+        std::string{"0"},
+        get_layout(pp.inputs[0], "NCHW"),
+        get_layout(pp.inputs[1], "NCHW"),
+        get_layout(pp.output, "NCHW"),
+        get_type(pp.inputs[0]),
+        std::string{"F"}
+    }, "-");
+}
+
+}
+
+std::string get_mlir_perf_for_conv(const problem_params& pp)
+{
+    const auto dbpath = fs::path{"opt"}/"rocm"/"share"/"miopen"/"db"/"miopen.db";
+    auto db = sqlite::read(dbpath);
+    std::string query = "select * from perf_db where config=${config}";
+    auto results = db.execute(interpolate_string(query, {{"config", generate_miopen_config(pp)}}));
+    if (results.empty())
+        return "";
+    return results.front().at("params");
+}
+
+} // namespace gpu
+} // namespace MIGRAPHX_INLINE_NS
+} // namespace migraphx
+