perfdb.cpp 2.92 KB
Newer Older
Paul's avatar
Paul committed
1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
31
32
33
34
35
36
37
38
39
40
41
42
43
44
45
46
47
48
49
50
51
52
53
54
55
56
57
58
59
60
61
62
63
64
65
66
67
68
69
70
71
72
73
74
75
76
77
78
79
80
81
82
83
84
85
86
87
88
89
90
91
92
93
94
#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