lowering.cpp 18.1 KB
Newer Older
Shucai Xiao's avatar
Shucai Xiao committed
1
#include <iterator>
Paul's avatar
Paul committed
2
3
4
#include <migraphx/gpu/lowering.hpp>
#include <migraphx/manage_ptr.hpp>
#include <migraphx/instruction.hpp>
5
6
7
8
9
10
11
12
#include <migraphx/make_op.hpp>

#include <migraphx/op/abs.hpp>
#include <migraphx/op/batch_norm_inference.hpp>
#include <migraphx/op/convolution.hpp>
#include <migraphx/op/deconvolution.hpp>
#include <migraphx/op/dot.hpp>
#include <migraphx/op/elu.hpp>
Shucai Xiao's avatar
Shucai Xiao committed
13
#include <migraphx/op/if_op.hpp>
14
15
16
17
18
19
20
21
22
#include <migraphx/op/leaky_relu.hpp>
#include <migraphx/op/lrn.hpp>
#include <migraphx/op/pooling.hpp>
#include <migraphx/op/reshape.hpp>
#include <migraphx/op/quant_convolution.hpp>
#include <migraphx/op/quant_dot.hpp>

#include <migraphx/gpu/abs.hpp>
#include <migraphx/gpu/batch_norm_inference.hpp>
Paul's avatar
Paul committed
23
24
#include <migraphx/gpu/context.hpp>
#include <migraphx/gpu/convolution.hpp>
kahmed10's avatar
kahmed10 committed
25
#include <migraphx/gpu/deconvolution.hpp>
Khalique's avatar
Khalique committed
26
#include <migraphx/gpu/elu.hpp>
27
#include <migraphx/gpu/equal.hpp>
Paul's avatar
Paul committed
28
#include <migraphx/gpu/gemm.hpp>
29
#include <migraphx/gpu/greater.hpp>
30
#include <migraphx/gpu/int8_conv_pack.hpp>
31
#include <migraphx/gpu/leaky_relu.hpp>
32
#include <migraphx/gpu/less.hpp>
Shucai Xiao's avatar
Shucai Xiao committed
33
34
35
#include <migraphx/gpu/logical_and.hpp>
#include <migraphx/gpu/logical_or.hpp>
#include <migraphx/gpu/logical_xor.hpp>
36
37
38
39
#include <migraphx/gpu/lrn.hpp>
#include <migraphx/gpu/miopen.hpp>
#include <migraphx/gpu/quant_convolution.hpp>
#include <migraphx/gpu/rocblas.hpp>
40
#include <migraphx/gpu/unary_not.hpp>
turneram's avatar
turneram committed
41
#include <migraphx/gpu/where.hpp>
42
#include <migraphx/iterator_for.hpp>
43
#include <migraphx/program.hpp>
Paul's avatar
Paul committed
44
#include <utility>
45
#include <functional>
Khalique's avatar
Khalique committed
46
#include <algorithm>
Shucai Xiao's avatar
Shucai Xiao committed
47
#include <map>
Paul's avatar
Paul committed
48

