lowering.cpp 8.13 KB
Newer Older
1
#include <rocblas.h>
Paul's avatar
Paul committed
2
#include <migraph/gpu/lowering.hpp>
Paul's avatar
Paul committed
3
4
5
#include <migraph/manage_ptr.hpp>
#include <migraph/instruction.hpp>
#include <migraph/operators.hpp>
Paul's avatar
Paul committed
6
#include <migraph/generate.hpp>
Paul's avatar
Paul committed
7
#include <migraph/shape_for_each.hpp>
Paul's avatar
Paul committed
8
9
#include <migraph/gpu/miopen.hpp>
#include <migraph/gpu/hip.hpp>
Paul's avatar
Paul committed
10
#include <migraph/dfor.hpp>
11
#include <migraph/gpu/device/contiguous.hpp>
Paul's avatar
Paul committed
12
#include <migraph/gpu/device/add.hpp>
Paul's avatar
Paul committed
13
#include <migraph/iterator_for.hpp>
Paul's avatar
Paul committed
14
15
#include <migraph/gpu/rocblas.hpp>
#include <migraph/gpu/context.hpp>
wsttiger's avatar
wsttiger committed
16
#include <migraph/gpu/convolution.hpp>
17
#include <migraph/gpu/contiguous.hpp>
Khalique's avatar
Khalique committed
18
#include <migraph/gpu/lrn.hpp>
19
#include <migraph/gpu/relu.hpp>
Khalique's avatar
Khalique committed
20
#include <migraph/gpu/leaky_relu.hpp>
21
22
#include <migraph/gpu/softmax.hpp>
#include <migraph/gpu/add.hpp>
Khalique's avatar
Khalique committed
23
#include <migraph/gpu/mul.hpp>
24
#include <migraph/gpu/batchnorm.hpp>
wsttiger's avatar
wsttiger committed
25
26
#include <migraph/gpu/pooling.hpp>
#include <migraph/gpu/gemm.hpp>
27
#include <migraph/gpu/concat.hpp>
Paul's avatar
Paul committed
28
#include <utility>
Paul's avatar
Paul committed
29
30

