"...composable_kernel_rocm.git" did not exist on "a24c5694ed30be912df38735d9b842bb4749e2aa"
program.cpp 25.8 KB
Newer Older
Paul's avatar
Paul committed
1
2
3
#include <migraphx/program.hpp>
#include <migraphx/stringutils.hpp>
#include <migraphx/instruction.hpp>
4
#include <migraphx/op/identity.hpp>
Paul's avatar
Paul committed
5
#include <migraphx/target.hpp>
Paul's avatar
Paul committed
6
7
8
#include <migraphx/env.hpp>
#include <migraphx/ranges.hpp>
#include <migraphx/time.hpp>
9
#include <migraphx/pass_manager.hpp>
10
#include <migraphx/register_target.hpp>
Shucai Xiao's avatar
Shucai Xiao committed
11
#include <migraphx/iterator_for.hpp>
12
#include <migraphx/iterator.hpp>
13
#include <migraphx/algorithm.hpp>
14
#include <migraphx/output_iterator.hpp>
Shucai Xiao's avatar
Shucai Xiao committed
15
#include <migraphx/make_op.hpp>
Paul's avatar
Paul committed
16
#include <iostream>
Paul's avatar
Paul committed
17
#include <sstream>
Paul's avatar
Paul committed
18
#include <algorithm>
19
#include <set>
Paul's avatar
Paul committed
20
#include <utility>
21

22
#include <unordered_set>
Shucai Xiao's avatar
Shucai Xiao committed
23
24
#include <map>
#include <cassert>
Paul's avatar
Paul committed
25

