lowering.cpp 7.9 KB
Newer Older
1
#include <rocblas.h>
Paul's avatar
Paul committed
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
#include <migraphx/gpu/lowering.hpp>
#include <migraphx/manage_ptr.hpp>
#include <migraphx/instruction.hpp>
#include <migraphx/operators.hpp>
#include <migraphx/generate.hpp>
#include <migraphx/shape_for_each.hpp>
#include <migraphx/gpu/miopen.hpp>
#include <migraphx/gpu/hip.hpp>
#include <migraphx/dfor.hpp>
#include <migraphx/gpu/device/contiguous.hpp>
#include <migraphx/gpu/device/add.hpp>
#include <migraphx/iterator_for.hpp>
#include <migraphx/gpu/rocblas.hpp>
#include <migraphx/gpu/context.hpp>
#include <migraphx/gpu/convolution.hpp>
#include <migraphx/gpu/contiguous.hpp>
#include <migraphx/gpu/relu.hpp>
Khalique's avatar
Khalique committed
19
20
#include <migraphx/gpu/sigmoid.hpp>
#include <migraphx/gpu/abs.hpp>
Paul's avatar
Paul committed
21
#include <migraphx/gpu/leaky_relu.hpp>
Khalique's avatar
Khalique committed
22
#include <migraphx/gpu/elu.hpp>
Paul's avatar
Paul committed
23
24
#include <migraphx/gpu/softmax.hpp>
#include <migraphx/gpu/add.hpp>
25
#include <migraphx/gpu/sin.hpp>
Shucai Xiao's avatar
Shucai Xiao committed
26
27
#include <migraphx/gpu/cos.hpp>
#include <migraphx/gpu/tan.hpp>
28
29
#include <migraphx/gpu/sinh.hpp>
#include <migraphx/gpu/cosh.hpp>
30
#include <migraphx/gpu/tanh.hpp>
31
32
33
#include <migraphx/gpu/asin.hpp>
#include <migraphx/gpu/acos.hpp>
#include <migraphx/gpu/atan.hpp>
Paul's avatar
Paul committed
34
35
36
37
38
#include <migraphx/gpu/mul.hpp>
#include <migraphx/gpu/batchnorm.hpp>
#include <migraphx/gpu/pooling.hpp>
#include <migraphx/gpu/gemm.hpp>
#include <migraphx/gpu/concat.hpp>
Paul's avatar
Paul committed
39
#include <utility>
40
#include <functional>
Paul's avatar
Paul committed
41

