"src/targets/vscode:/vscode.git/clone" did not exist on "fc9b2a7d83d9d9e3340cdfa707cdb03c4581bb4b"
fuse_ops.cpp 34 KB
Newer Older
kahmed10's avatar
kahmed10 committed
1
2
#include <migraphx/pass_manager.hpp>
#include <migraphx/dead_code_elimination.hpp>
Paul's avatar
Paul committed
3
4
5
#include <migraphx/gpu/fuse_ops.hpp>
#include <migraphx/matcher.hpp>
#include <migraphx/gpu/miopen.hpp>
kahmed10's avatar
kahmed10 committed
6
#include <migraphx/gpu/clip.hpp>
Paul's avatar
Paul committed
7
#include <migraphx/gpu/convolution.hpp>
8
#include <migraphx/gpu/device_name.hpp>
9
#include <migraphx/gpu/oper.hpp>
kahmed10's avatar
kahmed10 committed
10
11
#include <migraphx/gpu/add.hpp>
#include <migraphx/gpu/mul.hpp>
12
#include <migraphx/gpu/gemm.hpp>
kahmed10's avatar
kahmed10 committed
13
#include <migraphx/gpu/device/layernorm.hpp>
kahmed10's avatar
kahmed10 committed
14
#include <migraphx/gpu/device/gelu.hpp>
Paul's avatar
Paul committed
15
#include <migraphx/gpu/device/mul_add.hpp>
16
17
18
19
20
#include <migraphx/gpu/device/add_clip.hpp>
#include <migraphx/gpu/device/add_relu.hpp>
#include <migraphx/gpu/device/add_sigmoid.hpp>
#include <migraphx/gpu/device/add_tanh.hpp>
#include <migraphx/gpu/device/mul_add_relu.hpp>
Paul's avatar
Paul committed
21
#include <migraphx/gpu/device/add.hpp>
22
23
24
#include <migraphx/match/layernorm.hpp>
#include <migraphx/match/gelu_erf.hpp>
#include <migraphx/match/gelu_tanh.hpp>
Paul's avatar
Paul committed
25
#include <migraphx/instruction.hpp>
26
#include <migraphx/register_op.hpp>
Paul's avatar
Paul committed
27
#include <migraphx/array.hpp>
kahmed10's avatar
kahmed10 committed
28
#include <migraphx/op/clip.hpp>
kahmed10's avatar
kahmed10 committed
29
#include <cmath>
30
#include <set>
Paul's avatar
Paul committed
31
32

