miopen_target.cpp 5.13 KB
Newer Older
Paul's avatar
Paul committed
1
2
#include <rtg/miopen/miopen_target.hpp>
#include <rtg/manage_ptr.hpp>
Paul's avatar
Paul committed
3
4
#include <rtg/instruction.hpp>
#include <rtg/operators.hpp>
Paul's avatar
Paul committed
5
6
#include <rtg/miopen/miopen.hpp>
#include <rtg/miopen/hip.hpp>
Paul's avatar
Paul committed
7

Paul's avatar
Paul committed
8
9
namespace rtg {
namespace miopen {
Paul's avatar
Paul committed
10

Paul's avatar
Paul committed
11
12
13
struct miopen_convolution
{
    convolution op;
Paul's avatar
Paul committed
14
    shared<convolution_descriptor> cd;
Paul's avatar
Paul committed
15
16

    std::string name() const { return "miopen::convolution"; }
Paul's avatar
Paul committed
17
18
    shape compute_shape(std::vector<shape> inputs) const
    {
Paul's avatar
Paul committed
19
        check_shapes{inputs}.has(4);
Paul's avatar
Paul committed
20
        return op.compute_shape({inputs.at(1), inputs.at(2)});
Paul's avatar
Paul committed
21
22
23
24
25
26
27
    }
    argument compute(shape output_shape, std::vector<argument> args) const
    {
        auto x_desc = make_tensor(args[1].get_shape());
        auto w_desc = make_tensor(args[2].get_shape());
        auto y_desc = make_tensor(output_shape);

Paul's avatar
Paul committed
28
        float alpha = 1, beta = 0;
Paul's avatar
Paul committed
29
30
        int algo_count;
        miopenConvAlgoPerf_t perf;
31
        miopenFindConvolutionForwardAlgorithm(args[0].implicit(),
Paul's avatar
Paul committed
32
                                              x_desc.get(),
33
                                              args[1].implicit(),
Paul's avatar
Paul committed
34
                                              w_desc.get(),
35
                                              args[2].implicit(),
Paul's avatar
Paul committed
36
                                              cd.get(),
Paul's avatar
Paul committed
37
                                              y_desc.get(),
38
                                              args[3].implicit(),
Paul's avatar
Paul committed
39
40
41
                                              1,
                                              &algo_count,
                                              &perf,
Paul's avatar
Paul committed
42
43
                                              nullptr,
                                              0,
Paul's avatar
Paul committed
44
                                              false);
45
        miopenConvolutionForward(args[0].implicit(),
Paul's avatar
Paul committed
46
                                 &alpha,
Paul's avatar
Paul committed
47
                                 x_desc.get(),
48
                                 args[1].implicit(),
Paul's avatar
Paul committed
49
                                 w_desc.get(),
50
                                 args[2].implicit(),
Paul's avatar
Paul committed
51
52
53
                                 cd.get(),
                                 perf.fwd_algo,
                                 &beta,
Paul's avatar
Paul committed
54
                                 y_desc.get(),
55
                                 args[3].implicit(),
Paul's avatar
Paul committed
56
57
58
59
60
61
62
63
64
65
                                 nullptr,
                                 0);
        return args[3];
    }
};

struct miopen_relu
{
    shared<activation_descriptor> ad;
    std::string name() const { return "miopen::relu"; }
Paul's avatar
Paul committed
66
    shape compute_shape(std::vector<shape> inputs) const
Paul's avatar
Paul committed
67
    {
Paul's avatar
Paul committed
68
69
        check_shapes{inputs}.has(3);
        return inputs.at(1);
Paul's avatar
Paul committed
70
71
72
73
74
75
76
    }

    argument compute(shape output_shape, std::vector<argument> args) const
    {
        float alpha = 1, beta = 0;
        auto x_desc = make_tensor(args[1].get_shape());
        auto y_desc = make_tensor(output_shape);
77
        miopenActivationForward(args[0].implicit(),
Paul's avatar
Paul committed
78
79
80
                                ad.get(),
                                &alpha,
                                x_desc.get(),
81
                                args[1].implicit(),
Paul's avatar
Paul committed
82
83
                                &beta,
                                y_desc.get(),
84
                                args[2].implicit());
Paul's avatar
Paul committed
85
86

        return args[2];
Paul's avatar
Paul committed
87
88
89
    }
};

Paul's avatar
Paul committed
90
91
struct miopen_apply
{
Paul's avatar
Paul committed
92
93
    program* prog = nullptr;
    instruction_ref handle{};
Paul's avatar
Paul committed
94
95
96

    void apply()
    {
Paul's avatar
Paul committed
97
        handle = prog->add_parameter("handle", shape{shape::any_type});
Paul's avatar
Paul committed
98
99
100
101
        for(auto it = prog->begin(); it != prog->end(); it++)
        {
            if(it->op.name() == "convolution")
            {
Paul's avatar
Paul committed
102
                apply_convolution(it);
Paul's avatar
Paul committed
103
104
105
            }
            else if(it->op.name() == "activation")
            {
Paul's avatar
Paul committed
106
107
108
109
110
                apply_activation(it);
            }
        }
    }

Paul's avatar
Paul committed
111
112
    instruction_ref insert_allocation(instruction_ref ins, const shape& s)
    {
Paul's avatar
Paul committed
113
        if(ins == --prog->end())
Paul's avatar
Paul committed
114
115
116
117
118
        {
            return prog->add_parameter("output", s);
        }
        else
        {
Paul's avatar
Paul committed
119
            auto is     = prog->add_outline(s);
Paul's avatar
Paul committed
120
121
122
123
124
125
            auto result = prog->insert_instruction(ins, hip_allocate{}, is);
            prog->insert_instruction(++ins, hip_free{}, result);
            return result;
        }
    }

Paul's avatar
Paul committed
126
127
    void apply_convolution(instruction_ref ins)
    {
Paul's avatar
Paul committed
128
129
        auto&& op   = any_cast<convolution>(ins->op);
        auto cd     = make_conv(op);
Paul's avatar
Paul committed
130
131
        auto output = insert_allocation(ins, ins->result);

Paul's avatar
Paul committed
132
133
134
135
136
137
        prog->replace_instruction(ins,
                                  miopen_convolution{op, std::move(cd)},
                                  handle,
                                  ins->arguments.at(0),
                                  ins->arguments.at(1),
                                  output);
Paul's avatar
Paul committed
138
139
    }

Paul's avatar
Paul committed
140
    void apply_activation(instruction_ref ins)
Paul's avatar
Paul committed
141
142
    {
        auto&& op = any_cast<activation>(ins->op);
Paul's avatar
Paul committed
143
144
        auto ad   = make_relu();
        if(op.mode == "relu")
Paul's avatar
Paul committed
145
146
        {
            auto output = insert_allocation(ins, ins->result);
Paul's avatar
Paul committed
147
148
            prog->replace_instruction(
                ins, miopen_relu{std::move(ad)}, handle, ins->arguments.at(0), output);
Paul's avatar
Paul committed
149
150
        }
    }
Paul's avatar
Paul committed
151
152
};

Paul's avatar
Paul committed
153
std::string miopen_target::name() const { return "miopen"; }
Paul's avatar
Paul committed
154

Paul's avatar
Paul committed
155
void miopen_target::apply(program& p) const { miopen_apply{&p}.apply(); }
Paul's avatar
Paul committed
156
157
158
159

} // namespace miopen

} // namespace rtg