lowering.cpp 9.34 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
21
#include <migraphx/gpu/sigmoid.hpp>
#include <migraphx/gpu/tanh.hpp>
#include <migraphx/gpu/abs.hpp>
Paul's avatar
Paul committed
22
#include <migraphx/gpu/leaky_relu.hpp>
Khalique's avatar
Khalique committed
23
#include <migraphx/gpu/elu.hpp>
Paul's avatar
Paul committed
24
25
#include <migraphx/gpu/softmax.hpp>
#include <migraphx/gpu/add.hpp>
26
#include <migraphx/gpu/sin.hpp>
Paul's avatar
Paul committed
27
28
29
30
31
#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
32
#include <utility>
33
#include <functional>
Paul's avatar
Paul committed
34

Paul's avatar
Paul committed
35
namespace migraphx {
36
inline namespace MIGRAPH_INLINE_NS {
Paul's avatar
Paul committed
37
namespace gpu {
Paul's avatar
Paul committed
38
39
40

struct miopen_apply
{
Paul's avatar
Paul committed
41
    program* prog = nullptr;
Paul's avatar
Paul committed
42
    context ctx{};
Khalique's avatar
Khalique committed
43
44
    std::unordered_map<std::string, std::function<instruction_ref(miopen_apply&, instruction_ref)>>
        apply_map{};
Paul's avatar
Paul committed
45

Paul's avatar
Paul committed
46
47
48
49
50
51
52
    void check_shape(shape x, instruction_ref i)
    {
        assert(x == i->get_shape());
        (void)x;
        (void)i;
    }

53
54
    void init()
    {
Khalique's avatar
Khalique committed
55
56
57
58
59
60
61
62
63
64
65
66
67
68
69
70
        apply_map["convolution"]          = &miopen_apply::apply_convolution;
        apply_map["relu"]                 = &miopen_apply::apply_relu;
        apply_map["sigmoid"]              = &miopen_apply::apply_sigmoid;
        apply_map["tanh"]                 = &miopen_apply::apply_tanh;
        apply_map["abs"]                  = &miopen_apply::apply_abs;
        apply_map["leaky_relu"]           = &miopen_apply::apply_leaky_relu;
        apply_map["elu"]                  = &miopen_apply::apply_elu;
        apply_map["pooling"]              = &miopen_apply::apply_pooling;
        apply_map["add"]                  = &miopen_apply::apply_add;
        apply_map["sin"]                  = &miopen_apply::apply_sin;
        apply_map["mul"]                  = &miopen_apply::apply_mul;
        apply_map["dot"]                  = &miopen_apply::apply_dot;
        apply_map["contiguous"]           = &miopen_apply::apply_contiguous;
        apply_map["concat"]               = &miopen_apply::apply_concat;
        apply_map["batch_norm_inference"] = &miopen_apply::apply_batch_norm_inference;
        apply_map["softmax"]              = &miopen_apply::apply_softmax;
71
72
    }

Paul's avatar
Paul committed
73
74
    void apply()
    {
75
        init();
Paul's avatar
Paul committed
76
77
        for(auto it = prog->begin(); it != prog->end(); it++)
        {
Paul's avatar
Paul committed
78
            auto s = it->get_shape();
79
            if(apply_map.count(it->name()) > 0)
Paul's avatar
Paul committed
80
            {
81
                check_shape(s, apply_map.at(it->name())(*this, it));
Paul's avatar
Paul committed
82
            }
Paul's avatar
Paul committed
83
84
85
        }
    }

Paul's avatar
Paul committed
86
    instruction_ref insert_allocation(instruction_ref ins, const shape& s, std::string tag = "")
Paul's avatar
Paul committed
87
    {
Paul's avatar
Paul committed
88
        if(ins == --prog->end() and tag.empty())
Paul's avatar
Paul committed
89
90
91
92
93
        {
            return prog->add_parameter("output", s);
        }
        else
        {
Paul's avatar
Paul committed
94
            auto is     = prog->add_outline(s);
Paul's avatar
Paul committed
95
            auto result = prog->insert_instruction(ins, hip_allocate{std::move(tag)}, is);
Paul's avatar
Paul committed
96
97
98
99
            return result;
        }
    }

Paul's avatar
Paul committed
100
    instruction_ref apply_convolution(instruction_ref ins)
Paul's avatar
Paul committed
101
    {
wsttiger's avatar
wsttiger committed
102
        auto&& op = any_cast<op::convolution>(ins->get_operator());
Paul's avatar
Paul committed
103

Paul's avatar
Paul committed
104
        auto conv = miopen_convolution{op, make_conv(op)};
Paul's avatar
Paul committed
105
        auto ws   = conv.compile(ctx, ins->get_shape(), ins->inputs());
Paul's avatar
Paul committed
106

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

Paul's avatar
Paul committed
110
        return prog->replace_instruction(
Paul's avatar
Paul committed
111
            ins, conv, ins->inputs().at(0), ins->inputs().at(1), workspace, output);
Paul's avatar
Paul committed
112
113
    }

Paul's avatar
Paul committed
114
    instruction_ref apply_pooling(instruction_ref ins)
Paul's avatar
Paul committed
115
    {
wsttiger's avatar
wsttiger committed
116
        auto&& op   = any_cast<op::pooling>(ins->get_operator());
Paul's avatar
Paul committed
117
        auto pd     = make_pooling(op);
Paul's avatar
Paul committed
118
        auto output = insert_allocation(ins, ins->get_shape());
Paul's avatar
Paul committed
119

Paul's avatar
Paul committed
120
        return prog->replace_instruction(
Paul's avatar
Paul committed
121
            ins, miopen_pooling{op, std::move(pd)}, ins->inputs().at(0), output);
Paul's avatar
Paul committed
122
123
    }

Khalique's avatar
Khalique committed
124
    instruction_ref apply_relu(instruction_ref ins)
Paul's avatar
Paul committed
125
    {
Khalique's avatar
Khalique committed
126
        auto ad = make_relu();
Khalique's avatar
Khalique committed
127
128
129
130

        auto output = insert_allocation(ins, ins->get_shape());
        return prog->replace_instruction(
            ins, miopen_relu{std::move(ad)}, ins->inputs().at(0), output);
Paul's avatar
Paul committed
131
    }
132

Khalique's avatar
Khalique committed
133
134
135
136
137
138
139
140
141
142
143
144
145
146
147
148
149
150
151
152
153
154
155
156
157
158
159
    instruction_ref apply_sigmoid(instruction_ref ins)
    {
        auto ad = make_sigmoid();

        auto output = insert_allocation(ins, ins->get_shape());
        return prog->replace_instruction(
            ins, miopen_sigmoid{std::move(ad)}, ins->inputs().at(0), output);
    }

    instruction_ref apply_tanh(instruction_ref ins)
    {
        auto ad = make_tanh();

        auto output = insert_allocation(ins, ins->get_shape());
        return prog->replace_instruction(
            ins, miopen_tanh{std::move(ad)}, ins->inputs().at(0), output);
    }

    instruction_ref apply_abs(instruction_ref ins)
    {
        auto ad = make_abs();

        auto output = insert_allocation(ins, ins->get_shape());
        return prog->replace_instruction(
            ins, miopen_abs{std::move(ad)}, ins->inputs().at(0), output);
    }

Khalique's avatar
Khalique committed
160
161
162
163
    instruction_ref apply_leaky_relu(instruction_ref ins)
    {
        auto&& op = any_cast<op::leaky_relu>(ins->get_operator());
        auto ad   = make_leaky_relu(op.alpha);
164

Khalique's avatar
Khalique committed
165
166
167
168
        auto output = insert_allocation(ins, ins->get_shape());
        return prog->replace_instruction(
            ins, miopen_leaky_relu{std::move(ad)}, ins->inputs().at(0), output);
    }
Paul's avatar
Paul committed
169

Khalique's avatar
Khalique committed
170
171
172
173
174
175
176
177
178
179
    instruction_ref apply_elu(instruction_ref ins)
    {
        auto&& op = any_cast<op::leaky_relu>(ins->get_operator());
        auto ad   = make_elu(op.alpha);

        auto output = insert_allocation(ins, ins->get_shape());
        return prog->replace_instruction(
            ins, miopen_elu{std::move(ad)}, ins->inputs().at(0), output);
    }

Paul's avatar
Paul committed
180
181
    instruction_ref apply_softmax(instruction_ref ins)
    {
wsttiger's avatar
wsttiger committed
182
        auto&& op   = any_cast<op::softmax>(ins->get_operator());
Paul's avatar
Paul committed
183
        auto output = insert_allocation(ins, ins->get_shape());
Paul's avatar
Paul committed
184
        return prog->replace_instruction(ins, miopen_softmax{op}, ins->inputs().at(0), output);
Paul's avatar
Paul committed
185
186
    }

Paul's avatar
Paul committed
187
    instruction_ref apply_add(instruction_ref ins)
Paul's avatar
Paul committed
188
    {
Paul's avatar
Paul committed
189
        auto output = insert_allocation(ins, ins->get_shape());
Paul's avatar
Paul committed
190
        return prog->replace_instruction(
Paul's avatar
Paul committed
191
            ins, hip_add{}, ins->inputs().at(0), ins->inputs().at(1), output);
Paul's avatar
Paul committed
192
    }
Paul's avatar
Paul committed
193

194
195
196
    instruction_ref apply_sin(instruction_ref ins)
    {
        auto output = insert_allocation(ins, ins->get_shape());
Shucai Xiao's avatar
Shucai Xiao committed
197
        return prog->replace_instruction(ins, hip_sin{}, ins->inputs().at(0), output);
198
199
    }

Khalique's avatar
Khalique committed
200
201
202
203
204
205
206
    instruction_ref apply_mul(instruction_ref ins)
    {
        auto output = insert_allocation(ins, ins->get_shape());
        return prog->replace_instruction(
            ins, hip_mul{}, ins->inputs().at(0), ins->inputs().at(1), output);
    }

207
    instruction_ref apply_dot(instruction_ref ins)
Paul's avatar
Paul committed
208
    {
Shucai Xiao's avatar
Shucai Xiao committed
209
        auto&& op   = any_cast<op::dot>(ins->get_operator());
Paul's avatar
Paul committed
210
        auto output = insert_allocation(ins, ins->get_shape());
Paul's avatar
Paul committed
211
        return prog->replace_instruction(
Paul's avatar
Paul committed
212
            ins, miopen_gemm{op}, ins->inputs().at(0), ins->inputs().at(1), output);
Paul's avatar
Paul committed
213
    }
214

Paul's avatar
Paul committed
215
    instruction_ref apply_contiguous(instruction_ref ins)
216
    {
wsttiger's avatar
wsttiger committed
217
        auto&& op   = any_cast<op::contiguous>(ins->get_operator());
Paul's avatar
Paul committed
218
        auto output = insert_allocation(ins, ins->get_shape());
Paul's avatar
Paul committed
219
        return prog->replace_instruction(ins, miopen_contiguous{op}, ins->inputs().at(0), output);
220
    }
221

222
223
    instruction_ref apply_concat(instruction_ref ins)
    {
wsttiger's avatar
wsttiger committed
224
225
        auto&& op                         = any_cast<op::concat>(ins->get_operator());
        auto output                       = insert_allocation(ins, ins->get_shape());
226
227
228
229
230
        std::vector<instruction_ref> refs = ins->inputs();
        refs.push_back(output);
        return prog->replace_instruction(ins, hip_concat{op}, refs);
    }

Paul's avatar
Paul committed
231
    instruction_ref apply_batch_norm_inference(instruction_ref ins)
232
    {
wsttiger's avatar
wsttiger committed
233
        auto&& op       = any_cast<op::batch_norm_inference>(ins->get_operator());
Paul's avatar
Paul committed
234
        auto output     = insert_allocation(ins, ins->get_shape());
Paul's avatar
Paul committed
235
        shape old_shape = ins->inputs().at(1)->get_shape();
wsttiger's avatar
wsttiger committed
236
        std::vector<int64_t> new_shape{1, static_cast<int64_t>(old_shape.elements()), 1, 1};
wsttiger's avatar
wsttiger committed
237
        auto reshape_op = op::reshape{new_shape};
Paul's avatar
Paul committed
238
        std::vector<instruction_ref> reshapes;
Paul's avatar
Paul committed
239
240
        std::transform(ins->inputs().begin() + 1,
                       ins->inputs().end(),
Paul's avatar
Paul committed
241
242
                       std::back_inserter(reshapes),
                       [&](auto i) { return prog->insert_instruction(ins, reshape_op, i); });
Paul's avatar
Paul committed
243
        return prog->replace_instruction(ins,
Paul's avatar
Paul committed
244
                                         miopen_batch_norm_inference{op},
Paul's avatar
Paul committed
245
                                         ins->inputs().at(0),
Paul's avatar
Paul committed
246
247
248
249
250
                                         reshapes[0],
                                         reshapes[1],
                                         reshapes[2],
                                         reshapes[3],
                                         output);
251
    }
Paul's avatar
Paul committed
252
253
};

Paul's avatar
Paul committed
254
void lowering::apply(program& p) const { miopen_apply{&p, ctx}.apply(); }
Paul's avatar
Paul committed
255
} // namespace gpu
256
} // namespace MIGRAPH_INLINE_NS
Paul's avatar
Paul committed
257
} // namespace migraphx