namespace migraph {
31
inline namespace MIGRAPH_INLINE_NS {
Paul's avatar
Paul committed
32
namespace gpu {
Paul's avatar
Paul committed
33
34
35

struct miopen_apply
{
Paul's avatar
Paul committed
36
    program* prog = nullptr;
Paul's avatar
Paul committed
37
    context ctx{};
Paul's avatar
Paul committed
38

Paul's avatar
Paul committed
39
40
41
42
43
44
45
    void check_shape(shape x, instruction_ref i)
    {
        assert(x == i->get_shape());
        (void)x;
        (void)i;
    }

Paul's avatar
Paul committed
46
47
    void apply()
    {
Paul's avatar
Paul committed
48
49
        for(auto it = prog->begin(); it != prog->end(); it++)
        {
Paul's avatar
Paul committed
50
            auto s = it->get_shape();
Paul's avatar
Paul committed
51
            if(it->name() == "convolution")
Paul's avatar
Paul committed
52
            {
Paul's avatar
Paul committed
53
                check_shape(s, apply_convolution(it));
Paul's avatar
Paul committed
54
            }
Khalique's avatar
Khalique committed
55
            else if(it->name() == "relu")
Paul's avatar
Paul committed
56
            {
Khalique's avatar
Khalique committed
57
                check_shape(s, apply_relu(it));
Paul's avatar
Paul committed
58
            }
Khalique's avatar
Khalique committed
59
60
61
62
            else if(it->name() == "leaky_relu")
            {
                check_shape(s, apply_leaky_relu(it));
            }
Paul's avatar
Paul committed
63
            else if(it->name() == "pooling")
Paul's avatar
Paul committed
64
            {
Paul's avatar
Paul committed
65
                check_shape(s, apply_pooling(it));
Paul's avatar
Paul committed
66
            }
Khalique's avatar
Khalique committed
67
            else if(it->name() == "lrn")
Khalique's avatar
Khalique committed
68
            {
Khalique's avatar
Khalique committed
69
                check_shape(s, apply_lrn(it));
Khalique's avatar
Khalique committed
70
            }
Paul's avatar
Paul committed
71
            else if(it->name() == "add")
Paul's avatar
Paul committed
72
            {
Paul's avatar
Paul committed
73
                check_shape(s, apply_add(it));
Paul's avatar
Paul committed
74
            }
Khalique's avatar
Khalique committed
75
76
77
78
            else if(it->name() == "mul")
            {
                check_shape(s, apply_mul(it));
            }
Shucai Xiao's avatar
Shucai Xiao committed
79
            else if(it->name() == "dot")
Paul's avatar
Paul committed
80
            {
Paul's avatar
Paul committed
81
                check_shape(s, apply_gemm(it));
Paul's avatar
Paul committed
82
            }
Paul's avatar
Paul committed
83
            else if(it->name() == "contiguous")
84
            {
Paul's avatar
Paul committed
85
                check_shape(s, apply_contiguous(it));
86
            }
87
88
89
90
            else if(it->name() == "concat")
            {
                check_shape(s, apply_concat(it));
            }
Paul's avatar
Paul committed
91
            else if(it->name() == "batch_norm_inference")
92
            {
Paul's avatar
Paul committed
93
                check_shape(s, apply_batch_norm_inference(it));
94
            }
Paul's avatar
Paul committed
95
96
97
98
            else if(it->name() == "softmax")
            {
                check_shape(s, apply_softmax(it));
            }
Paul's avatar
Paul committed
99
100
101
        }
    }

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

Paul's avatar
Paul committed
116
    instruction_ref apply_convolution(instruction_ref ins)
Paul's avatar
Paul committed
117
    {
wsttiger's avatar
wsttiger committed
118
        auto&& op = any_cast<op::convolution>(ins->get_operator());
Paul's avatar
Paul committed
119

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

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

Paul's avatar
Paul committed
126
        return prog->replace_instruction(
Paul's avatar
Paul committed
127
            ins, conv, ins->inputs().at(0), ins->inputs().at(1), workspace, output);
Paul's avatar
Paul committed
128
129
    }

Paul's avatar
Paul committed
130
    instruction_ref apply_pooling(instruction_ref ins)
Paul's avatar
Paul committed
131
    {
wsttiger's avatar
wsttiger committed
132
        auto&& op   = any_cast<op::pooling>(ins->get_operator());
Paul's avatar
Paul committed
133
        auto pd     = make_pooling(op);
Paul's avatar
Paul committed
134
        auto output = insert_allocation(ins, ins->get_shape());
Paul's avatar
Paul committed
135

Paul's avatar
Paul committed
136
        return prog->replace_instruction(
Paul's avatar
Paul committed
137
            ins, miopen_pooling{op, std::move(pd)}, ins->inputs().at(0), output);
Paul's avatar
Paul committed
138
139
    }

Khalique's avatar
Khalique committed
140
    instruction_ref apply_lrn(instruction_ref ins)
Khalique's avatar
Khalique committed
141
    {
Khalique's avatar
Khalique committed
142
143
        auto&& op   = any_cast<op::lrn>(ins->get_operator());
        auto ldesc  = make_lrn(op);
Khalique's avatar
Khalique committed
144
        auto output = insert_allocation(ins, ins->get_shape());
Khalique's avatar
Khalique committed
145
        return prog->replace_instruction(
Khalique's avatar
Khalique committed
146
            ins, miopen_lrn{std::move(ldesc)}, ins->inputs().at(0), output);
Khalique's avatar
Khalique committed
147
148
    }

Khalique's avatar
Khalique committed
149
    instruction_ref apply_relu(instruction_ref ins)
Paul's avatar
Paul committed
150
    {
Khalique's avatar
Khalique committed
151
        auto ad = make_relu();
Khalique's avatar
Khalique committed
152
153
154
155

        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
156
    }
157

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

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

Paul's avatar
Paul committed
168
169
    instruction_ref apply_softmax(instruction_ref ins)
    {
wsttiger's avatar
wsttiger committed
170
        auto&& op   = any_cast<op::softmax>(ins->get_operator());
Paul's avatar
Paul committed
171
        auto output = insert_allocation(ins, ins->get_shape());
Paul's avatar
Paul committed
172
        return prog->replace_instruction(ins, miopen_softmax{op}, ins->inputs().at(0), output);
Paul's avatar
Paul committed
173
174
    }

Paul's avatar
Paul committed
175
    instruction_ref apply_add(instruction_ref ins)
Paul's avatar
Paul committed
176
    {
Paul's avatar
Paul committed
177
        auto output = insert_allocation(ins, ins->get_shape());
Paul's avatar
Paul committed
178
        return prog->replace_instruction(
Paul's avatar
Paul committed
179
            ins, hip_add{}, ins->inputs().at(0), ins->inputs().at(1), output);
Paul's avatar
Paul committed
180
    }
Paul's avatar
Paul committed
181

Khalique's avatar
Khalique committed
182
183
184
185
186
187
188
    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);
    }

Paul's avatar
Paul committed
189
    instruction_ref apply_gemm(instruction_ref ins)
Paul's avatar
Paul committed
190
    {
Shucai Xiao's avatar
Shucai Xiao committed
191
        auto&& op   = any_cast<op::dot>(ins->get_operator());
Paul's avatar
Paul committed
192
        auto output = insert_allocation(ins, ins->get_shape());
Paul's avatar
Paul committed
193
        return prog->replace_instruction(
Paul's avatar
Paul committed
194
            ins, miopen_gemm{op}, ins->inputs().at(0), ins->inputs().at(1), output);
Paul's avatar
Paul committed
195
    }
196

Paul's avatar
Paul committed
197
    instruction_ref apply_contiguous(instruction_ref ins)
198
    {
wsttiger's avatar
wsttiger committed
199
        auto&& op   = any_cast<op::contiguous>(ins->get_operator());
Paul's avatar
Paul committed
200
        auto output = insert_allocation(ins, ins->get_shape());
Paul's avatar
Paul committed
201
        return prog->replace_instruction(ins, miopen_contiguous{op}, ins->inputs().at(0), output);
202
    }
203

204
205
    instruction_ref apply_concat(instruction_ref ins)
    {
wsttiger's avatar
wsttiger committed
206
207
        auto&& op                         = any_cast<op::concat>(ins->get_operator());
        auto output                       = insert_allocation(ins, ins->get_shape());
208
209
210
211
212
        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
213
    instruction_ref apply_batch_norm_inference(instruction_ref ins)
214
    {
wsttiger's avatar
wsttiger committed
215
        auto&& op       = any_cast<op::batch_norm_inference>(ins->get_operator());
Paul's avatar
Paul committed
216
        auto output     = insert_allocation(ins, ins->get_shape());
Paul's avatar
Paul committed
217
        shape old_shape = ins->inputs().at(1)->get_shape();
wsttiger's avatar
wsttiger committed
218
        std::vector<int64_t> new_shape{1, static_cast<int64_t>(old_shape.elements()), 1, 1};
wsttiger's avatar
wsttiger committed
219
        auto reshape_op = op::reshape{new_shape};
Paul's avatar
Paul committed
220
        std::vector<instruction_ref> reshapes;
Paul's avatar
Paul committed
221
222
        std::transform(ins->inputs().begin() + 1,
                       ins->inputs().end(),
Paul's avatar
Paul committed
223
224
                       std::back_inserter(reshapes),
                       [&](auto i) { return prog->insert_instruction(ins, reshape_op, i); });
Paul's avatar
Paul committed
225
        return prog->replace_instruction(ins,
Paul's avatar
Paul committed
226
                                         miopen_batch_norm_inference{op},
Paul's avatar
Paul committed
227
                                         ins->inputs().at(0),
Paul's avatar
Paul committed
228
229
230
231
232
                                         reshapes[0],
                                         reshapes[1],
                                         reshapes[2],
                                         reshapes[3],
                                         output);
233
    }
Paul's avatar
Paul committed
234
235
};

Paul's avatar
Paul committed
236
void lowering::apply(program& p) const { miopen_apply{&p, ctx}.apply(); }
Paul's avatar
Paul committed
237
} // namespace gpu
238
} // namespace MIGRAPH_INLINE_NS
Paul's avatar
Paul committed
239
} // namespace migraph