Paul's avatar
Paul committed
26
namespace migraphx {
Paul's avatar
Paul committed
27
inline namespace MIGRAPHX_INLINE_NS {
Paul's avatar
Paul committed
28

29
30
using milliseconds = std::chrono::duration<double, std::milli>;

Paul's avatar
Paul committed
31
32
struct program_impl
{
Shucai Xiao's avatar
Shucai Xiao committed
33
    // A map is used to keep references to modules of the program
34
    std::unordered_map<std::string, module> modules;
Paul's avatar
Paul committed
35
    context ctx;
36
    std::string target_name;
Paul's avatar
Paul committed
37
38
};

39
program::program() : impl(std::make_unique<program_impl>()) { this->create_module("main"); }
Paul's avatar
Paul committed
40

Paul's avatar
Paul committed
41
program::program(program&&) noexcept = default;
Shucai Xiao's avatar
Shucai Xiao committed
42
program::~program() noexcept         = default;
Paul's avatar
Paul committed
43

44
// copy constructor
Shucai Xiao's avatar
Shucai Xiao committed
45
program::program(const program& p) { assign(p); }
46
47

// copy assignment operator
Shucai Xiao's avatar
Shucai Xiao committed
48
program& program::operator=(program p)
49
{
Shucai Xiao's avatar
Shucai Xiao committed
50
    std::swap(p.impl, this->impl);
51
52
53
    return *this;
}

Shucai Xiao's avatar
Shucai Xiao committed
54
void program::assign(const program& p)
55
{
Shucai Xiao's avatar
Shucai Xiao committed
56
    if(!impl)
57
58
59
    {
        impl = std::make_unique<program_impl>();
    }
Shucai Xiao's avatar
Shucai Xiao committed
60
    else if(!impl->modules.empty())
61
    {
Shucai Xiao's avatar
Shucai Xiao committed
62
        impl->modules.clear();
63
    }
Shucai Xiao's avatar
Shucai Xiao committed
64

65
    impl->ctx         = p.impl->ctx;
Shucai Xiao's avatar
Shucai Xiao committed
66
67
    impl->target_name = p.impl->target_name;
    impl->modules     = p.impl->modules;
Shucai Xiao's avatar
Shucai Xiao committed
68
69
70
71

    // build a map from old ins to new ins
    // Build a map from old module to new module
    std::unordered_map<module_ref, module_ref> mod_map;
Paul's avatar
Paul committed
72
73
74
75
76
    std::transform(
        impl->modules.begin(),
        impl->modules.end(),
        std::inserter(mod_map, mod_map.begin()),
        [&](auto&& xp) { return std::make_pair(&p.impl->modules.at(xp.first), &xp.second); });
Shucai Xiao's avatar
Shucai Xiao committed
77
78
79
80
81
82
83
84
85
86
87
88
89
90
91
92

    std::unordered_map<instruction_ref, instruction_ref> ins_map;
    for(auto&& pp : mod_map)
    {
        auto old_ins = iterator_for(*pp.first);
        auto new_ins = iterator_for(*pp.second);
        std::transform(old_ins.begin(),
                       old_ins.end(),
                       new_ins.begin(),
                       std::inserter(ins_map, ins_map.begin()),
                       [](auto x, auto y) { return std::make_pair(x, y); });
    }

    // Update all references from all modules
    for(auto&& mp : impl->modules)
    {
93
        for(auto ins : iterator_for(mp.second))
Shucai Xiao's avatar
Shucai Xiao committed
94
95
            instruction::replace_refs(ins, ins_map, mod_map);
    }
96
97
}

Paul's avatar
Paul committed
98
shape program::get_parameter_shape(std::string name) const
Paul's avatar
Paul committed
99
{
Shucai Xiao's avatar
Shucai Xiao committed
100
101
    const auto* mm = this->get_main_module();
    return mm->get_parameter_shape(std::move(name));
Paul's avatar
Paul committed
102
103
}

104
105
std::vector<std::string> program::get_parameter_names() const
{
Shucai Xiao's avatar
Shucai Xiao committed
106
107
    const auto* mm = this->get_main_module();
    return mm->get_parameter_names();
108
109
}

mei-ye's avatar
mei-ye committed
110
111
instruction_ref program::get_parameter(std::string name) const
{
Shucai Xiao's avatar
Shucai Xiao committed
112
113
    const auto* mm = this->get_main_module();
    return mm->get_parameter(std::move(name));
mei-ye's avatar
mei-ye committed
114
115
}

Paul's avatar
Paul committed
116
117
std::unordered_map<std::string, shape> program::get_parameter_shapes() const
{
Shucai Xiao's avatar
Shucai Xiao committed
118
119
    const auto* mm = this->get_main_module();
    return mm->get_parameter_shapes();
Paul's avatar
Paul committed
120
121
}

Shucai Xiao's avatar
Shucai Xiao committed
122
std::size_t program::size() const { return impl->modules.size(); }
123

124
125
std::vector<shape> program::get_output_shapes() const
{
Shucai Xiao's avatar
Shucai Xiao committed
126
127
    const auto* mm = this->get_main_module();
    return mm->get_output_shapes();
128
}
Paul's avatar
Paul committed
129

Paul's avatar
Paul committed
130
131
context& program::get_context() const { return impl->ctx; }

Paul's avatar
Paul committed
132
133
instruction_ref program::validate() const
{
Shucai Xiao's avatar
Shucai Xiao committed
134
135
    const auto* mm = this->get_main_module();
    return mm->validate();
Paul's avatar
Paul committed
136
137
}

138
139
bool program::is_compiled() const { return not this->impl->target_name.empty(); }

140
void program::compile(const target& t, compile_options options)
Paul's avatar
Paul committed
141
{
142
143
144
    assert(not this->is_compiled());
    this->impl->target_name = t.name();
    this->impl->ctx         = t.get_context();
Paul's avatar
Paul committed
145
    if(enabled(MIGRAPHX_TRACE_COMPILE{}))
146
        options.trace = tracer{std::cout};
Shucai Xiao's avatar
Shucai Xiao committed
147

148
149
    options.trace(*this);
    options.trace();
Shucai Xiao's avatar
Shucai Xiao committed
150

Shucai Xiao's avatar
Shucai Xiao committed
151
    auto&& passes = t.get_passes(this->impl->ctx, options);
152
153
154
    run_passes(*this, passes, options.trace);

    auto mods = this->get_modules();
Shucai Xiao's avatar
Shucai Xiao committed
155

156
157
    // Validate and finalize
    for(const auto& mod : reverse(mods))
Paul's avatar
Paul committed
158
    {
Shucai Xiao's avatar
Shucai Xiao committed
159
160
161
162
163
164
        auto invalid = mod->validate();
        if(invalid != mod->end())
        {
            MIGRAPHX_THROW("Invalid module " + mod->name() + " from compilation at instruction " +
                           std::to_string(std::distance(mod->begin(), invalid)));
        }
165
166
167
168
169
170
171
        auto dangling = mod->find_dangling_reference();
        if(dangling != mod->end())
        {
            auto index = std::distance(mod->begin(), dangling);
            MIGRAPHX_THROW("Dangling reference in module " + mod->name() + " from instruction " +
                           std::to_string(index));
        }
Shucai Xiao's avatar
Shucai Xiao committed
172
        mod->finalize(this->impl->ctx);
Paul's avatar
Paul committed
173
    }
Paul's avatar
Paul committed
174
175
176
177
}

void program::finalize()
{
Shucai Xiao's avatar
Shucai Xiao committed
178
179
    auto* mm = this->get_main_module();
    mm->finalize(this->impl->ctx);
Paul's avatar
Paul committed
180
181
}

Paul's avatar
Paul committed
182
template <class F>
Shucai Xiao's avatar
Shucai Xiao committed
183
std::vector<argument> generic_eval(const module* mod,
184
185
                                   context& ctx,
                                   std::unordered_map<std::string, argument> params,
Shucai Xiao's avatar
Shucai Xiao committed
186
                                   std::unordered_map<instruction_ref, argument> results,
187
                                   F trace)
Paul's avatar
Paul committed
188
{
Shucai Xiao's avatar
Shucai Xiao committed
189
190
    assert(mod->validate() == mod->end());
    results.reserve(mod->size() * 2);
Paul's avatar
Paul committed
191
192
    std::vector<argument> values;
    values.reserve(16);
Shucai Xiao's avatar
Shucai Xiao committed
193
    for(auto ins : iterator_for(*mod))
Paul's avatar
Paul committed
194
    {
195
196
        const auto& name = ins->name();
        if(name == "@literal")
Paul's avatar
Paul committed
197
        {
Paul's avatar
Paul committed
198
            results.emplace(ins, trace(ins, [&] { return ins->get_literal().get_argument(); }));
Paul's avatar
Paul committed
199
        }
200
        else if(name == "@param")
Paul's avatar
Paul committed
201
        {
Paul's avatar
Paul committed
202
203
204
205
206
            results.emplace(
                ins, trace(ins, [&] {
                    auto param_name = any_cast<builtin::param>(ins->get_operator()).parameter;
                    if(not contains(params, param_name))
                        MIGRAPHX_THROW("Parameter not found: " + param_name);
207
                    auto param = params[param_name];
Paul's avatar
Paul committed
208
209
210
211
212
                    if(param.get_shape() != ins->get_shape())
                        MIGRAPHX_THROW("Incorrect shape {" + to_string(param.get_shape()) +
                                       "} for parameter: " + param_name);
                    return param;
                }));
Paul's avatar
Paul committed
213
        }
214
        else if(name == "@outline")
Paul's avatar
Paul committed
215
        {
Paul's avatar
Paul committed
216
            results.emplace(ins, trace(ins, [&] { return argument{ins->get_shape(), nullptr}; }));
Paul's avatar
Paul committed
217
        }
218
219
220
221
222
223
224
225
226
227
228
229
230
        else if(name == "@return")
        {
            std::vector<argument> prog_outputs;
            std::transform(ins->inputs().begin(),
                           ins->inputs().end(),
                           std::back_inserter(prog_outputs),
                           [&](instruction_ref i) {
                               assert(results.find(i) != results.end());
                               return results[i];
                           });

            return prog_outputs;
        }
Paul's avatar
Paul committed
231
232
        else
        {
Paul's avatar
Paul committed
233
            values.resize(ins->inputs().size());
Paul's avatar
Paul committed
234
235
236
237
238
            std::transform(
                ins->inputs().begin(), ins->inputs().end(), values.begin(), [&](instruction_ref i) {
                    assert(results.find(i) != results.end());
                    return results[i];
                });
Shucai Xiao's avatar
Shucai Xiao committed
239
240
241
242
243
244
245

            const auto& mod_args = ins->module_inputs();
            auto module_eval     = [&](module_ref smod,
                                   const std::unordered_map<std::string, argument>& inputs) {
                return generic_eval(smod, ctx, inputs, results, trace);
            };

Shucai Xiao's avatar
Shucai Xiao committed
246
247
248
249
            results.emplace(ins, trace(ins, [&] {
                                return ins->normalized_operator().compute(
                                    ctx, ins->get_shape(), values, mod_args, module_eval);
                            }));
Paul's avatar
Paul committed
250
        }
251
        assert(results.find(ins) != results.end());
Paul's avatar
Paul committed
252
    }
Shucai Xiao's avatar
Shucai Xiao committed
253
    return {results.at(std::prev(mod->end()))};
Paul's avatar
Paul committed
254
255
}

Shucai Xiao's avatar
Shucai Xiao committed
256
257
258
259
260
261
template <class F>
std::vector<argument> generic_eval(const program& p,
                                   context& ctx,
                                   std::unordered_map<std::string, argument> params,
                                   F trace)
{
Shucai Xiao's avatar
Shucai Xiao committed
262
263
    const module* mm = p.get_main_module();
    return generic_eval(mm, ctx, params, {}, trace);
Shucai Xiao's avatar
Shucai Xiao committed
264
265
}

266
std::vector<argument> program::eval(parameter_map params) const
Paul's avatar
Paul committed
267
{
Paul's avatar
Paul committed
268
269
    auto& ctx = this->impl->ctx;
#ifndef NDEBUG
Paul's avatar
Paul committed
270
    auto sctx          = ctx;
Paul's avatar
Paul committed
271
272
273
    auto check_context = [&](auto f) {
        assert(is_shared(ctx, sctx));
        auto x = f();
Paul's avatar
Paul committed
274
        sctx   = ctx;
Paul's avatar
Paul committed
275
276
277
        return x;
    };
#else
Paul's avatar
Paul committed
278
    auto check_context = [](auto f) { return f(); };
Paul's avatar
Paul committed
279
#endif
Paul's avatar
Paul committed
280
281
282
283

    auto trace_level = value_of(MIGRAPHX_TRACE_EVAL{});

    if(trace_level > 0)
Paul's avatar
Paul committed
284
    {
Paul's avatar
Paul committed
285
        return generic_eval(*this, ctx, std::move(params), [&](auto& ins, auto f) {
Paul's avatar
Paul committed
286
            ctx.finish();
Paul's avatar
Paul committed
287
288
            std::cout << "Run instruction: ";
            this->debug_print(ins);
289
            timer t{};
Paul's avatar
Paul committed
290
            auto result = check_context(f);
291
            double t1   = t.record<milliseconds>();
Paul's avatar
Paul committed
292
            ctx.finish();
293
294
            double t2 = t.record<milliseconds>();
            std::cout << "Time: " << t1 << "ms, " << t2 << "ms" << std::endl;
Paul's avatar
Paul committed
295
            if(trace_level > 1 and ins->name().front() != '@' and ins->name() != "load")
296
                std::cout << "Output: " << result << std::endl;
Paul's avatar
Paul committed
297
            return result;
Paul's avatar
Paul committed
298
        });
Paul's avatar
Paul committed
299
300
301
302
    }
    else
    {
        return generic_eval(
Paul's avatar
Paul committed
303
            *this, ctx, std::move(params), [&](auto&, auto f) { return check_context(f); });
Paul's avatar
Paul committed
304
    }
Paul's avatar
Paul committed
305
306
}

307
const int program_file_version = 5;
308
309
310
311
312
313
314
315

value program::to_value() const
{
    value result;
    result["version"] = program_file_version;
    result["target"]  = this->impl->target_name;
    if(not this->impl->target_name.empty())
        result["context"] = this->impl->ctx.to_value();
Shucai Xiao's avatar
Shucai Xiao committed
316

317
    value module_vals = value::object{};
Shucai Xiao's avatar
Shucai Xiao committed
318
    std::unordered_map<instruction_ref, std::string> names;
319
    for(auto& mod : this->get_modules())
Shucai Xiao's avatar
Shucai Xiao committed
320
    {
Shucai Xiao's avatar
Shucai Xiao committed
321
322
        value mod_val;
        value nodes;
323
324
        mod_val["name"] = mod->name();
        names           = mod->print(
Shucai Xiao's avatar
Shucai Xiao committed
325
326
327
328
329
330
331
332
333
334
335
336
337
338
339
340
341
342
343
344
345
346
347
348
349
350
351
352
353
354
355
356
357
358
            [&](auto ins, auto ins_names) {
                value node;
                node["output"]     = ins_names.at(ins);
                node["name"]       = ins->name();
                node["shape"]      = migraphx::to_value(ins->get_shape());
                node["normalized"] = ins->is_normalized();
                if(ins->name() == "@literal")
                    node["literal"] = migraphx::to_value(ins->get_literal());
                node["operator"] = ins->get_operator().to_value();
                std::vector<std::string> inputs;
                std::transform(ins->inputs().begin(),
                               ins->inputs().end(),
                               std::back_inserter(inputs),
                               [&](auto i) {
                                   assert(contains(ins_names, i));
                                   return ins_names.at(i);
                               });
                node["inputs"]   = inputs;
                auto module_args = ins->module_inputs();
                if(not module_args.empty())
                {
                    std::vector<std::string> module_inputs;
                    std::transform(module_args.begin(),
                                   module_args.end(),
                                   std::back_inserter(module_inputs),
                                   [&](auto mod_ref) { return mod_ref->name(); });
                    node["module_inputs"] = module_inputs;
                }

                nodes.push_back(node);
            },
            names);
        mod_val["nodes"] = nodes;

359
        module_vals[mod->name()] = mod_val;
Shucai Xiao's avatar
Shucai Xiao committed
360
    }
Shucai Xiao's avatar
Shucai Xiao committed
361
362
363

    result["modules"] = module_vals;

364
365
    return result;
}
Shucai Xiao's avatar
Shucai Xiao committed
366

Shucai Xiao's avatar
Shucai Xiao committed
367
368
369
370
371
static void mod_from_val(module_ref mod,
                         const value& v,
                         std::unordered_map<std::string, instruction_ref>& instructions,
                         const std::unordered_map<std::string, module_ref>& map_mods)
{
372
    const auto& module_val = v.at(mod->name());
Shucai Xiao's avatar
Shucai Xiao committed
373
374
375
376
377
378
379
380
381
382
383
384
385
386
387
388
389
390
391
392
393
394
395
396
397
398
399
400
401
402
403
404
405
406
407
408
409
410
411
412
413
414
415
416
417
418
419
420
421
422
423
424
425
426
427
428
429
430
431
432
433
    for(const value& node : module_val.at("nodes"))
    {
        instruction_ref output;
        auto name       = node.at("name").to<std::string>();
        auto fields     = node.at("operator");
        auto normalized = node.at("normalized").to<bool>();

        if(name == "@param")
        {
            output = mod->add_parameter(fields["parameter"].to<std::string>(),
                                        migraphx::from_value<shape>(node.at("shape")));
        }
        else if(name == "@literal")
        {
            output = mod->add_literal(migraphx::from_value<literal>(node.at("literal")));
        }
        else
        {
            auto op = make_op(name, fields);
            std::vector<instruction_ref> inputs;
            std::transform(node.at("inputs").begin(),
                           node.at("inputs").end(),
                           std::back_inserter(inputs),
                           [&](const value& i) {
                               auto i_name = i.to<std::string>();
                               assert(contains(instructions, i_name));
                               return instructions.at(i_name);
                           });

            std::vector<module_ref> module_inputs;
            if(node.contains("module_inputs"))
            {
                std::transform(node.at("module_inputs").begin(),
                               node.at("module_inputs").end(),
                               std::back_inserter(module_inputs),
                               [&](const value& i) { return map_mods.at(i.to<std::string>()); });

                for(auto& smod : module_inputs)
                {
                    mod_from_val(smod, v, instructions, map_mods);
                }
            }

            if(name == "@return")
            {
                output = mod->add_return(inputs);
            }
            else if(module_inputs.empty())
            {
                output = mod->add_instruction(op, inputs);
            }
            else
            {
                output = mod->add_instruction(op, inputs, module_inputs);
            }
        }
        output->set_normalized(normalized);
        instructions[node.at("output").to<std::string>()] = output;
    }
}

434
435
436
437
void program::from_value(const value& v)
{
    auto version = v.at("version").to<int>();
    if(version != program_file_version)
Shucai Xiao's avatar
Shucai Xiao committed
438
439
440
441
    {
        MIGRAPHX_THROW("Warning: Program version mismatch");
    }

442
443
444
445
446
447
448
449
    this->impl->target_name = v.at("target").to<std::string>();
    if(not this->impl->target_name.empty())
    {
        target t        = make_target(this->impl->target_name);
        this->impl->ctx = t.get_context();
        this->impl->ctx.from_value(v.at("context"));
    }

Shucai Xiao's avatar
Shucai Xiao committed
450
451
    auto module_vals = v.at("modules");
    for(const auto& vv : module_vals)
452
    {
453
        const auto& name = vv.get_key();
Shucai Xiao's avatar
Shucai Xiao committed
454
455
        if(name == "main")
            continue;
456
        impl->modules.emplace(name, name);
457
    }
458
    std::unordered_map<std::string, module_ref> map_mods;
Paul's avatar
Paul committed
459
460
461
462
    std::transform(impl->modules.begin(),
                   impl->modules.end(),
                   std::inserter(map_mods, map_mods.end()),
                   [&](auto&& pp) { return std::make_pair(pp.first, &pp.second); });
Shucai Xiao's avatar
Shucai Xiao committed
463
464
465
466
467

    std::unordered_map<std::string, instruction_ref> map_insts;
    auto* mm = get_main_module();
    mod_from_val(mm, module_vals, map_insts, map_mods);

468
469
470
    this->finalize();
}

Paul's avatar
Paul committed
471
472
473
double common_average(const std::vector<double>& v)
{
    std::size_t n = v.size() / 4;
Paul's avatar
Paul committed
474
475
    double total  = std::accumulate(v.begin() + n, v.end() - n, 0.0);
    return total / std::distance(v.begin() + n, v.end() - n);
Paul's avatar
Paul committed
476
477
}

Paul Fultz II's avatar
Paul Fultz II committed
478
479
480
481
482
483
484
485
std::string perf_group(const operation& op)
{
    auto attr = op.attributes();
    if(attr.contains("group"))
        return attr.at("group").to<std::string>();
    return op.name();
}

Paul's avatar
Paul committed
486
487
void program::perf_report(std::ostream& os, std::size_t n, parameter_map params) const
{
488
    auto& ctx = this->impl->ctx;
Paul's avatar
Paul committed
489
490
    // Run once by itself
    eval(params);
Paul's avatar
Paul committed
491
    ctx.finish();
Paul's avatar
Paul committed
492
    // Run and time entire program
Paul's avatar
Paul committed
493
494
    std::vector<double> total_vec;
    total_vec.reserve(n);
Paul's avatar
Paul committed
495
    for(std::size_t i = 0; i < n; i++)
Paul's avatar
Paul committed
496
    {
Paul's avatar
Paul committed
497
498
499
500
        total_vec.push_back(time<milliseconds>([&] {
            eval(params);
            ctx.finish();
        }));
Paul's avatar
Paul committed
501
    }
Paul's avatar
Paul committed
502
503
    std::sort(total_vec.begin(), total_vec.end());
    std::unordered_map<instruction_ref, std::vector<double>> ins_vec;
Paul's avatar
Paul committed
504
    // Fill the map
Paul's avatar
Paul committed
505
    generic_eval(*this, ctx, params, [&](auto ins, auto) {
Paul's avatar
Paul committed
506
        ins_vec[ins].reserve(n);
Paul's avatar
Paul committed
507
508
        return argument{};
    });
Paul's avatar
Paul committed
509
    // Run and time each instruction
Paul's avatar
Paul committed
510
    for(std::size_t i = 0; i < n; i++)
Paul's avatar
Paul committed
511
    {
Paul's avatar
Paul committed
512
        generic_eval(*this, ctx, params, [&](auto ins, auto f) {
513
            argument result;
Paul's avatar
Paul committed
514
515
516
517
            ins_vec[ins].push_back(time<milliseconds>([&] {
                result = f();
                ctx.finish();
            }));
518
            return result;
Paul's avatar
Paul committed
519
520
        });
    }
Paul's avatar
Paul committed
521
522
    for(auto&& p : ins_vec)
        std::sort(p.second.begin(), p.second.end());
Paul's avatar
Paul committed
523
    // Run and time implicit overhead
Paul's avatar
Paul committed
524
525
    std::vector<double> overhead_vec;
    overhead_vec.reserve(n);
Paul's avatar
Paul committed
526
    for(std::size_t i = 0; i < n; i++)
Paul's avatar
Paul committed
527
    {
Paul's avatar
Paul committed
528
        overhead_vec.push_back(time<milliseconds>([&] { dry_run(params); }));
Paul's avatar
Paul committed
529
530
    }

Paul's avatar
Paul committed
531
    double total_time             = common_average(total_vec);
Paul's avatar
Paul committed
532
    double rate                   = 1000.0 / total_time;
Paul's avatar
Paul committed
533
    double overhead_time          = common_average(overhead_vec);
Paul's avatar
Paul committed
534
    double overhead_percent       = overhead_time * 100.0 / total_time;
Paul's avatar
Paul committed
535
    double total_instruction_time = 0.0;
Paul's avatar
Paul committed
536
    std::unordered_map<std::string, double> op_times;
Paul's avatar
Paul committed
537
    for(auto&& p : ins_vec)
Paul's avatar
Paul committed
538
539
    {
        double avg = common_average(p.second);
Paul Fultz II's avatar
Paul Fultz II committed
540
        op_times[perf_group(p.first->get_operator())] += avg;
Paul's avatar
Paul committed
541
542
        total_instruction_time += avg;
    }
Paul's avatar
Paul committed
543
544
    double calculate_overhead_time    = total_time - total_instruction_time;
    double calculate_overhead_percent = calculate_overhead_time * 100.0 / total_time;
Paul's avatar
Paul committed
545

Shucai Xiao's avatar
Shucai Xiao committed
546
547
548
    std::unordered_map<instruction_ref, std::string> names;
    this->print(names, [&](auto ins, auto ins_names) {
        instruction::print(std::cout, ins, ins_names);
549
550
551
552
553

        // skip return instruction
        if(ins->name() == "@return")
            return;

Paul's avatar
Paul committed
554
555
556
        double avg     = common_average(ins_vec[ins]);
        double percent = std::ceil(100.0 * avg / total_instruction_time);
        os << ": " << avg << "ms, " << percent << "%";
557
        os << std::endl;
Paul's avatar
Paul committed
558
    });
Paul's avatar
Paul committed
559
560
561

    os << std::endl;
    os << "Summary:" << std::endl;
562
563
564
565
566
567
568
    std::vector<std::pair<double, std::string>> op_times_sorted;
    std::transform(op_times.begin(),
                   op_times.end(),
                   std::back_inserter(op_times_sorted),
                   [](auto p) { return std::make_pair(p.second, p.first); });
    std::sort(op_times_sorted.begin(), op_times_sorted.end(), std::greater<>{});
    for(auto&& p : op_times_sorted)
Paul's avatar
Paul committed
569
    {
570
571
        auto&& name    = p.second;
        double avg     = p.first;
Paul's avatar
Paul committed
572
573
574
575
576
        double percent = std::ceil(100.0 * avg / total_instruction_time);
        os << name << ": " << avg << "ms, " << percent << "%" << std::endl;
    }

    os << std::endl;
Paul's avatar
Paul committed
577

Paul's avatar
Paul committed
578
    os << "Rate: " << rate << "/sec" << std::endl;
Paul's avatar
Paul committed
579
580
    os << "Total time: " << total_time << "ms" << std::endl;
    os << "Total instructions time: " << total_instruction_time << "ms" << std::endl;
Paul's avatar
Paul committed
581
582
583
584
    os << "Overhead time: " << overhead_time << "ms"
       << ", " << calculate_overhead_time << "ms" << std::endl;
    os << "Overhead: " << std::round(overhead_percent) << "%"
       << ", " << std::round(calculate_overhead_percent) << "%" << std::endl;
Paul's avatar
Paul committed
585
586
}

Paul's avatar
Paul committed
587
588
void program::debug_print() const { std::cout << *this << std::endl; }
void program::debug_print(instruction_ref ins) const
Paul's avatar
Paul committed
589
{
Shucai Xiao's avatar
Shucai Xiao committed
590
    std::unordered_map<instruction_ref, std::string> names;
591
    if(std::any_of(this->impl->modules.begin(), this->impl->modules.end(), [&](const auto& pp) {
592
           return is_end(pp.second.end(), ins);
Shucai Xiao's avatar
Shucai Xiao committed
593
       }))
Paul's avatar
Paul committed
594
595
596
597
    {
        std::cout << "End instruction" << std::endl;
        return;
    }
598
599
    else if(std::none_of(this->impl->modules.begin(),
                         this->impl->modules.end(),
600
                         [&](const auto& pp) { return pp.second.has_instruction(ins); }))
Paul's avatar
Paul committed
601
602
603
604
    {
        std::cout << "Instruction not part of program" << std::endl;
        return;
    }
Shucai Xiao's avatar
Shucai Xiao committed
605

Paul's avatar
Paul committed
606
    std::stringstream ss;
Shucai Xiao's avatar
Shucai Xiao committed
607
    this->print(names, [&](auto x, auto ins_names) {
Paul's avatar
Paul committed
608
        if(x == ins)
Paul's avatar
Paul committed
609
        {
Shucai Xiao's avatar
Shucai Xiao committed
610
            instruction::print(std::cout, x, ins_names);
Paul's avatar
Paul committed
611
612
613
614
615
            std::cout << std::endl;
        }
    });
}

Shucai Xiao's avatar
Shucai Xiao committed
616
617
618
619
void program::print(
    std::unordered_map<instruction_ref, std::string>& names,
    const std::function<void(instruction_ref, std::unordered_map<instruction_ref, std::string>)>&
        print_func) const
620
{
621
    for(const auto& pp : this->impl->modules)
622
    {
623
        names = pp.second.print(print_func, names);
624
625
626
    }
}

Shucai Xiao's avatar
Shucai Xiao committed
627
void program::print_graph(std::ostream& os, bool brief) const
628
{
Shucai Xiao's avatar
Shucai Xiao committed
629
630
    const auto* mm = this->get_main_module();
    mm->print_graph(os, brief);
631
632
633
634
}

void program::print_cpp(std::ostream& os) const
{
Shucai Xiao's avatar
Shucai Xiao committed
635
636
637
638
639
640
641
642
    auto vec_modules = this->get_modules();
    std::unordered_map<instruction_ref, std::string> names;
    for(auto& mod : vec_modules)
    {
        os << "module: \"" << mod->name() << "\"" << std::endl;
        names = mod->print_cpp(os, names);
        os << std::endl;
    }
643
644
}

Paul's avatar
Paul committed
645
646
void program::dry_run(std::unordered_map<std::string, argument> params) const
{
Paul's avatar
Paul committed
647
    auto& ctx = this->impl->ctx;
Paul's avatar
Paul committed
648
    generic_eval(*this, ctx, std::move(params), [](auto&&...) { return argument{}; });
Paul's avatar
Paul committed
649
650
}

Shucai Xiao's avatar
Shucai Xiao committed
651
void program::annotate(std::ostream& os, const std::function<void(instruction_ref)>& a) const
Paul's avatar
Paul committed
652
{
653
    for(auto& pp : this->impl->modules)
Shucai Xiao's avatar
Shucai Xiao committed
654
    {
655
656
        std::cout << pp.first << ":" << std::endl;
        pp.second.annotate(os, a);
Shucai Xiao's avatar
Shucai Xiao committed
657
    }
Paul's avatar
Paul committed
658
659
}

Paul's avatar
Paul committed
660
const module* program::get_module(const std::string& name) const { return &impl->modules.at(name); }
Shucai Xiao's avatar
Shucai Xiao committed
661
662
663

module* program::create_module(const std::string& name)
{
664
    assert(not contains(impl->modules, name));
665
666
    auto r = impl->modules.emplace(name, name);
    return &(r.first->second);
Shucai Xiao's avatar
Shucai Xiao committed
667
668
}

Paul's avatar
Paul committed
669
module* program::get_module(const std::string& name) { return &impl->modules.at(name); }
Shucai Xiao's avatar
Shucai Xiao committed
670
671
672
673
674

module* program::get_main_module() { return get_module("main"); }

const module* program::get_main_module() const { return get_module("main"); }

Paul's avatar
Paul committed
675
template <class T>
676
std::vector<T*> generic_get_modules(T* mm)
Shucai Xiao's avatar
Shucai Xiao committed
677
{
678
    std::vector<T*> vec_modules;
Shucai Xiao's avatar
Shucai Xiao committed
679
680
681
682
683
    vec_modules.push_back(mm);
    auto sub_modules = mm->get_sub_modules();
    vec_modules.insert(vec_modules.end(), sub_modules.begin(), sub_modules.end());
    return vec_modules;
}
Shucai Xiao's avatar
Shucai Xiao committed
684

Paul's avatar
Paul committed
685
template <class Map, class T, class OutputIterator>
686
void generic_get_unused_modules(Map& m, const std::vector<T*>& mods, OutputIterator out)
Shucai Xiao's avatar
Shucai Xiao committed
687
{
688
689
690
691
    std::unordered_set<std::string> used;
    std::transform(mods.begin(), mods.end(), std::inserter(used, used.end()), [](auto&& mod) {
        return mod->name();
    });
Paul's avatar
Paul committed
692
693
694
695
696
    transform_if(m.begin(),
                 m.end(),
                 out,
                 [&](auto&& pp) { return not contains(used, pp.first); },
                 [](auto&& pp) { return &pp.second; });
697
}
Shucai Xiao's avatar
Shucai Xiao committed
698

699
700
701
702
703
704
705
706
707
708
709
710
std::vector<const module*> program::get_modules() const
{
    auto result = generic_get_modules(this->get_main_module());
    generic_get_unused_modules(impl->modules, result, std::back_inserter(result));
    return result;
}

std::vector<module*> program::get_modules()
{
    auto result = generic_get_modules(this->get_main_module());
    generic_get_unused_modules(impl->modules, result, std::back_inserter(result));
    return result;
Shucai Xiao's avatar
Shucai Xiao committed
711
712
}

713
714
715
716
717
718
719
720
721
722
723
724
725
726
727
728
729
730
731
732
733
734
735
736
737
738
739
740
741
742
743
744
745
746
747
748
749
750
751
752
753
754
755
756
757
758
759
template <class Map, class T>
bool is_unused_module(Map& m, const std::vector<T*>& mods, const std::string& name)
{
    bool is_unused = false;
    generic_get_unused_modules(m, mods, make_function_output_iterator([&](auto* mod) {
                                   if(mod->name() == name)
                                       is_unused = true;
                               }));
    return is_unused;
}

template <class Map>
bool references_instruction(Map& m, const instruction& ins, const std::string& name)
{
    return std::any_of(m.begin(), m.end(), [&](auto&& p) {
        if(p.first == name)
            return false;
        return std::any_of(p.second.begin(), p.second.end(), [&](auto&& i) {
            return std::any_of(i.inputs().begin(), i.inputs().end(), [&](auto&& j) {
                return std::addressof(*j) == std::addressof(ins);
            });
        });
    });
}

void program::remove_module(const std::string& name)
{
    // cppcheck-suppress assertWithSideEffect
    assert(is_unused_module(impl->modules, generic_get_modules(this->get_main_module()), name) &&
           "Module used in program");
    assert(std::none_of(
               impl->modules.at(name).begin(),
               impl->modules.at(name).end(),
               [&](auto&& ins) { return references_instruction(impl->modules, ins, name); }) &&
           "Instruction referenced in another module");
    impl->modules.erase(name);
}

void program::remove_unused_modules()
{
    std::vector<module*> unused;
    generic_get_unused_modules(
        impl->modules, generic_get_modules(this->get_main_module()), std::back_inserter(unused));
    for(auto* m : unused)
        this->remove_module(m->name());
}

760
761
program& program::sort()
{
762
    for(auto& pp : this->impl->modules)
Shucai Xiao's avatar
Shucai Xiao committed
763
    {
764
        pp.second.sort();
Shucai Xiao's avatar
Shucai Xiao committed
765
766
    }

767
768
769
    return *this;
}

Paul's avatar
Paul committed
770
bool operator==(const program& x, const program& y) { return to_string(x) == to_string(y); }
Paul's avatar
Paul committed
771

Paul's avatar
Paul committed
772
std::ostream& operator<<(std::ostream& os, const program& p)
Paul's avatar
Paul committed
773
{
Shucai Xiao's avatar
Shucai Xiao committed
774
775
776
    auto vec_modules = p.get_modules();
    std::unordered_map<instruction_ref, std::string> names;
    for(auto& mod : vec_modules)
Shucai Xiao's avatar
Shucai Xiao committed
777
    {
Shucai Xiao's avatar
Shucai Xiao committed
778
779
780
781
782
783
784
        os << "module: \"" << mod->name() << "\"" << std::endl;
        names = mod->print(
            [&](auto ins, auto ins_names) {
                instruction::print(os, ins, ins_names);
                os << std::endl;
            },
            names);
785
        os << std::endl;
Shucai Xiao's avatar
Shucai Xiao committed
786
787
    }

Paul's avatar
Paul committed
788
    return os;
Paul's avatar
Paul committed
789
}
Paul's avatar
Paul committed
790

Paul's avatar
Paul committed
791
} // namespace MIGRAPHX_INLINE_NS
Paul's avatar
Paul committed
792
} // namespace migraphx