lowering.cpp 7.91 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 {
43
inline namespace MIGRAPH_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{};
50
    std::unordered_map<std::string, std::function<instruction_ref(instruction_ref)>>
Khalique's avatar
Khalique committed
51
        apply_map{};
Paul's avatar
Paul committed
52

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

60
61
    void init()
    {
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
88
        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();
89
90
    }

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

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

118
119
120
    void add_convolution_op() {
        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
    }

133
134
135
136
137
    void add_pooling_op() {
        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
138

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

144
145
    template<class T>
    void add_generic_op(std::string name, T x)
Khalique's avatar
Khalique committed
146
    {
147
148
149
150
        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
151

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

157
158
    template<class T, class Op>
    void add_extend_op(std::string name, T x, Op o)
159
    {
160
161
162
163
164
        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);
165

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

172
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) {
            auto&& op = any_cast<Op>(ins->get_operator());
            auto ad   = f(op.alpha);
177

178
179
180
181
182
183
184
185
186
187
188
189
190
191
192
193
194
195
196
197
198
199
200
201
202
203
204
205
206
207
208
209
210
211
212
213
214
215
216
217
218
219
            auto output = insert_allocation(ins, ins->get_shape());
            return prog->replace_instruction(
                ins, T{std::move(ad)}, ins->inputs().at(0), output);
        });
        (void)x;
        (void)o;
        (void)f;
    }

    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();
            auto output = insert_allocation(ins, ins->get_shape());
            return prog->replace_instruction(
                ins, T{std::move(ad)}, ins->inputs().at(0), output);
        });
        (void)x;
        (void)f;
    }

    void add_batch_norm_inference_op() {
        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,
                        ins->inputs().end(),
                        std::back_inserter(reshapes),
                        [&](auto i) { return prog->insert_instruction(ins, reshape_op, i); });
            return prog->replace_instruction(ins,
                                            miopen_batch_norm_inference{op},
                                            ins->inputs().at(0),
                                            reshapes[0],
                                            reshapes[1],
                                            reshapes[2],
                                            reshapes[3],
                                            output);
        });
220
    }
Paul's avatar
Paul committed
221
222
};

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