Paul's avatar
Paul committed
49
namespace migraphx {
Paul's avatar
Paul committed
50
inline namespace MIGRAPHX_INLINE_NS {
Paul's avatar
Paul committed
51
namespace gpu {
Paul's avatar
Paul committed
52
53
54

struct miopen_apply
{
Shucai Xiao's avatar
Shucai Xiao committed
55
    module* mod          = nullptr;
56
    const lowering* pass = nullptr;
Shucai Xiao's avatar
Shucai Xiao committed
57
    std::unordered_map<std::string, std::function<instruction_ref(instruction_ref)>> apply_map{};
Shucai Xiao's avatar
Shucai Xiao committed
58
    instruction_ref last{};
59
    std::unordered_map<instruction_ref, std::string> prog_output_names{};
Shucai Xiao's avatar
Shucai Xiao committed
60
61
    bool offload_copy   = false;
    bool int8_x4_format = true;
Paul's avatar
Paul committed
62

63
    context& get_context() const
64
65
66
67
68
69
    {
        assert(pass != nullptr);
        assert(pass->ctx != nullptr);
        return *pass->ctx;
    }

Paul's avatar
Paul committed
70
71
72
73
74
75
76
    void check_shape(shape x, instruction_ref i)
    {
        assert(x == i->get_shape());
        (void)x;
        (void)i;
    }

77
78
    void create_output_names()
    {
Shucai Xiao's avatar
Shucai Xiao committed
79
        this->last = instruction::get_output_alias(std::prev(mod->end()));
80
81
        if(this->last->name() == "@return")
        {
82
            const auto& prog_outputs = last->inputs();
83
84
85
86
87
88
89
90
91
92
            std::vector<instruction_ref> outputs_alias(prog_outputs.size());

            std::transform(prog_outputs.begin(),
                           prog_outputs.end(),
                           outputs_alias.begin(),
                           [](const auto& i) { return instruction::get_output_alias(i); });

            std::size_t index = 0;
            for(auto ins : outputs_alias)
            {
Shucai Xiao's avatar
Shucai Xiao committed
93
                prog_output_names[ins] = mod->name() + ":#output_" + std::to_string(index++);
94
95
96
97
            }
        }
    }

98
99
    void init()
    {
Shucai Xiao's avatar
Shucai Xiao committed
100
        assert(mod != nullptr);
101
        assert(pass != nullptr);
102

Shucai Xiao's avatar
Shucai Xiao committed
103
104
105
106
107
108
109
#if ROCBLAS_VERSION_MAJOR >= 2 && ROCBLAS_VERSION_MINOR >= 38
        auto& ctx = get_context();
        rocblas_gemm_flags flag;
        rocblas_query_int8_layout_flag(ctx.get_stream().get_rocblas(), &flag);
        int8_x4_format = (flag == rocblas_gemm_flags_pack_int8x4);
#endif

Shucai Xiao's avatar
Shucai Xiao committed
110
        offload_copy = (mod->name() == "main") ? pass->offload_copy : false;
111
        create_output_names();
Paul's avatar
Paul committed
112

113
114
115
116
117
118
119
120
121
122
123
124
        add_generic_op("acos");
        add_generic_op("acosh");
        add_generic_op("add");
        add_generic_op("asin");
        add_generic_op("asinh");
        add_generic_op("atan");
        add_generic_op("atanh");
        add_generic_op("ceil");
        add_generic_op("contiguous");
        add_generic_op("cos");
        add_generic_op("cosh");
        add_generic_op("div");
125
        add_generic_op("equal");
126
127
128
        add_generic_op("erf");
        add_generic_op("exp");
        add_generic_op("floor");
129
130
        add_generic_op("greater");
        add_generic_op("less");
131
        add_generic_op("log");
Shucai Xiao's avatar
Shucai Xiao committed
132
133
134
        add_generic_op("logical_and");
        add_generic_op("logical_or");
        add_generic_op("logical_xor");
135
136
137
        add_generic_op("max");
        add_generic_op("min");
        add_generic_op("mul");
138
        add_generic_op("not");
139
140
141
142
143
144
145
146
147
148
149
150
151
152
153
        add_generic_op("pow");
        add_generic_op("prelu");
        add_generic_op("recip");
        add_generic_op("relu");
        add_generic_op("round");
        add_generic_op("rsqrt");
        add_generic_op("sigmoid");
        add_generic_op("sign");
        add_generic_op("sin");
        add_generic_op("sinh");
        add_generic_op("sqdiff");
        add_generic_op("sqrt");
        add_generic_op("sub");
        add_generic_op("tan");
        add_generic_op("tanh");
turneram's avatar
turneram committed
154
        add_generic_op("where");
155

Shucai Xiao's avatar
Shucai Xiao committed
156
        add_extend_op("abs");
157
158
159
160
161
        add_extend_op("argmax");
        add_extend_op("argmin");
        add_extend_op("clip");
        add_extend_op("concat");
        add_extend_op("convert");
Shucai Xiao's avatar
Shucai Xiao committed
162
        add_extend_op("elu");
163
        add_extend_op("gather");
Shucai Xiao's avatar
Shucai Xiao committed
164
        add_extend_op("leaky_relu");
165
        add_extend_op("logsoftmax");
Shucai Xiao's avatar
Shucai Xiao committed
166
        add_extend_op("lrn");
turneram's avatar
turneram committed
167
        add_extend_op("multinomial");
168
        add_extend_op("pad");
169
        add_extend_op("pooling");
170
        add_extend_op("prefix_scan_sum");
171
172
173
174
175
        add_extend_op("reduce_max");
        add_extend_op("reduce_mean");
        add_extend_op("reduce_min");
        add_extend_op("reduce_prod");
        add_extend_op("reduce_sum");
Cagri Eryilmaz's avatar
Cagri Eryilmaz committed
176
        add_extend_op("reverse");
177
178
179
        add_extend_op("rnn_var_sl_last_output");
        add_extend_op("rnn_var_sl_shift_output");
        add_extend_op("rnn_var_sl_shift_sequence");
180
        add_extend_op("scatter");
181
        add_extend_op("softmax");
Shucai Xiao's avatar
Shucai Xiao committed
182
        add_extend_op("topk");
183

184
185
        add_gemm_op<op::dot>("dot");
        add_gemm_op<op::quant_dot>("quant_dot");
186
        add_convolution_op();
kahmed10's avatar
kahmed10 committed
187
        add_deconvolution_op();
188
        add_quant_convolution_op();
189
        add_batch_norm_inference_op();
Shucai Xiao's avatar
Shucai Xiao committed
190
        add_neg_op();
Shucai Xiao's avatar
Shucai Xiao committed
191
        add_if_op();
Shucai Xiao's avatar
Shucai Xiao committed
192
        add_loop_op();
193
194
    }

195
196
    void copy_params()
    {
Shucai Xiao's avatar
Shucai Xiao committed
197
        if(not offload_copy)
198
            return;
199

Shucai Xiao's avatar
Shucai Xiao committed
200
        for(auto ins : iterator_for(*mod))
201
202
203
        {
            if(ins->name() != "@param")
                continue;
204

Shucai Xiao's avatar
Shucai Xiao committed
205
206
207
208
            // parameter no outputs, no need to insert copy to gpu
            if(ins->outputs().empty())
                continue;

209
210
            auto pos = std::next(ins);
            auto a   = insert_allocation(pos, ins->get_shape());
211
            auto c   = mod->insert_instruction(pos, make_op("hip::copy_to_gpu"), ins, a);
Shucai Xiao's avatar
Shucai Xiao committed
212
            mod->replace_instruction(ins, c);
213
        }
214
215

        // return instruction
Shucai Xiao's avatar
Shucai Xiao committed
216
        auto ret = std::prev(mod->end());
217
218
        if(ret->name() == "@return")
        {
219
            const auto& inputs = ret->inputs();
220
221
222

            // each input of ret need to be copied from gpu to host, and replace
            // output with copy output
223
            for(const auto& in : inputs)
224
            {
225
                auto p_output = mod->insert_instruction(ret, make_op("hip::copy_from_gpu"), in);
226
227
228
229
230
231
                instruction::replace_argument(ret, in, p_output);
            }
        }
        // else branch to handle legacy program without the return instruction
        else
        {
232
            mod->add_instruction(make_op("hip::copy_from_gpu"), ret);
233
        }
234
235
    }

Paul's avatar
Paul committed
236
237
    void apply()
    {
238
        init();
Shucai Xiao's avatar
Shucai Xiao committed
239
        for(auto it = mod->begin(); it != mod->end(); it++)
Paul's avatar
Paul committed
240
        {
Paul's avatar
Paul committed
241
            auto s = it->get_shape();
242
            if(apply_map.count(it->name()) > 0)
243
            {
244
                check_shape(s, apply_map.at(it->name())(it));
Paul's avatar
Paul committed
245
            }
Paul's avatar
Paul committed
246
        }
247

248
        copy_params();
Paul's avatar
Paul committed
249
250
    }

Paul's avatar
Paul committed
251
    instruction_ref insert_allocation(instruction_ref ins, const shape& s, std::string tag = "")
Paul's avatar
Paul committed
252
    {
253
        // Instruction's output is an input of the ret instruction
Shucai Xiao's avatar
Shucai Xiao committed
254
        if(offload_copy)
Paul's avatar
Paul committed
255
        {
256
257
            auto result = mod->insert_instruction(
                ins, make_op("hip::allocate", {{"shape", to_value(s)}, {"tag", std::move(tag)}}));
Paul's avatar
Paul committed
258
259
            return result;
        }
260
261
262
263

        auto ins_alias = instruction::get_output_alias(ins);
        if(last->name() == "@return" and tag.empty() and prog_output_names.count(ins_alias) > 0)
        {
Shucai Xiao's avatar
Shucai Xiao committed
264
            return mod->add_parameter(prog_output_names[ins_alias], s);
265
266
267
        }
        else if(ins == last and tag.empty())
        {
Shucai Xiao's avatar
Shucai Xiao committed
268
            return mod->add_parameter("output", s);
269
270
        }

271
272
        return mod->insert_instruction(
            ins, make_op("hip::allocate", {{"shape", to_value(s)}, {"tag", std::move(tag)}}));
Paul's avatar
Paul committed
273
274
    }

Shucai Xiao's avatar
Shucai Xiao committed
275
    void add_convolution_op()
Paul's avatar
Paul committed
276
    {
277
278
        apply_map.emplace("convolution", [=](instruction_ref ins) {
            auto&& op = any_cast<op::convolution>(ins->get_operator());
Paul's avatar
Paul committed
279

280
            auto conv = miopen_convolution{op, make_conv(op)};
281
            auto ws   = conv.find(get_context(), ins->get_shape(), to_shapes(ins->inputs()));
Paul's avatar
Paul committed
282

283
284
            auto workspace = insert_allocation(ins, ws, "workspace");
            auto output    = insert_allocation(ins, ins->get_shape());
kahmed10's avatar
kahmed10 committed
285

Shucai Xiao's avatar
Shucai Xiao committed
286
            return mod->replace_instruction(
kahmed10's avatar
kahmed10 committed
287
288
289
290
291
292
293
294
295
296
297
298
299
300
                ins, conv, ins->inputs().at(0), ins->inputs().at(1), workspace, output);
        });
    }

    void add_deconvolution_op()
    {
        apply_map.emplace("deconvolution", [=](instruction_ref ins) {
            auto&& op = any_cast<op::deconvolution>(ins->get_operator());

            auto conv = miopen_deconvolution{op, make_deconv(op)};
            auto ws   = conv.compile(get_context(), ins->get_shape(), to_shapes(ins->inputs()));

            auto workspace = insert_allocation(ins, ws, "workspace");
            auto output    = insert_allocation(ins, ins->get_shape());
Paul's avatar
Paul committed
301

Shucai Xiao's avatar
Shucai Xiao committed
302
            return mod->replace_instruction(
303
304
                ins, conv, ins->inputs().at(0), ins->inputs().at(1), workspace, output);
        });
Paul's avatar
Paul committed
305
306
    }

307
308
    template <typename Op>
    void add_gemm_op(const std::string& name)
309
310
    {
        apply_map.emplace(name, [=](instruction_ref ins) {
311
            std::vector<instruction_ref> refs = ins->inputs();
Shucai Xiao's avatar
Shucai Xiao committed
312
            if(refs.size() == 2)
313
314
            {
                auto output = insert_allocation(ins, ins->get_shape());
Shucai Xiao's avatar
Shucai Xiao committed
315
316
317
318
319
320
                refs.push_back(output);
            }
            else
            {
                auto c_alias = instruction::get_output_alias(refs.back());
                if(ins == last or refs.back()->outputs().size() > 1 or c_alias->inputs().empty())
321
                {
322
323
324
325
                    auto output = insert_allocation(ins, ins->get_shape());
                    auto copy_out =
                        mod->insert_instruction(ins, make_op("hip::copy"), refs.back(), output);
                    refs.back() = copy_out;
326
327
                    refs.push_back(copy_out);
                }
Shucai Xiao's avatar
Shucai Xiao committed
328
329
330
331
                else
                {
                    refs.push_back(refs.back());
                }
332
            }
Shucai Xiao's avatar
Shucai Xiao committed
333
            return mod->replace_instruction(
334
                ins, rocblas_gemm<Op>{Op{}, 1, 0, int8_x4_format}, refs);
335
336
337
        });
    }

338
339
340
341
342
    void add_quant_convolution_op()
    {
        apply_map.emplace("quant_convolution", [=](instruction_ref ins) {
            auto&& op = any_cast<op::quant_convolution>(ins->get_operator());
            auto conv = miopen_quant_convolution{op, make_conv(op)};
343
            auto ws   = conv.compile(get_context(), ins->get_shape(), to_shapes(ins->inputs()));
344

Shucai Xiao's avatar
Shucai Xiao committed
345
            auto args      = ins->inputs();
346
            auto workspace = insert_allocation(ins, ws, "workspace");
Shucai Xiao's avatar
Shucai Xiao committed
347
348
            auto output    = insert_allocation(ins, ins->get_shape());

Shucai Xiao's avatar
Shucai Xiao committed
349
            return mod->replace_instruction(ins, conv, args[0], args[1], workspace, output);
Shucai Xiao's avatar
Shucai Xiao committed
350
351
352
        });
    }

353
354
355
    void add_generic_op(const std::string& name) { add_generic_op(name, "gpu::" + name); }

    void add_generic_op(const std::string& op_name, const std::string& gpu_name)
Paul's avatar
Paul committed
356
    {
357
        apply_map.emplace(op_name, [=](instruction_ref ins) {
358
359
360
            auto output                       = insert_allocation(ins, ins->get_shape());
            std::vector<instruction_ref> refs = ins->inputs();
            refs.push_back(output);
Paul's avatar
Paul committed
361

Shucai Xiao's avatar
Shucai Xiao committed
362
            return mod->replace_instruction(ins, make_op(gpu_name), refs);
363
        });
Paul's avatar
Paul committed
364
    }
Paul's avatar
Paul committed
365

366
367
368
    void add_extend_op(const std::string& name) { add_extend_op(name, "gpu::" + name); }

    void add_extend_op(const std::string& op_name, const std::string& gpu_name)
Khalique's avatar
Khalique committed
369
    {
370
371
        apply_map.emplace(op_name, [=](instruction_ref ins) {
            auto&& op                         = ins->get_operator();
372
373
374
            auto output                       = insert_allocation(ins, ins->get_shape());
            std::vector<instruction_ref> refs = ins->inputs();
            refs.push_back(output);
Paul's avatar
Paul committed
375

Shucai Xiao's avatar
Shucai Xiao committed
376
            return mod->replace_instruction(ins, make_op(gpu_name, op.to_value()), refs);
377
        });
Khalique's avatar
Khalique committed
378
379
    }

Shucai Xiao's avatar
Shucai Xiao committed
380
    void add_batch_norm_inference_op()
381
    {
382
383
384
385
        apply_map.emplace("batch_norm_inference", [=](instruction_ref ins) {
            auto&& op       = any_cast<op::batch_norm_inference>(ins->get_operator());
            auto output     = insert_allocation(ins, ins->get_shape());
            shape old_shape = ins->inputs().at(1)->get_shape();
Shucai Xiao's avatar
Shucai Xiao committed
386
387
388
389
390
391
392
393
394
395
396
397
398
399
            auto input      = ins->inputs()[0];
            auto input_lens = input->get_shape().lens();
            std::vector<int64_t> rsp_lens(input_lens.size(), 1);
            // for per_activation case, also need to reshape input
            if(op.bn_mode == op::batch_norm_inference::per_activation)
            {
                std::copy(input_lens.begin() + 1, input_lens.end(), rsp_lens.begin() + 1);
            }
            else
            {
                rsp_lens[1] = static_cast<int64_t>(old_shape.elements());
            }

            auto reshape_op = op::reshape{rsp_lens};
400
401
            std::vector<instruction_ref> reshapes;
            std::transform(ins->inputs().begin() + 1,
Shucai Xiao's avatar
Shucai Xiao committed
402
403
                           ins->inputs().end(),
                           std::back_inserter(reshapes),
Shucai Xiao's avatar
Shucai Xiao committed
404
                           [&](auto i) { return mod->insert_instruction(ins, reshape_op, i); });
Shucai Xiao's avatar
Shucai Xiao committed
405

Shucai Xiao's avatar
Shucai Xiao committed
406
407
408
409
410
411
412
413
            return mod->replace_instruction(ins,
                                            miopen_batch_norm_inference{op},
                                            input,
                                            reshapes[0],
                                            reshapes[1],
                                            reshapes[2],
                                            reshapes[3],
                                            output);
Shucai Xiao's avatar
Shucai Xiao committed
414

415
        });
416
    }
Shucai Xiao's avatar
Shucai Xiao committed
417
418
419
420
421
422
423

    // use 0 - input to represent neg
    void add_neg_op()
    {
        apply_map.emplace("neg", [=](instruction_ref ins) {
            auto s = ins->get_shape();
            std::vector<float> zeros(s.elements(), 0.0f);
Shucai Xiao's avatar
Shucai Xiao committed
424
            auto l0     = mod->add_literal(literal(s, zeros));
Shucai Xiao's avatar
Shucai Xiao committed
425
            auto output = insert_allocation(ins, s);
Shucai Xiao's avatar
Shucai Xiao committed
426
            return mod->replace_instruction(
427
                ins, make_op("gpu::sub"), l0, ins->inputs().front(), output);
Shucai Xiao's avatar
Shucai Xiao committed
428
429
        });
    }
Shucai Xiao's avatar
Shucai Xiao committed
430

Shucai Xiao's avatar
Shucai Xiao committed
431
    // add input and output argument for the if operator
Shucai Xiao's avatar
Shucai Xiao committed
432
433
434
435
    void add_if_op()
    {
        apply_map.emplace("if", [=](instruction_ref ins) {
            std::vector<instruction_ref> inputs = ins->inputs();
436
437
438
            auto cpu_cond =
                mod->insert_instruction(ins, make_op("hip::copy_from_gpu"), inputs.front());
            auto sync_cond = mod->insert_instruction(ins, make_op("hip::sync_stream"), cpu_cond);
Shucai Xiao's avatar
Shucai Xiao committed
439
440
441
442
443
444
445
446
447
448
449
450
451
452
453
454
455
456
457
458
459
460
            inputs.front() = sync_cond;

            std::vector<module_ref> mod_args = ins->module_inputs();
            std::map<std::string, shape> name_shapes;
            for(const auto& smod : mod_args)
            {
                auto ps = smod->get_parameter_shapes();
                name_shapes.insert(ps.begin(), ps.end());
            }

            bool ins_output_allocated = false;
            for(auto& pn : name_shapes)
            {
                const auto& s = pn.second;
                instruction_ref output{};
                if(s == ins->get_shape() and not ins_output_allocated)
                {
                    output               = insert_allocation(ins, s);
                    ins_output_allocated = true;
                }
                else
                {
461
462
                    output = mod->insert_instruction(
                        ins, make_op("hip::allocate", {{"shape", to_value(s)}}));
Shucai Xiao's avatar
Shucai Xiao committed
463
464
465
466
467
468
469
                }
                inputs.push_back(output);
            }

            return mod->replace_instruction(ins, ins->get_operator(), inputs, mod_args);
        });
    }
Shucai Xiao's avatar
Shucai Xiao committed
470
471
472
473
474
475
476
477
478
479
480
481
482
483
484
485
486
487
488
489
490
491
492
493
494
495
496
497
498
499
500
501
502
503
504
505
506
507

    // replace the loop operator with gpu_loop operator
    void add_loop_op()
    {
        apply_map.emplace("loop", [=](instruction_ref ins) {
            std::vector<instruction_ref> inputs = ins->inputs();
            // copy max_iter from gpu to cpu
            auto cpu_max_iter =
                mod->insert_instruction(ins, make_op("hip::copy_from_gpu"), inputs.at(0));
            auto cpu_cond =
                mod->insert_instruction(ins, make_op("hip::copy_from_gpu"), inputs.at(1));
            auto synced_max_iter =
                mod->insert_instruction(ins, make_op("hip::sync_stream"), cpu_max_iter, cpu_cond);
            inputs.at(0)     = synced_max_iter;
            inputs.at(1)     = cpu_cond;
            auto copy_inputs = inputs;
            std::transform(
                copy_inputs.begin(), copy_inputs.end(), std::back_inserter(inputs), [&](auto in) {
                    return mod->insert_instruction(
                        ins, make_op("hip::allocate", {{"shape", to_value(in->get_shape())}}));
                });

            auto mod_args = ins->module_inputs();
            auto output   = insert_allocation(ins, ins->get_shape());

            const auto* sub_mod = mod_args.front();
            auto cond_out       = mod->insert_instruction(
                ins,
                make_op("hip::allocate",
                        {{"shape", to_value(sub_mod->get_output_shapes().front())}}));
            // add cond and mod outputs to the argument list
            inputs.push_back(cond_out);
            inputs.push_back(output);

            return mod->replace_instruction(
                ins, make_op("gpu::loop", ins->get_operator().to_value()), inputs, mod_args);
        });
    }
Paul's avatar
Paul committed
508
509
};

Shucai Xiao's avatar
Shucai Xiao committed
510
void lowering::apply(module& m) const { miopen_apply{&m, this}.apply(); }
Shucai Xiao's avatar
Shucai Xiao committed
511

Paul's avatar
Paul committed
512
} // namespace gpu
Paul's avatar
Paul committed
513
} // namespace MIGRAPHX_INLINE_NS
Paul's avatar
Paul committed
514
} // namespace migraphx