lowering.cpp 6.69 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
18
#include <migraph/gpu/contiguous.hpp>
#include <migraph/gpu/relu.hpp>
Khalique's avatar
Khalique committed
19
#include <migraph/gpu/leaky_relu.hpp>
20
21
22
#include <migraph/gpu/softmax.hpp>
#include <migraph/gpu/add.hpp>
#include <migraph/gpu/batchnorm.hpp>
wsttiger's avatar
wsttiger committed
23
24
#include <migraph/gpu/pooling.hpp>
#include <migraph/gpu/gemm.hpp>
Paul's avatar
Paul committed
25
#include <utility>
Paul's avatar
Paul committed
26
27

namespace migraph {
Paul's avatar
Paul committed
28
namespace gpu {
Paul's avatar
Paul committed
29
30
31

struct miopen_apply
{
Paul's avatar
Paul committed
32
    program* prog = nullptr;
Paul's avatar
Paul committed
33
    context ctx{};
Paul's avatar
Paul committed
34

Paul's avatar
Paul committed
35
36
37
38
39
40
41
    void check_shape(shape x, instruction_ref i)
    {
        assert(x == i->get_shape());
        (void)x;
        (void)i;
    }

Paul's avatar
Paul committed
42
43
    void apply()
    {
Paul's avatar
Paul committed
44
45
        for(auto it = prog->begin(); it != prog->end(); it++)
        {
Paul's avatar
Paul committed
46
            auto s = it->get_shape();
Paul's avatar
Paul committed
47
            if(it->name() == "convolution")
Paul's avatar
Paul committed
48
            {
Paul's avatar
Paul committed
49
                check_shape(s, apply_convolution(it));
Paul's avatar
Paul committed
50
            }
Paul's avatar
Paul committed
51
            else if(it->name() == "activation")
Paul's avatar
Paul committed
52
            {
Paul's avatar
Paul committed
53
                check_shape(s, apply_activation(it));
Paul's avatar
Paul committed
54
            }
Paul's avatar
Paul committed
55
            else if(it->name() == "pooling")
Paul's avatar
Paul committed
56
            {
Paul's avatar
Paul committed
57
                check_shape(s, apply_pooling(it));
Paul's avatar
Paul committed
58
            }
Paul's avatar
Paul committed
59
            else if(it->name() == "add")
Paul's avatar
Paul committed
60
            {
Paul's avatar
Paul committed
61
                check_shape(s, apply_add(it));
Paul's avatar
Paul committed
62
            }
Paul's avatar
Paul committed
63
            else if(it->name() == "gemm")
Paul's avatar
Paul committed
64
            {
Paul's avatar
Paul committed
65
                check_shape(s, apply_gemm(it));
Paul's avatar
Paul committed
66
            }
Paul's avatar
Paul committed
67
            else if(it->name() == "contiguous")
68
            {
Paul's avatar
Paul committed
69
                check_shape(s, apply_contiguous(it));
70
            }
Paul's avatar
Paul committed
71
            else if(it->name() == "batch_norm_inference")
72
            {
Paul's avatar
Paul committed
73
                check_shape(s, apply_batch_norm_inference(it));
74
            }
Paul's avatar
Paul committed
75
76
77
78
            else if(it->name() == "softmax")
            {
                check_shape(s, apply_softmax(it));
            }
Paul's avatar
Paul committed
79
80
81
        }
    }

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

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

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

103
        auto workspace = insert_allocation(ins, ws, "workspace");
Paul's avatar
Paul committed
104
        auto output    = insert_allocation(ins, ins->get_shape());
Paul's avatar
Paul committed
105

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

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

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

Paul's avatar
Paul committed
120
    instruction_ref apply_activation(instruction_ref ins)
Paul's avatar
Paul committed
121
    {
wsttiger's avatar
wsttiger committed
122
        auto&& op = any_cast<op::activation>(ins->get_operator());
Paul's avatar
Paul committed
123
124
        auto ad   = make_relu();
        if(op.mode == "relu")
Paul's avatar
Paul committed
125
        {
Paul's avatar
Paul committed
126
            auto output = insert_allocation(ins, ins->get_shape());
Paul's avatar
Paul committed
127
            return prog->replace_instruction(
Paul's avatar
Paul committed
128
                ins, miopen_relu{std::move(ad)}, ins->inputs().at(0), output);
Paul's avatar
Paul committed
129
        }
Paul's avatar
Paul committed
130
        return ins;
Paul's avatar
Paul committed
131
    }
132

Khalique's avatar
Khalique committed
133
134
135
136
    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);
137

Khalique's avatar
Khalique committed
138
139
140
141
        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
142

Paul's avatar
Paul committed
143
144
    instruction_ref apply_softmax(instruction_ref ins)
    {
wsttiger's avatar
wsttiger committed
145
        auto&& op   = any_cast<op::softmax>(ins->get_operator());
Paul's avatar
Paul committed
146
        auto output = insert_allocation(ins, ins->get_shape());
Paul's avatar
Paul committed
147
        return prog->replace_instruction(ins, miopen_softmax{op}, ins->inputs().at(0), output);
Paul's avatar
Paul committed
148
149
    }

Paul's avatar
Paul committed
150
    instruction_ref apply_add(instruction_ref ins)
Paul's avatar
Paul committed
151
    {
Paul's avatar
Paul committed
152
        auto output = insert_allocation(ins, ins->get_shape());
Paul's avatar
Paul committed
153
        return prog->replace_instruction(
Paul's avatar
Paul committed
154
            ins, hip_add{}, ins->inputs().at(0), ins->inputs().at(1), output);
Paul's avatar
Paul committed
155
    }
Paul's avatar
Paul committed
156

Paul's avatar
Paul committed
157
    instruction_ref apply_gemm(instruction_ref ins)
Paul's avatar
Paul committed
158
    {
wsttiger's avatar
wsttiger committed
159
        auto&& op   = any_cast<op::gemm>(ins->get_operator());
Paul's avatar
Paul committed
160
        auto output = insert_allocation(ins, ins->get_shape());
Paul's avatar
Paul committed
161
        return prog->replace_instruction(
Paul's avatar
Paul committed
162
            ins, miopen_gemm{op}, ins->inputs().at(0), ins->inputs().at(1), output);
Paul's avatar
Paul committed
163
    }
164

Paul's avatar
Paul committed
165
    instruction_ref apply_contiguous(instruction_ref ins)
166
    {
wsttiger's avatar
wsttiger committed
167
        auto&& op   = any_cast<op::contiguous>(ins->get_operator());
Paul's avatar
Paul committed
168
        auto output = insert_allocation(ins, ins->get_shape());
Paul's avatar
Paul committed
169
        return prog->replace_instruction(ins, miopen_contiguous{op}, ins->inputs().at(0), output);
170
    }
171

Paul's avatar
Paul committed
172
    instruction_ref apply_batch_norm_inference(instruction_ref ins)
173
    {
wsttiger's avatar
wsttiger committed
174
        auto&& op       = any_cast<op::batch_norm_inference>(ins->get_operator());
Paul's avatar
Paul committed
175
        auto output     = insert_allocation(ins, ins->get_shape());
Paul's avatar
Paul committed
176
        shape old_shape = ins->inputs().at(1)->get_shape();
wsttiger's avatar
wsttiger committed
177
        std::vector<int64_t> new_shape{1, static_cast<int64_t>(old_shape.elements()), 1, 1};
wsttiger's avatar
wsttiger committed
178
        auto reshape_op = op::reshape{new_shape};
Paul's avatar
Paul committed
179
        std::vector<instruction_ref> reshapes;
Paul's avatar
Paul committed
180
181
        std::transform(ins->inputs().begin() + 1,
                       ins->inputs().end(),
Paul's avatar
Paul committed
182
183
                       std::back_inserter(reshapes),
                       [&](auto i) { return prog->insert_instruction(ins, reshape_op, i); });
Paul's avatar
Paul committed
184
        return prog->replace_instruction(ins,
Paul's avatar
Paul committed
185
                                         miopen_batch_norm_inference{op},
Paul's avatar
Paul committed
186
                                         ins->inputs().at(0),
Paul's avatar
Paul committed
187
188
189
190
191
                                         reshapes[0],
                                         reshapes[1],
                                         reshapes[2],
                                         reshapes[3],
                                         output);
192
    }
Paul's avatar
Paul committed
193
194
};

Paul's avatar
Paul committed
195
void lowering::apply(program& p) const { miopen_apply{&p, ctx}.apply(); }
Paul's avatar
Paul committed
196
} // namespace gpu
Paul's avatar
Paul committed
197
} // namespace migraph