namespace migraphx {
Paul's avatar
Paul committed
33
inline namespace MIGRAPHX_INLINE_NS {
Paul's avatar
Paul committed
34
35
namespace gpu {

36
37
MIGRAPHX_DECLARE_ENV_VAR(MIGRAPHX_DISABLE_MIOPEN_FUSION)

Paul's avatar
Paul committed
38
39
40
41
42
43
44
45
struct fusion
{
    using op_t = miopenFusionOpDescriptor_t;
    shared<fusion_plan_descriptor> fp;

    // Used as a temporary hack to keep descriptor references alive
    std::vector<std::shared_ptr<void>> storage;

Paul's avatar
Paul committed
46
    template <class T>
Paul's avatar
Paul committed
47
48
49
50
51
52
53
    auto keep_alive(T x)
    {
        auto result = share(std::move(x));
        storage.push_back(result);
        return result;
    }

54
55
    fusion() = default;

Paul's avatar
Paul committed
56
57
    fusion(const shape& input)
    {
58
        assert(input.standard());
Paul's avatar
Paul committed
59
        auto t = make_tensor(input);
Paul's avatar
Paul committed
60
        fp     = make_fusion_plan(t);
61
        assert(fp);
Paul's avatar
Paul committed
62
63
64
        keep_alive(std::move(t));
    }

65
66
    bool empty() const { return fp == nullptr; }

Paul's avatar
Paul committed
67
68
    op_t operator[](std::size_t i) const
    {
69
        assert(fp);
Paul's avatar
Paul committed
70
71
72
        op_t result;
        auto status = miopenFusionPlanGetOp(fp.get(), i, &result);
        if(status != miopenStatusSuccess)
Paul's avatar
Paul committed
73
            MIGRAPHX_THROW("Failed retrieving operator at " + std::to_string(i));
Paul's avatar
Paul committed
74
75
76
        return result;
    }

77
78
79
80
81
    auto get() const
    {
        assert(fp);
        return fp.get();
    }
Paul's avatar
Paul committed
82
83
84

    op_t create_bias(const shape& bias)
    {
85
        assert(fp);
Paul's avatar
Paul committed
86
        op_t result;
Paul's avatar
Paul committed
87
88
        auto b      = shape{bias.type(), {1, bias.lens().at(1), 1, 1}};
        auto t      = keep_alive(make_tensor(b));
Paul's avatar
Paul committed
89
90
        auto status = miopenCreateOpBiasForward(fp.get(), &result, t.get());
        if(status != miopenStatusSuccess)
Paul's avatar
Paul committed
91
            MIGRAPHX_THROW("Creating operator failed");
Paul's avatar
Paul committed
92
93
94
95
96
        return result;
    }

    op_t create_relu()
    {
97
        assert(fp);
Paul's avatar
Paul committed
98
99
100
        op_t result;
        auto status = miopenCreateOpActivationForward(fp.get(), &result, miopenActivationRELU);
        if(status != miopenStatusSuccess)
Paul's avatar
Paul committed
101
            MIGRAPHX_THROW("Creating operator failed");
Paul's avatar
Paul committed
102
103
104
105
106
        return result;
    }

    op_t create_conv(const op::convolution& op, const shape& weights)
    {
107
        assert(fp);
Paul's avatar
Paul committed
108
        op_t result;
Paul's avatar
Paul committed
109
110
        auto cd     = keep_alive(make_conv(op));
        auto t      = keep_alive(make_tensor(weights));
Paul's avatar
Paul committed
111
112
        auto status = miopenCreateOpConvForward(fp.get(), &result, cd.get(), t.get());
        if(status != miopenStatusSuccess)
Paul's avatar
Paul committed
113
            MIGRAPHX_THROW("Creating operator failed");
Paul's avatar
Paul committed
114
115
        return result;
    }
Paul's avatar
Paul committed
116
117
118

    shape get_workspace(context&)
    {
119
        // assert(fp);
Paul's avatar
Paul committed
120
121
122
123
124
        // TODO: Use zero workspace for now
        std::size_t ws_size = 0;
        // int algo_count = 1;
        // miopenConvFwdAlgorithm_t algo;
        // miopenFusionPlanConvolutionGetAlgo(fp.get(), 1, &algo_count, &algo);
Paul's avatar
Paul committed
125
126
        // miopenFusionPlanGetWorkSpaceSize(ctx.get_stream().get_miopen(), fp.get(), &ws_size,
        // algo);
Paul's avatar
Paul committed
127
128
129
        return shape{shape::int8_type, {ws_size}};
    }

130
    bool compile(context& ctx)
Paul's avatar
Paul committed
131
    {
132
        assert(fp);
133
134
        return miopenCompileFusionPlan(ctx.get_stream().get_miopen(), fp.get()) ==
               miopenStatusSuccess;
Paul's avatar
Paul committed
135
136
    }

Paul's avatar
Paul committed
137
138
139
140
    argument execute(context& ctx,
                     const fused_operator_args& fargs,
                     const argument& x,
                     const argument& y) const
Paul's avatar
Paul committed
141
    {
142
        assert(fp);
Paul's avatar
Paul committed
143
144
        auto x_td   = make_tensor(x.get_shape());
        auto y_td   = make_tensor(y.get_shape());
Paul's avatar
Paul committed
145
        auto status = miopenExecuteFusionPlan(ctx.get_stream().get_miopen(),
Paul's avatar
Paul committed
146
147
148
149
150
151
                                              fp.get(),
                                              x_td.get(),
                                              x.implicit(),
                                              y_td.get(),
                                              y.implicit(),
                                              fargs.get());
Paul's avatar
Paul committed
152
        if(status != miopenStatusSuccess)
Paul's avatar
Paul committed
153
            MIGRAPHX_THROW("Failed to execute fusion plan");
Paul's avatar
Paul committed
154
155
        return y;
    }
Paul's avatar
Paul committed
156
157
};

158
159
160
161
162
163
const std::unordered_set<std::string>& get_supported_archs()
{
    static std::unordered_set<std::string> supported_archs{"gfx900", "gfx906", "gfx908", "gfx1030"};
    return supported_archs;
}

Paul's avatar
Paul committed
164
MIGRAPHX_PRED_MATCHER(bias_shape, instruction_ref ins)
Paul's avatar
Paul committed
165
166
{
    auto&& s = ins->get_shape();
Paul's avatar
Paul committed
167
168
    return s.broadcasted() and s.strides().size() == 4 and s.strides()[0] == 0 and
           s.strides()[1] != 0 and s.strides()[2] == 0 and s.strides()[3] == 0;
Paul's avatar
Paul committed
169
170
}

Paul's avatar
Paul committed
171
MIGRAPHX_PRED_MATCHER(fusable_conv, instruction_ref ins)
Paul's avatar
Paul committed
172
{
173
    const auto device_name = trim(split_string(get_device_name(), ':').front());
174
175
    if(not contains(get_supported_archs(), device_name))
        return false;
176
177
    if(enabled(MIGRAPHX_DISABLE_MIOPEN_FUSION{}))
        return false;
Paul's avatar
Paul committed
178
179
    if(ins->name() != "gpu::convolution")
        return false;
Paul's avatar
Paul committed
180
181
    if(ins->get_shape().type() != shape::float_type)
        return false;
Paul's avatar
Paul committed
182
183
184
    auto wei = ins->inputs().at(1)->get_shape();
    assert(wei.lens().size() == 4);
    auto conv = any_cast<miopen_convolution>(ins->get_operator());
Khalique's avatar
Khalique committed
185
    if(conv.op.group > 1)
Khalique's avatar
Khalique committed
186
        return false;
Paul's avatar
Paul committed
187
    if(wei.lens()[1] > 512 and conv.algo != miopenConvolutionFwdAlgoWinograd)
Paul's avatar
Paul committed
188
        return false;
189
190
191
192
193
194

    // Do not fuse non-symmetric input
    auto input_lens = ins->inputs().at(0)->get_shape().lens();
    if(input_lens[2] != input_lens[3] or wei.lens()[2] != wei.lens()[3])
        return false;

Paul's avatar
Paul committed
195
    auto op = conv.op;
196
197
    // Dont fuse winograd for non-3x3s since there is no fused windograd for those configs
    if(conv.algo == miopenConvolutionFwdAlgoWinograd and wei.lens()[2] != 3 and
198
       wei.lens()[3] != 3 and contains({{1, 1}}, op.stride))
199
        return false;
kahmed10's avatar
kahmed10 committed
200
    return contains({{0, 0, 0, 0}, {1, 1, 1, 1}, {2, 2, 2, 2}}, op.padding) and
201
           contains({{0, 0}, {1, 1}}, op.stride) and contains({{1, 1}}, op.dilation);
Paul's avatar
Paul committed
202
203
}

204
struct hip_triadd : ternary_device<hip_triadd, &device::add>
Paul's avatar
Paul committed
205
206
{
};
207
MIGRAPHX_REGISTER_OP(hip_triadd)
Paul's avatar
Paul committed
208

209
struct hip_triadd_clip : quinary_device<hip_triadd_clip, &device::add_clip>
kahmed10's avatar
kahmed10 committed
210
211
{
};
212
MIGRAPHX_REGISTER_OP(hip_triadd_clip)
kahmed10's avatar
kahmed10 committed
213

214
struct hip_add_clip : quaternary_device<hip_add_clip, &device::add_clip>
kahmed10's avatar
kahmed10 committed
215
216
{
};
217
MIGRAPHX_REGISTER_OP(hip_add_clip)
kahmed10's avatar
kahmed10 committed
218

219
struct hip_triadd_relu : ternary_device<hip_triadd_relu, &device::add_relu>
Paul's avatar
Paul committed
220
221
{
};
222
MIGRAPHX_REGISTER_OP(hip_triadd_relu)
Paul's avatar
Paul committed
223

224
225
226
struct hip_triadd_sigmoid : ternary_device<hip_triadd_sigmoid, &device::add_sigmoid>
{
};
227
MIGRAPHX_REGISTER_OP(hip_triadd_sigmoid)
228
229
230
231

struct hip_triadd_tanh : ternary_device<hip_triadd_tanh, &device::add_tanh>
{
};
232
MIGRAPHX_REGISTER_OP(hip_triadd_tanh)
233
234
235
236

struct hip_add_relu : binary_device<hip_add_relu, &device::add_relu>
{
};
237
MIGRAPHX_REGISTER_OP(hip_add_relu)
238
239
240
241

struct hip_add_sigmoid : binary_device<hip_add_relu, &device::add_sigmoid>
{
};
242
MIGRAPHX_REGISTER_OP(hip_add_sigmoid)
243
244

struct hip_add_tanh : binary_device<hip_add_tanh, &device::add_tanh>
Paul's avatar
Paul committed
245
246
{
};
247
MIGRAPHX_REGISTER_OP(hip_add_tanh)
Paul's avatar
Paul committed
248

kahmed10's avatar
kahmed10 committed
249
250
struct hip_layernorm : unary_device<hip_layernorm, &device::layernorm>
{
251
252
    // Empty finalize to skip dimension reduction
    void finalize(context&, const shape&, const std::vector<shape>&) {}
kahmed10's avatar
kahmed10 committed
253
};
254
MIGRAPHX_REGISTER_OP(hip_layernorm)
kahmed10's avatar
kahmed10 committed
255

Paul Fultz II's avatar
Paul Fultz II committed
256
257
struct hip_triadd_layernorm : ternary_device<hip_triadd_layernorm, &device::triadd_layernorm>
{
258
259
260
261
262
    shape compute_shape(const std::vector<shape>& inputs) const
    {
        check_shapes{inputs, *this}.has(4).standard();
        return inputs[0];
    }
Paul Fultz II's avatar
Paul Fultz II committed
263
264
265
266
267
    // Empty finalize to skip dimension reduction
    void finalize(context&, const shape&, const std::vector<shape>&) {}
};
MIGRAPHX_REGISTER_OP(hip_triadd_layernorm)

kahmed10's avatar
kahmed10 committed
268
269
270
struct hip_gelu : unary_device<hip_gelu, &device::gelu>
{
};
271
MIGRAPHX_REGISTER_OP(hip_gelu)
kahmed10's avatar
kahmed10 committed
272
273
274
275

struct hip_add_gelu : binary_device<hip_add_gelu, &device::add_gelu>
{
};
276
MIGRAPHX_REGISTER_OP(hip_add_gelu)
kahmed10's avatar
kahmed10 committed
277
278
279
280

struct hip_gelu_new : unary_device<hip_gelu_new, &device::gelu_new>
{
};
281
MIGRAPHX_REGISTER_OP(hip_gelu_new)
kahmed10's avatar
kahmed10 committed
282
283
284
285

struct hip_add_gelu_new : binary_device<hip_add_gelu_new, &device::add_gelu_new>
{
};
286
MIGRAPHX_REGISTER_OP(hip_add_gelu_new)
kahmed10's avatar
kahmed10 committed
287

288
struct hip_mul_add : ternary_device<hip_mul_add, &device::mul_add>
Paul's avatar
Paul committed
289
290
{
};
291
MIGRAPHX_REGISTER_OP(hip_mul_add)
Paul's avatar
Paul committed
292

293
struct hip_mul_add_relu : ternary_device<hip_mul_add_relu, &device::mul_add_relu>
Paul's avatar
Paul committed
294
295
{
};
296
MIGRAPHX_REGISTER_OP(hip_mul_add_relu)
Paul's avatar
Paul committed
297

Paul's avatar
Paul committed
298
299
300
void move_broadcasted_back(std::vector<instruction_ref>& args)
{
    // Ensure the last arguments is the broadcasted one
Paul's avatar
Paul committed
301
    auto last = std::prev(args.end());
Paul's avatar
Paul committed
302
303
    auto it =
        std::find_if(args.begin(), last, [](auto arg) { return arg->get_shape().broadcasted(); });
Paul's avatar
Paul committed
304
305
    if(it != last)
        std::swap(*it, *std::prev(last));
Paul's avatar
Paul committed
306
307
308
309
310
}

void move_standard_front(std::vector<instruction_ref>& args)
{
    // Ensure the first arguments is the standard one
Paul's avatar
Paul committed
311
    auto last = std::prev(args.end());
Paul's avatar
Paul committed
312
313
    auto it =
        std::find_if(args.begin(), last, [](auto arg) { return arg->get_shape().standard(); });
Paul's avatar
Paul committed
314
    if(it != last)
Paul's avatar
Paul committed
315
316
317
        std::swap(*it, args.front());
}

318
319
auto gpu_name(const std::string& s) { return match::name("gpu::" + s); }

kahmed10's avatar
kahmed10 committed
320
321
struct find_layernorm
{
322
    auto matcher() const { return match::layernorm(&gpu_name); }
kahmed10's avatar
kahmed10 committed
323

324
    void apply(module& m, const match::matcher_result& r) const
kahmed10's avatar
kahmed10 committed
325
326
327
328
329
    {
        auto ins   = r.result;
        auto x_ins = r.instructions["x"];
        auto args  = ins->inputs();

330
331
332
333
334
335
336
337
338
        // We dont fuse for non-standard layouts
        if(not x_ins->get_shape().standard())
            return;

        auto relements = x_ins->get_shape().lens().back();

        if(relements > 1024 or (relements % 4 != 0 and relements > 256))
            return;

339
        m.replace_instruction(ins, hip_layernorm{}, x_ins, args.back());
kahmed10's avatar
kahmed10 committed
340
341
342
    }
};

Paul Fultz II's avatar
Paul Fultz II committed
343
344
345
346
347
348
349
350
struct find_triadd_layernorm
{
    auto matcher() const
    {
        return match::name("gpu::layernorm")(match::arg(0)(match::name("gpu::triadd")(
            match::used_once(), match::all_of[match::inputs()](match::standard_shape()))));
    }

351
    void apply(module& m, const match::matcher_result& r) const
Paul Fultz II's avatar
Paul Fultz II committed
352
353
354
    {
        auto ins    = r.result;
        auto triadd = ins->inputs().front();
355
        m.replace_instruction(ins, hip_triadd_layernorm{}, triadd->inputs());
Paul Fultz II's avatar
Paul Fultz II committed
356
357
358
    }
};

kahmed10's avatar
kahmed10 committed
359
360
struct find_gelu
{
361
    auto matcher() const { return match::gelu_erf(&gpu_name); }
kahmed10's avatar
kahmed10 committed
362

363
    void apply(module& m, const match::matcher_result& r) const
kahmed10's avatar
kahmed10 committed
364
365
366
367
368
    {
        auto ins   = r.result;
        auto x_ins = r.instructions["x"];
        auto args  = ins->inputs();

369
        m.replace_instruction(ins, hip_gelu{}, x_ins, args.back());
kahmed10's avatar
kahmed10 committed
370
371
372
373
374
375
376
377
378
379
    }
};

struct find_add_gelu
{
    auto matcher() const
    {
        return match::name("gpu::gelu")(match::arg(0)(match::name("gpu::add").bind("add")));
    }

380
    void apply(module& m, const match::matcher_result& r) const
kahmed10's avatar
kahmed10 committed
381
382
383
384
385
386
387
388
    {
        auto add_ins = r.instructions["add"];
        auto ins     = r.result;
        auto args    = add_ins->inputs();
        move_standard_front(args);
        move_broadcasted_back(args);

        args.back() = ins->inputs().back();
389
        m.replace_instruction(ins, hip_add_gelu{}, args);
kahmed10's avatar
kahmed10 committed
390
391
392
393
394
    }
};

struct find_gelu_new
{
kahmed10's avatar
kahmed10 committed
395
    bool fast_math = true;
kahmed10's avatar
kahmed10 committed
396

397
    auto matcher() const { return match::gelu_tanh(&gpu_name); }
kahmed10's avatar
kahmed10 committed
398

399
    void apply(module& m, const match::matcher_result& r) const
kahmed10's avatar
kahmed10 committed
400
401
402
403
404
    {
        auto ins   = r.result;
        auto x_ins = r.instructions["x"];
        auto args  = ins->inputs();

Paul Fultz II's avatar
Paul Fultz II committed
405
        if(fast_math)
406
            m.replace_instruction(ins, hip_gelu{}, x_ins, args.back());
Paul Fultz II's avatar
Paul Fultz II committed
407
        else
408
            m.replace_instruction(ins, hip_gelu_new{}, x_ins, args.back());
kahmed10's avatar
kahmed10 committed
409
410
411
412
413
414
415
416
417
418
    }
};

struct find_add_gelu_new
{
    auto matcher() const
    {
        return match::name("gpu::gelu_new")(match::arg(0)(match::name("gpu::add").bind("add")));
    }

419
    void apply(module& m, const match::matcher_result& r) const
kahmed10's avatar
kahmed10 committed
420
421
422
423
424
425
426
427
    {
        auto add_ins = r.instructions["add"];
        auto ins     = r.result;
        auto args    = add_ins->inputs();
        move_standard_front(args);
        move_broadcasted_back(args);

        args.back() = ins->inputs().back();
428
        m.replace_instruction(ins, hip_add_gelu_new{}, args);
kahmed10's avatar
kahmed10 committed
429
430
431
    }
};

kahmed10's avatar
kahmed10 committed
432
433
434
435
436
437
struct find_add_clip
{
    auto matcher() const
    {
        return match::name(std::unordered_set<std::string>{"gpu::clip", "gpu::clipped_relu"})(
            match::arg(0)(match::any_of(match::name("gpu::add"),
kahmed10's avatar
kahmed10 committed
438
                                        match::name("gpu::triadd"),
kahmed10's avatar
kahmed10 committed
439
440
441
442
                                        match::any_of[match::inputs()](match::standard_shape()))
                              .bind("add")));
    }

443
    void apply(module& m, const match::matcher_result& r) const
kahmed10's avatar
kahmed10 committed
444
    {
kahmed10's avatar
kahmed10 committed
445
446
447
448
449
450
451
452
453
454
        auto add_ins  = r.instructions["add"];
        auto ins      = r.result;
        auto ins_args = ins->inputs();
        auto add_args = add_ins->inputs();
        move_standard_front(add_args);
        move_broadcasted_back(add_args);

        // Use the allocation from the clip operator
        add_args.pop_back();
        add_args.insert(add_args.end(), std::next(ins_args.begin()), ins_args.end());
kahmed10's avatar
kahmed10 committed
455
        if(add_ins->name() == "gpu::add")
456
            m.replace_instruction(ins, hip_add_clip{}, add_args);
kahmed10's avatar
kahmed10 committed
457
        else if(add_ins->name() == "gpu::triadd")
458
            m.replace_instruction(ins, hip_triadd_clip{}, add_args);
kahmed10's avatar
kahmed10 committed
459
460
461
    }
};

462
struct find_add_unary
Paul's avatar
Paul committed
463
{
464
465
466
    std::string op_name;
    operation binary_add_op;
    operation ternary_add_op;
Paul's avatar
Paul committed
467
468
    auto matcher() const
    {
469
        return match::name(op_name)(match::arg(0)(
Paul's avatar
Paul committed
470
            match::used_once(),
Paul's avatar
Paul committed
471
            match::any_of(match::name("gpu::add"),
kahmed10's avatar
kahmed10 committed
472
                          match::name("gpu::triadd"),
Paul's avatar
Paul committed
473
474
475
                          match::any_of(match::name("@literal"),
                                        match::any_of[match::inputs()](match::standard_shape())))
                .bind("add")));
Paul's avatar
Paul committed
476
    }
Paul's avatar
Paul committed
477

478
    void apply(module& m, const match::matcher_result& r) const
Paul's avatar
Paul committed
479
    {
Paul's avatar
Paul committed
480
        auto add_ins = r.instructions["add"];
Paul's avatar
Paul committed
481
482
        auto ins     = r.result;
        auto args    = add_ins->inputs();
Paul's avatar
Paul committed
483
484
485
        move_standard_front(args);
        move_broadcasted_back(args);

Paul's avatar
Paul committed
486
        // Use the allocation from the relu operator
Paul's avatar
Paul committed
487
        args.back() = ins->inputs().back();
Paul's avatar
Paul committed
488
        if(add_ins->name() == "gpu::add")
489
            m.replace_instruction(ins, binary_add_op, args);
kahmed10's avatar
kahmed10 committed
490
        else if(add_ins->name() == "gpu::triadd")
491
            m.replace_instruction(ins, ternary_add_op, args);
Paul's avatar
Paul committed
492
493
494
    }
};

Paul's avatar
Paul committed
495
struct find_triadd
Paul's avatar
Paul committed
496
497
498
{
    auto matcher() const
    {
Paul's avatar
Paul committed
499
        return match::name("gpu::add")(match::either_arg(0, 1)(
Paul's avatar
Paul committed
500
            match::name("gpu::add")(match::used_once()).bind("add"),
Paul's avatar
Paul committed
501
502
503
            match::any(match::any_of(match::name("@literal"),
                                     match::any_of[match::inputs()](match::standard_shape())))
                .bind("input")));
Paul's avatar
Paul committed
504
505
    }

506
    void apply(module& m, const match::matcher_result& r) const
Paul's avatar
Paul committed
507
    {
Paul's avatar
Paul committed
508
509
510
511
        auto add_ins   = r.instructions["add"];
        auto input_ins = r.instructions["input"];
        auto ins       = r.result;
        auto args      = add_ins->inputs();
512

Paul's avatar
Paul committed
513
        auto is_broadcasted = [](auto arg) { return arg->get_shape().broadcasted(); };
514
        if(std::count_if(args.begin(), args.end(), is_broadcasted) > 2)
Paul's avatar
Paul committed
515
516
            return;
        args.insert(args.begin(), input_ins);
Paul's avatar
Paul committed
517
518
519
        move_standard_front(args);
        move_broadcasted_back(args);

Paul's avatar
Paul committed
520
        args.back() = ins->inputs().back();
521
        m.replace_instruction(ins, hip_triadd{}, args);
Paul's avatar
Paul committed
522
    }
Paul's avatar
Paul committed
523
524
};

Paul's avatar
Paul committed
525
526
527
528
struct find_mul_add
{
    auto matcher() const
    {
Paul's avatar
Paul committed
529
530
        return match::name("gpu::add")(match::either_arg(0, 1)(
            match::name("gpu::mul")(match::used_once()).bind("mul"), match::any().bind("b")));
Paul's avatar
Paul committed
531
532
    }

533
    void apply(module& m, const match::matcher_result& r) const
Paul's avatar
Paul committed
534
    {
Paul's avatar
Paul committed
535
536
537
538
        auto mul_ins = r.instructions["mul"];
        auto b_ins   = r.instructions["b"];
        auto ins     = r.result;
        auto args    = mul_ins->inputs();
Paul's avatar
Paul committed
539
540
541
542
543
544
545
        assert(mul_ins != b_ins);

        move_standard_front(args);
        move_broadcasted_back(args);
        args.insert(std::prev(args.end()), b_ins);

        args.back() = ins->inputs().back();
546
        m.replace_instruction(ins, hip_mul_add{}, args);
Paul's avatar
Paul committed
547
548
549
    }
};

Paul's avatar
Paul committed
550
551
552
553
struct find_mul_add_relu
{
    auto matcher() const
    {
Paul's avatar
Paul committed
554
        return match::name("gpu::relu")(
kahmed10's avatar
kahmed10 committed
555
            match::arg(0)(match::name("gpu::mul_add")(match::used_once()).bind("mul_add")));
Paul's avatar
Paul committed
556
557
    }

558
    void apply(module& m, const match::matcher_result& r) const
Paul's avatar
Paul committed
559
560
    {
        auto mul_add_ins = r.instructions["mul_add"];
Paul's avatar
Paul committed
561
562
        auto ins         = r.result;
        auto args        = mul_add_ins->inputs();
Paul's avatar
Paul committed
563
564
565

        // Use the allocation from the relu operator
        args.back() = ins->inputs().back();
566
        m.replace_instruction(ins, hip_mul_add_relu{}, args);
Paul's avatar
Paul committed
567
568
569
    }
};

570
571
572
573
574
575
576
577
578
579
580
581
582
583
584
585
586
587
588
589
590
591
592
593
594
struct miopen_fusion
{
    struct fuse_op_data
    {
        operation op;
        float alpha = 1;
        float beta  = 0;
    };
    struct fuse_op : fuse_op_data, reflect_equality<fuse_op>, reflect_stream<fuse_op>
    {
        template <class Self, class F>
        static auto reflect(Self& self, F f)
        {
            return pack(f(self.op, "op"), f(self.alpha, "alpha"), f(self.beta, "beta"));
        }
    };
    std::vector<fuse_op> ops = {};
    fusion f                 = {};
    std::function<void(context&, const fusion&, const std::vector<argument>&)> execute;
    template <class Self, class F>
    static auto reflect(Self& self, F f)
    {
        return pack(f(self.ops, "ops"));
    }

595
596
597
598
599
    std::ptrdiff_t output_alias(const std::vector<shape>& shapes) const
    {
        return shapes.size() - 1;
    }

600
601
602
603
604
605
606
607
608
609
610
611
612
613
614
615
616
617
618
619
620
621
622
623
624
625
626
627
628
629
630
631
632
633
634
635
636
637
638
639
640
641
642
643
644
645
646
647
648
649
650
651
652
653
654
655
656
657
658
659
660
661
662
663
664
665
666
667
668
669
670
671
672
673
674
675
676
677
678
679
680
681
682
683
684
685
    value compile(context& ctx, const shape&, std::vector<shape> inputs)
    {
        // Compensate for allocation
        inputs.pop_back();
        std::size_t i = 0;
        f             = fusion(inputs[i]);
        i++;
        std::vector<std::function<void(const fused_operator_args&, const std::vector<argument>&)>>
            invokers;
        for(auto&& fop : ops)
        {
            if(i > inputs.size())
            {
                f = {};
                return {};
            }
            if(fop.op.name() == "convolution")
            {
                auto* mop = f.create_conv(any_cast<op::convolution>(fop.op), inputs[i]);
                invokers.push_back(
                    [=](const fused_operator_args& fargs, const std::vector<argument>& args) {
                        miopenSetOpArgsConvForward(
                            fargs.get(), mop, &fop.alpha, &fop.beta, args[i].implicit());
                    });
                i++;
            }
            else if(fop.op.name() == "add")
            {
                auto* mop = f.create_bias(inputs[i]);
                invokers.push_back(
                    [=](const fused_operator_args& fargs, const std::vector<argument>& args) {
                        miopenSetOpArgsBiasForward(
                            fargs.get(), mop, &fop.alpha, &fop.beta, args[i].implicit());
                    });
                i++;
            }
            else if(fop.op.name() == "relu")
            {
                auto* mop = f.create_relu();
                invokers.push_back([=](const fused_operator_args& fargs,
                                       const std::vector<argument>&) {
                    miopenSetOpArgsActivForward(fargs.get(), mop, &fop.alpha, &fop.beta, 0, 0, 0);
                });
            }
            else
            {
                f = {};
                return {};
            }
        }
        if(not f.compile(ctx))
        {
            f = {};
            return {};
        }
        execute = [invokers](context& c, const fusion& ff, const std::vector<argument>& args) {
            auto fargs = make_fused_args();
            for(auto&& invoker : invokers)
                invoker(fargs, args);
            ff.execute(c, fargs, args.front(), args.back());
        };
        return {{"workspace", f.get_workspace(ctx).bytes()}};
    }
    void finalize(context& ctx, const shape& output_shape, const std::vector<shape>& inputs)
    {
        if(not f.empty())
            return;
        auto v = compile(ctx, output_shape, inputs);
        if(not v.is_object())
            MIGRAPHX_THROW("Failed to compile fusion plan");
    }
    std::string name() const { return "gpu::miopen_fusion"; }
    shape compute_shape(const std::vector<shape>& inputs) const
    {
        if(ops.empty())
            return {};
        // TODO: Check number of arguments
        return ops.front().op.compute_shape({inputs[0], inputs[1]});
    }
    argument compute(context& ctx, const shape&, const std::vector<argument>& args) const
    {
        execute(ctx, f, args);
        return args.back();
    }
};

Paul's avatar
Paul committed
686
687
688
struct miopen_conv_bias
{
    op::convolution op;
689
    fusion fp         = {};
690
691
    fusion::op_t conv = {};
    fusion::op_t bias = {};
Paul's avatar
Paul committed
692

Paul's avatar
Paul committed
693
694
695
696
697
698
    template <class Self, class F>
    static auto reflect(Self& self, F f)
    {
        return op::convolution::reflect(self.op, f);
    }

Paul's avatar
Paul committed
699
700
701
702
703
    std::string name() const { return "gpu::conv_bias"; }
    shape compute_shape(const std::vector<shape>& inputs) const
    {
        check_shapes{inputs, *this}.has(5);
        // TODO: Check slices
kahmed10's avatar
kahmed10 committed
704
        return op.normalize_compute_shape({inputs.at(0), inputs.at(1)});
Paul's avatar
Paul committed
705
    }
Paul's avatar
Paul committed
706
    argument compute(context& ctx, const shape&, const std::vector<argument>& args) const
Paul's avatar
Paul committed
707
    {
Paul's avatar
Paul committed
708
        auto fargs  = make_fused_args();
Paul's avatar
Paul committed
709
        float alpha = 1;
Paul's avatar
Paul committed
710
        float beta  = 0;
Paul's avatar
Paul committed
711
712
        miopenSetOpArgsConvForward(fargs.get(), conv, &alpha, &beta, args[1].implicit());
        miopenSetOpArgsBiasForward(fargs.get(), bias, &alpha, &beta, args[3].implicit());
713
        return fp.execute(ctx, fargs, args[0], args[4]);
Paul's avatar
Paul committed
714
715
    }

716
717
    void finalize(context& ctx, const shape&, const std::vector<shape>& inputs)
    {
718
719
720
721
        fp   = fusion(inputs[0]);
        conv = fp.create_conv(op, inputs[1]);
        bias = fp.create_bias(inputs[3]);
        if(not fp.compile(ctx))
722
            MIGRAPHX_THROW("Failed to compile fusion plan");
723
724
    }

725
    shape get_workspace(context& ctx) { return fp.get_workspace(ctx); }
Paul's avatar
Paul committed
726
727
728
729
    std::ptrdiff_t output_alias(const std::vector<shape>& shapes) const
    {
        return shapes.size() - 1;
    }
Paul's avatar
Paul committed
730
};
731
MIGRAPHX_REGISTER_OP(miopen_conv_bias)
Paul's avatar
Paul committed
732

Paul's avatar
Add cbr  
Paul committed
733
734
735
struct miopen_conv_bias_relu
{
    op::convolution op;
736
    fusion fp         = {};
737
738
739
    fusion::op_t conv = {};
    fusion::op_t bias = {};
    fusion::op_t relu = {};
Paul's avatar
Add cbr  
Paul committed
740

Paul's avatar
Paul committed
741
742
743
744
745
746
    template <class Self, class F>
    static auto reflect(Self& self, F f)
    {
        return op::convolution::reflect(self.op, f);
    }

Paul's avatar
Add cbr  
Paul committed
747
748
749
750
751
    std::string name() const { return "gpu::conv_bias_relu"; }
    shape compute_shape(const std::vector<shape>& inputs) const
    {
        check_shapes{inputs, *this}.has(5);
        // TODO: Check slices
kahmed10's avatar
kahmed10 committed
752
        return op.normalize_compute_shape({inputs.at(0), inputs.at(1)});
Paul's avatar
Add cbr  
Paul committed
753
    }
Paul's avatar
Paul committed
754
    argument compute(context& ctx, const shape&, const std::vector<argument>& args) const
Paul's avatar
Add cbr  
Paul committed
755
756
    {
        auto fargs  = make_fused_args();
Paul's avatar
Paul committed
757
        float alpha = 1;
Paul's avatar
Paul committed
758
        float beta  = 0;
Paul's avatar
Add cbr  
Paul committed
759
760
        miopenSetOpArgsConvForward(fargs.get(), conv, &alpha, &beta, args[1].implicit());
        miopenSetOpArgsBiasForward(fargs.get(), bias, &alpha, &beta, args[3].implicit());
Paul's avatar
Paul committed
761
        miopenSetOpArgsActivForward(fargs.get(), relu, &alpha, &beta, 0, 0, 0);
762
        return fp.execute(ctx, fargs, args[0], args[4]);
Paul's avatar
Add cbr  
Paul committed
763
    }
764
765
    void finalize(context& ctx, const shape&, const std::vector<shape>& inputs)
    {
766
767
768
769
770
        fp   = fusion(inputs[0]);
        conv = fp.create_conv(op, inputs[1]);
        bias = fp.create_bias(inputs[3]);
        relu = fp.create_relu();
        fp.compile(ctx);
771
772
    }

773
    shape get_workspace(context& ctx) { return fp.get_workspace(ctx); }
Paul's avatar
Paul committed
774
775
776
777
    std::ptrdiff_t output_alias(const std::vector<shape>& shapes) const
    {
        return shapes.size() - 1;
    }
Paul's avatar
Add cbr  
Paul committed
778
};
779
MIGRAPHX_REGISTER_OP(miopen_conv_bias_relu)
Paul's avatar
Add cbr  
Paul committed
780

Paul's avatar
Paul committed
781
template <class... Ms>
Paul's avatar
Add cbr  
Paul committed
782
783
auto conv_bias(Ms... ms)
{
Paul's avatar
Paul committed
784
    return match::name("gpu::add")(
Paul's avatar
Paul committed
785
786
        match::either_arg(0, 1)(bias_shape(match::used_once()).bind("bias"),
                                fusable_conv(match::used_once()).bind("conv")),
Paul's avatar
Paul committed
787
        ms...);
Paul's avatar
Paul committed
788
789
}

Paul's avatar
Paul committed
790
template <class Op>
791
void apply_conv_bias(context& ctx, module& m, const match::matcher_result& r)
Paul's avatar
Paul committed
792
793
794
795
796
797
798
799
800
801
{
    auto conv_ins    = r.instructions["conv"];
    auto bias_ins    = r.instructions["bias"];
    auto ins         = r.result;
    auto input_ins   = conv_ins->inputs().at(0);
    auto weights_ins = conv_ins->inputs().at(1);
    auto conv_op     = any_cast<miopen_convolution>(conv_ins->get_operator()).op;
    auto alloc_ins   = ins->inputs().back();
    auto old_ws_ins  = conv_ins->inputs().at(2);

802
    Op cb{conv_op};
Paul's avatar
Paul committed
803
    // TODO: Insert ws allocation
Paul's avatar
Paul committed
804
    auto ws = cb.get_workspace(ctx);
Paul's avatar
Paul committed
805
    (void)ws;
806
    m.replace_instruction(ins, cb, input_ins, weights_ins, old_ws_ins, bias_ins, alloc_ins);
Paul's avatar
Add cbr  
Paul committed
807
808
}

809
810
811
812
813
814
815
816
817
818
819
820
821
822
823
824
825
826
827
inline auto precompile_name(std::string s) // NOLINT
{
    return match::make_basic_pred_matcher([=](instruction_ref ins) {
        if(ins->name() != "gpu::precompile_op")
            return false;
        auto op = from_value<operation>(ins->get_operator().to_value().at("op"));
        return (op.name() == s);
    });
}

template <class... Ms>
auto conv_bias_pointwise(Ms... ms)
{
    return precompile_name("pointwise")(
        match::either_arg(0, 1)(bias_shape(match::used_once()).bind("bias"),
                                fusable_conv(match::used_once()).bind("conv")),
        ms...);
}

Paul's avatar
Paul committed
828
struct find_conv_bias
Paul's avatar
Paul committed
829
{
Paul's avatar
Paul committed
830
    context* ctx = nullptr;
Paul's avatar
Paul committed
831
832
    auto matcher() const
    {
kahmed10's avatar
kahmed10 committed
833
834
        return conv_bias(match::none_of(
            match::output(match::name(std::unordered_set<std::string>{"gpu::relu"}))));
Paul's avatar
Paul committed
835
836
    }

837
    void apply(module& m, const match::matcher_result& r) const
Paul's avatar
Paul committed
838
    {
839
        apply_conv_bias<miopen_conv_bias>(*ctx, m, r);
Paul's avatar
Paul committed
840
841
842
    }
};

Paul's avatar
Paul committed
843
struct find_conv_bias_relu
Paul's avatar
Add cbr  
Paul committed
844
845
{
    context* ctx = nullptr;
Paul's avatar
Paul committed
846
    auto matcher() const { return match::name("gpu::relu")(match::arg(0)(conv_bias())); }
Paul's avatar
Add cbr  
Paul committed
847

848
    void apply(module& m, const match::matcher_result& r) const
Paul's avatar
Add cbr  
Paul committed
849
    {
850
        apply_conv_bias<miopen_conv_bias_relu>(*ctx, m, r);
Paul's avatar
Add cbr  
Paul committed
851
852
    }
};
853

854
855
856
857
858
859
860
861
862
863
864
struct find_conv_pointwise
{
    context* ctx = nullptr;
    auto matcher() const
    {
        return precompile_name("pointwise")(
            match::nargs(3),
            match::either_arg(0, 1)(bias_shape(match::used_once()).bind("bias"),
                                    fusable_conv(match::used_once()).bind("conv")));
    }

865
    void apply(module& m, const match::matcher_result& r) const
866
867
868
869
870
871
872
873
874
875
876
877
878
879
880
881
882
883
884
885
886
887
888
889
890
891
892
    {
        auto conv_ins    = r.instructions["conv"];
        auto bias_ins    = r.instructions["bias"];
        auto ins         = r.result;
        auto input_ins   = conv_ins->inputs().at(0);
        auto weights_ins = conv_ins->inputs().at(1);
        auto conv_op     = any_cast<miopen_convolution>(conv_ins->get_operator()).op;
        auto alloc_ins   = ins->inputs().back();

        module_ref pm = ins->module_inputs().front();

        miopen_fusion op{};
        op.ops.push_back({{conv_op}});
        for(auto&& i : *pm)
        {
            if(i.name()[0] == '@')
                continue;
            op.ops.push_back({{i.get_operator()}});
        }
        std::vector<instruction_ref> inputs = {input_ins, weights_ins, bias_ins, alloc_ins};
        auto v                              = op.compile(*ctx, ins->get_shape(), to_shapes(inputs));
        if(not v.is_object())
            return;
        m.replace_instruction(ins, op, inputs);
    }
};

893
894
895
896
897
898
899
900
901
902
struct find_gemm_add
{
    auto matcher() const
    {
        return match::name("gpu::add")(
            match::all_of[match::inputs()](match::standard_shape()),
            match::either_arg(0, 1)(match::used_once().bind("c"),
                                    match::name("gpu::gemm")(match::nargs(3)).bind("gemm")));
    }

903
    void apply(module& m, const match::matcher_result& r) const
904
905
906
907
908
909
910
911
    {
        auto ins      = r.result;
        auto gemm_ins = r.instructions["gemm"];
        auto c_ins    = r.instructions["c"];

        auto gemm = any_cast<rocblas_gemm<op::dot>>(gemm_ins->get_operator());

        // Already fused gemm
912
        if(not float_equal(gemm.beta, 0))
913
914
915
916
917
918
919
920
            return;

        auto inputs = gemm_ins->inputs();
        inputs.pop_back();

        auto copy_ins = c_ins;

        // Insert copy
921
        if(ins == m.end() or c_ins->outputs().size() > 1 or c_ins->inputs().empty())
922
        {
923
            copy_ins = m.insert_instruction(ins, hip_copy{}, c_ins, ins->inputs().back());
924
925
926
927
        }
        inputs.push_back(copy_ins);
        inputs.push_back(copy_ins);

928
        gemm.beta = 1;
929
        m.replace_instruction(ins, gemm, inputs);
930
931
932
    }
};

933
934
935
936
937
938
939
940
941
942
943
944
945
946
947
948
949
auto pointwise_name(const std::string& s)
{
    return precompile_name("pointwise")(match::make_basic_pred_matcher([=](auto ins) {
        module_ref pm = ins->module_inputs().front();
        auto n = std::count_if(pm->begin(), pm->end(), [&](auto& i) { return i.name() == s; });
        if(n != 1)
            return false;
        return std::all_of(pm->begin(), pm->end(), [&](auto& i) {
            return starts_with(i.name(), "@") or i.name() == s;
        });
    }));
}

struct find_gemm_pointwise
{
    auto matcher() const
    {
Paul's avatar
Paul committed
950
        return precompile_name("pointwise")(
951
952
953
954
955
956
            match::nargs(3),
            match::all_of[match::inputs()](match::standard_shape()),
            match::either_arg(0, 1)(match::used_once().bind("c"),
                                    match::name("gpu::gemm")(match::nargs(3)).bind("gemm")));
    }

Paul's avatar
Paul committed
957
958
959
960
961
962
963
964
965
966
967
968
969
970
971
972
973
974
975
976
977
978
979
980
981
982
983
984
985
986
987
988
989
990
991
992
993
994
995
996
997
998
999
1000
1001
1002
1003
1004
1005
1006
1007
1008
1009
1010
1011
1012
1013
1014
    // TODO: Move to matcher.hpp
    static auto match_param(const std::string& name)
    {
        return match::make_basic_pred_matcher([=](auto ins) {
            if (ins->name() != "@param")
                return false;
            auto p = any_cast<builtin::param>(ins->get_operator());
            return p.parameter == name;
        });
    }

    template<class M>
    static auto match_mul_const(M m, const std::string& var)
    {
        return match::name("mul")(match::either_arg(0, 1)(match::name("@literal").bind(var), m)).bind(var+"_mul");
    }

    static auto match_add(const std::string& input, const std::string& output)
    {
        auto param = match::name("@param");
        auto add = match::name("add")(match::args(param, param));
        auto inner_mul = match::any_of(
            match_mul_const(match_param(input), "alpha"),
            match_mul_const(match_param(output), "beta")
            );
        auto mul_add = match::name("add")(match::either_arg(0, 1)(inner_mul, param));
        auto add_mul = match_mul_const(add, "gamma");
        return match::name("@return")(match::args(match::any_of(add, mul_add, add_mul)));
    }

    static float get_float(instruction_ref ins)
    {
        return ins->get_literal().at<float>();
    }

    template<class Gemm>
    static bool update_gemm(Gemm& gemm, module_ref pm, unsigned input)
    {
        auto names = pm->get_parameter_names();
        if(names.size() != 2)
            return false;
        std::sort(names.begin(), names.end());
        unsigned output = input == 0 ? 1 : 0;
        auto mr = match::match_instruction(*pm, std::prev(pm->end()), match_add(names[input], names[output]));
        if (mr.result == pm->end())
            return false;
        if (contains(mr.instructions, "alpha_mul"))
            gemm.alpha *= get_float(mr.instructions["alpha"]);
        else if (contains(mr.instructions, "beta_mul"))
            gemm.beta *= get_float(mr.instructions["beta"]);
        else if (contains(mr.instructions, "gamma_mul"))
        {
            gemm.alpha *= get_float(mr.instructions["gamma"]);
            gemm.beta *= get_float(mr.instructions["gamma"]);
        }
        return true;
    }

1015
1016
1017
1018
1019
1020
1021
1022
1023
1024
1025
    void apply(module& m, const match::matcher_result& r) const
    {
        auto ins      = r.result;
        auto gemm_ins = r.instructions["gemm"];
        auto c_ins    = r.instructions["c"];

        auto gemm = any_cast<rocblas_gemm<op::dot>>(gemm_ins->get_operator());

        // Already fused gemm
        if(not float_equal(gemm.beta, 0))
            return;
Paul's avatar
Paul committed
1026
1027
1028
1029
        gemm.beta = 1;

        if (not update_gemm(gemm, ins->module_inputs().front(), ins->inputs().front() == gemm_ins ? 0 : 1))
            return;
1030
1031
1032
1033
1034

        auto inputs = gemm_ins->inputs();
        inputs.pop_back();

        inputs.push_back(c_ins);
1035
        inputs.push_back(ins->inputs().back());
1036
1037
1038
1039
1040

        m.replace_instruction(ins, gemm, inputs);
    }
};

1041
1042
1043
1044
1045
1046
1047
struct find_commutative_broadcast
{
    auto matcher() const
    {
        return match::name("gpu::add", "gpu::mul")(match::arg(1)(match::broadcast_shape()));
    }

1048
    void apply(module& m, const match::matcher_result& r) const
1049
1050
1051
1052
1053
    {
        auto ins  = r.result;
        auto args = ins->inputs();
        move_broadcasted_back(args);

1054
        m.replace_instruction(ins, ins->get_operator(), args);
1055
1056
1057
    }
};

1058
void fuse_ops::apply(module& m) const
Paul's avatar
Paul committed
1059
{
1060
1061
1062
1063
    match::find_matches(m, find_gelu{}, find_gelu_new{fast_math});
    run_passes(m, {dead_code_elimination{}});
    match::find_matches(m, find_triadd{});
    match::find_matches(m,
kahmed10's avatar
kahmed10 committed
1064
                        find_layernorm{},
1065
                        find_conv_pointwise{ctx},
1066
1067
1068
1069
1070
1071
1072
1073
1074
1075
                        find_conv_bias_relu{ctx},
                        find_conv_bias{ctx},
                        find_add_gelu{},
                        find_add_gelu_new{},
                        find_mul_add{},
                        find_mul_add_relu{},
                        find_add_unary{"gpu::relu", hip_add_relu{}, hip_triadd_relu{}},
                        find_add_unary{"gpu::sigmoid", hip_add_sigmoid{}, hip_triadd_sigmoid{}},
                        find_add_unary{"gpu::tanh", hip_add_tanh{}, hip_triadd_tanh{}},
                        find_add_clip{});
1076
    run_passes(m, {dead_code_elimination{}});
1077
1078
1079
1080
1081
    match::find_matches(m,
                        find_triadd_layernorm{},
                        find_gemm_add{},
                        find_gemm_pointwise{},
                        find_commutative_broadcast{});
Paul's avatar
Paul committed
1082
}
Paul's avatar
Paul committed
1083
1084

} // namespace gpu
Paul's avatar
Paul committed
1085
} // namespace MIGRAPHX_INLINE_NS
Paul's avatar
Paul committed
1086
} // namespace migraphx