lowering.cpp 9.14 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{};
43
    std::unordered_map<std::string, std::function<instruction_ref(miopen_apply&, instruction_ref)>> apply_map{};
Paul's avatar
Paul committed
44

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

52
53
54
55
56
57
58
59
60
61
62
63
64
65
66
67
68
69
70
71
    void init()
    {
        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;
    }

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

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

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

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

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

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

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

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

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

        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
130
    }
131

Khalique's avatar
Khalique committed
132
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
    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
159
160
161
162
    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);
163

Khalique's avatar
Khalique committed
164
165
166
167
        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
168

Khalique's avatar
Khalique committed
169
170
171
172
173
174
175
176
177
178
    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
179
180
    instruction_ref apply_softmax(instruction_ref ins)
    {
wsttiger's avatar
wsttiger committed
181
        auto&& op   = any_cast<op::softmax>(ins->get_operator());
Paul's avatar
Paul committed
182
        auto output = insert_allocation(ins, ins->get_shape());
Paul's avatar
Paul committed
183
        return prog->replace_instruction(ins, miopen_softmax{op}, ins->inputs().at(0), output);
Paul's avatar
Paul committed
184
185
    }

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

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

Khalique's avatar
Khalique committed
199
200
201
202
203
204
205
    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);
    }

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

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

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

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