Paul's avatar
Paul committed
42
namespace migraphx {
Paul's avatar
Paul committed
43
inline namespace MIGRAPHX_INLINE_NS {
Paul's avatar
Paul committed
44
namespace gpu {
Paul's avatar
Paul committed
45
46
47

struct miopen_apply
{
Paul's avatar
Paul committed
48
    program* prog = nullptr;
Paul's avatar
Paul committed
49
    context ctx{};
Shucai Xiao's avatar
Shucai Xiao committed
50
    std::unordered_map<std::string, std::function<instruction_ref(instruction_ref)>> apply_map{};
Paul's avatar
Paul committed
51

Paul's avatar
Paul committed
52
53
54
55
56
57
58
    void check_shape(shape x, instruction_ref i)
    {
        assert(x == i->get_shape());
        (void)x;
        (void)i;
    }

59
60
    void init()
    {
61
62
63
64
65
66
67
68
69
70
71
72
73
74
75
76
77
78
79
80
81
82
83
84
85
86
87
        add_miopen_simple_op("relu", miopen_relu{}, make_relu);
        add_miopen_simple_op("sigmoid", miopen_sigmoid{}, make_sigmoid);
        add_miopen_simple_op("abs", miopen_abs{}, make_abs);
        add_miopen_simple_op("tanh", miopen_tanh{}, make_tanh);

        add_miopen_extend_op("leaky_relu", miopen_leaky_relu{}, op::leaky_relu{}, make_leaky_relu);
        add_miopen_extend_op("elu", miopen_elu{}, op::elu{}, make_elu);

        add_generic_op("add", hip_add{});
        add_generic_op("sin", hip_sin{});
        add_generic_op("cos", hip_cos{});
        add_generic_op("tan", hip_tan{});
        add_generic_op("sinh", hip_sinh{});
        add_generic_op("cosh", hip_cosh{});
        add_generic_op("asin", hip_asin{});
        add_generic_op("acos", hip_acos{});
        add_generic_op("atan", hip_atan{});
        add_generic_op("mul", hip_mul{});

        add_extend_op("dot", miopen_gemm{}, op::dot{});
        add_extend_op("contiguous", miopen_contiguous{}, op::contiguous{});
        add_extend_op("concat", hip_concat{}, op::concat{});
        add_extend_op("softmax", miopen_softmax{}, op::softmax{});

        add_convolution_op();
        add_pooling_op();
        add_batch_norm_inference_op();
88
89
    }

Paul's avatar
Paul committed
90
91
    void apply()
    {
92
        init();
Paul's avatar
Paul committed
93
94
        for(auto it = prog->begin(); it != prog->end(); it++)
        {
Paul's avatar
Paul committed
95
            auto s = it->get_shape();
96
            if(apply_map.count(it->name()) > 0)
Paul's avatar
Paul committed
97
            {
98
                check_shape(s, apply_map.at(it->name())(it));
Paul's avatar
Paul committed
99
            }
Paul's avatar
Paul committed
100
101
102
        }
    }

Paul's avatar
Paul committed
103
    instruction_ref insert_allocation(instruction_ref ins, const shape& s, std::string tag = "")
Paul's avatar
Paul committed
104
    {
Paul's avatar
Paul committed
105
        if(ins == --prog->end() and tag.empty())
Paul's avatar
Paul committed
106
107
108
109
110
        {
            return prog->add_parameter("output", s);
        }
        else
        {
Paul's avatar
Paul committed
111
            auto is     = prog->add_outline(s);
Paul's avatar
Paul committed
112
            auto result = prog->insert_instruction(ins, hip_allocate{std::move(tag)}, is);
Paul's avatar
Paul committed
113
114
115
116
            return result;
        }
    }

Shucai Xiao's avatar
Shucai Xiao committed
117
118
    void add_convolution_op()
    {
119
120
        apply_map.emplace("convolution", [=](instruction_ref ins) {
            auto&& op = any_cast<op::convolution>(ins->get_operator());
Khalique's avatar
Khalique committed
121

122
123
            auto conv = miopen_convolution{op, make_conv(op)};
            auto ws   = conv.compile(ctx, ins->get_shape(), ins->inputs());
Khalique's avatar
Khalique committed
124

125
126
            auto workspace = insert_allocation(ins, ws, "workspace");
            auto output    = insert_allocation(ins, ins->get_shape());
Khalique's avatar
Khalique committed
127

128
129
130
            return prog->replace_instruction(
                ins, conv, ins->inputs().at(0), ins->inputs().at(1), workspace, output);
        });
Khalique's avatar
Khalique committed
131
132
    }

Shucai Xiao's avatar
Shucai Xiao committed
133
134
    void add_pooling_op()
    {
135
136
137
138
        apply_map.emplace("pooling", [=](instruction_ref ins) {
            auto&& op   = any_cast<op::pooling>(ins->get_operator());
            auto pd     = make_pooling(op);
            auto output = insert_allocation(ins, ins->get_shape());
Khalique's avatar
Khalique committed
139

140
141
142
            return prog->replace_instruction(
                ins, miopen_pooling{op, std::move(pd)}, ins->inputs().at(0), output);
        });
Khalique's avatar
Khalique committed
143
144
    }

Shucai Xiao's avatar
Shucai Xiao committed
145
    template <class T>
146
    void add_generic_op(std::string name, T x)
Khalique's avatar
Khalique committed
147
    {
148
149
150
151
        apply_map.emplace(name, [=](instruction_ref ins) {
            auto output                       = insert_allocation(ins, ins->get_shape());
            std::vector<instruction_ref> refs = ins->inputs();
            refs.push_back(output);
Paul's avatar
Paul committed
152

153
154
155
            return prog->replace_instruction(ins, T{}, refs);
        });
        (void)x;
156
157
    }

Shucai Xiao's avatar
Shucai Xiao committed
158
    template <class T, class Op>
159
    void add_extend_op(std::string name, T x, Op o)
160
    {
161
162
163
164
165
        apply_map.emplace(name, [=](instruction_ref ins) {
            auto&& op                         = any_cast<Op>(ins->get_operator());
            auto output                       = insert_allocation(ins, ins->get_shape());
            std::vector<instruction_ref> refs = ins->inputs();
            refs.push_back(output);
166

167
168
169
170
            return prog->replace_instruction(ins, T{op}, refs);
        });
        (void)x;
        (void)o;
171
    }
172

Shucai Xiao's avatar
Shucai Xiao committed
173
174
175
176
    template <class T, class Op, class F>
    void add_miopen_extend_op(std::string name, T x, Op o, F f)
    {
        apply_map.emplace(name, [=](instruction_ref ins) {
177
178
            auto&& op = any_cast<Op>(ins->get_operator());
            auto ad   = f(op.alpha);
179

180
            auto output = insert_allocation(ins, ins->get_shape());
Shucai Xiao's avatar
Shucai Xiao committed
181
            return prog->replace_instruction(ins, T{std::move(ad)}, ins->inputs().at(0), output);
182
183
184
185
186
187
        });
        (void)x;
        (void)o;
        (void)f;
    }

Shucai Xiao's avatar
Shucai Xiao committed
188
189
190
191
192
    template <class T, class F>
    void add_miopen_simple_op(std::string name, T x, F f)
    {
        apply_map.emplace(name, [=](instruction_ref ins) {
            auto ad     = f();
193
            auto output = insert_allocation(ins, ins->get_shape());
Shucai Xiao's avatar
Shucai Xiao committed
194
            return prog->replace_instruction(ins, T{std::move(ad)}, ins->inputs().at(0), output);
195
196
197
198
199
        });
        (void)x;
        (void)f;
    }

Shucai Xiao's avatar
Shucai Xiao committed
200
201
    void add_batch_norm_inference_op()
    {
202
203
204
205
206
207
208
209
        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();
            std::vector<int64_t> new_shape{1, static_cast<int64_t>(old_shape.elements()), 1, 1};
            auto reshape_op = op::reshape{new_shape};
            std::vector<instruction_ref> reshapes;
            std::transform(ins->inputs().begin() + 1,
Shucai Xiao's avatar
Shucai Xiao committed
210
211
212
                           ins->inputs().end(),
                           std::back_inserter(reshapes),
                           [&](auto i) { return prog->insert_instruction(ins, reshape_op, i); });
213
            return prog->replace_instruction(ins,
Shucai Xiao's avatar
Shucai Xiao committed
214
215
216
217
218
219
220
                                             miopen_batch_norm_inference{op},
                                             ins->inputs().at(0),
                                             reshapes[0],
                                             reshapes[1],
                                             reshapes[2],
                                             reshapes[3],
                                             output);
221
        });
222
    }
Paul's avatar
Paul committed
223
224
};

Paul's avatar
Paul committed
225
void lowering::apply(program& p) const { miopen_apply{&p, ctx}.apply(); }
Paul's avatar
Paul committed
226
} // namespace gpu
Paul's avatar
Paul committed
227
} // namespace MIGRAPHX_INLINE_NS
Paul's avatar
Paul committed
228
} // namespace migraphx