miopen_target.cpp 4.87 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
7

#include <miopen/miopen.h>

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

Paul's avatar
Paul committed
11
using miopen_handle     = RTG_MANAGE_PTR(miopenHandle_t, miopenDestroy);
Paul's avatar
Paul committed
12
using tensor_descriptor = RTG_MANAGE_PTR(miopenTensorDescriptor_t, miopenDestroyTensorDescriptor);
Paul's avatar
Paul committed
13
14
15
16
using convolution_descriptor = RTG_MANAGE_PTR(miopenConvolutionDescriptor_t,
                                              miopenDestroyConvolutionDescriptor);
using activation_descriptor  = RTG_MANAGE_PTR(miopenActivationDescriptor_t,
                                             miopenDestroyActivationDescriptor);
Paul's avatar
Paul committed
17

Paul's avatar
Paul committed
18
template <class Result, class F, class... Ts>
Paul's avatar
Paul committed
19
20
21
Result make_obj(F f, Ts... xs)
{
    typename Result::pointer x = nullptr;
Paul's avatar
Paul committed
22
    auto status                = f(&x, xs...);
Paul's avatar
Paul committed
23
    Result r{x};
Paul's avatar
Paul committed
24
    if(status != miopenStatusSuccess)
Paul's avatar
Paul committed
25
26
27
28
29
30
31
32
33
34
35
        RTG_THROW("MIOpen call failed");
    return r;
}

tensor_descriptor make_tensor(const rtg::shape& s)
{
    auto t = make_obj<tensor_descriptor>(&miopenCreateTensorDescriptor);
    // Convert to ints
    std::vector<int> lens(s.lens().begin(), s.lens().end());
    std::vector<int> strides(s.strides().begin(), s.strides().end());
    miopenDataType_t d;
Paul's avatar
Paul committed
36
37
38
39
    if(s.type() == shape::float_type)
        d = miopenFloat;
    else
        RTG_THROW("Unsupported type");
Paul's avatar
Paul committed
40
41
42
43
44
45
46
    miopenSetTensorDescriptor(t.get(), d, s.lens().size(), lens.data(), strides.data());
    return t;
}

convolution_descriptor make_conv(const rtg::convolution& op)
{
    auto c = make_obj<convolution_descriptor>(&miopenCreateConvolutionDescriptor);
Paul's avatar
Paul committed
47
48
49
50
51
52
53
54
    miopenInitConvolutionDescriptor(c.get(),
                                    miopenConvolution,
                                    op.padding[0],
                                    op.padding[1],
                                    op.stride[0],
                                    op.stride[1],
                                    op.dilation[0],
                                    op.dilation[1]);
Paul's avatar
Paul committed
55
56
57
58
59
60
61
62
63
    return c;
}

struct miopen_convolution
{
    convolution op;
    convolution_descriptor cd;

    std::string name() const { return "miopen::convolution"; }
Paul's avatar
Paul committed
64
65
66
    shape compute_shape(std::vector<shape> inputs) const
    {
        return op.compute_shape({inputs.at(1), inputs.at(2)});
Paul's avatar
Paul committed
67
68
69
70
71
72
73
74
    }
    argument compute(shape output_shape, std::vector<argument> args) const
    {
        argument result;
        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
75
        float alpha = 1, beta = 0;
Paul's avatar
Paul committed
76
77
        int algo_count;
        miopenConvAlgoPerf_t perf;
Paul's avatar
Paul committed
78
        miopenFindConvolutionForwardAlgorithm(args[0].get(),
Paul's avatar
Paul committed
79
                                              x_desc.get(),
Paul's avatar
Paul committed
80
81
82
                                              args[1].get(),
                                              w_desc.get(),
                                              args[2].get(),
Paul's avatar
Paul committed
83
                                              cd.get(),
Paul's avatar
Paul committed
84
85
                                              y_desc.get(),
                                              args[4].get(),
Paul's avatar
Paul committed
86
87
88
                                              1,
                                              &algo_count,
                                              &perf,
Paul's avatar
Paul committed
89
                                              args[3].get(),
Paul's avatar
Paul committed
90
91
                                              args[3].get_shape().bytes(),
                                              false);
Paul's avatar
Paul committed
92
        miopenConvolutionForward(args[0].get(),
Paul's avatar
Paul committed
93
                                 &alpha,
Paul's avatar
Paul committed
94
95
96
97
                                 x_desc.get(),
                                 args[1].get(),
                                 w_desc.get(),
                                 args[2].get(),
Paul's avatar
Paul committed
98
99
100
                                 cd.get(),
                                 perf.fwd_algo,
                                 &beta,
Paul's avatar
Paul committed
101
102
103
                                 y_desc.get(),
                                 args[4].get(),
                                 args[3].get(),
Paul's avatar
Paul committed
104
                                 args[3].get_shape().bytes());
Paul's avatar
Paul committed
105
106
107
108
        return result;
    }
};

Paul's avatar
Paul committed
109
110
struct miopen_apply
{
Paul's avatar
Paul committed
111
    program* prog;
Paul's avatar
Paul committed
112
113
114

    void apply()
    {
Paul's avatar
Paul committed
115
116
117
118
        for(auto it = prog->begin(); it != prog->end(); it++)
        {
            if(it->op.name() == "convolution")
            {
Paul's avatar
Paul committed
119
                apply_convolution(it);
Paul's avatar
Paul committed
120
121
122
            }
            else if(it->op.name() == "activation")
            {
Paul's avatar
Paul committed
123
124
125
126
127
128
129
130
131
132
133
                apply_activation(it);
            }
        }
    }

    void apply_convolution(instruction_ref ins)
    {
        // auto&& op = any_cast<convolution>(ins->op);
        // prog->replace_instruction(ins, miopen_convolution{op}, ins->arguments);
    }

Paul's avatar
Paul committed
134
    void apply_activation(instruction_ref ins) {}
Paul's avatar
Paul committed
135
136
};

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

Paul's avatar
Paul committed
139
void miopen_target::apply(program& p) const { miopen_apply{&p}.apply(); }
Paul's avatar
Paul committed
140
141
142
143

} // namespace miopen

} // namespace rtg