miopen.cpp 10.2 KB
Newer Older
Paul's avatar
Paul committed
1

Paul's avatar
Paul committed
2
3
4
5
#include <migraph/program.hpp>
#include <migraph/operators.hpp>
#include <migraph/generate.hpp>
#include <migraph/cpu/cpu_target.hpp>
Paul's avatar
Paul committed
6
7
8
#include <migraph/gpu/target.hpp>
#include <migraph/gpu/miopen.hpp>
#include <migraph/gpu/hip.hpp>
Paul's avatar
Paul committed
9
#include <migraph/manage_ptr.hpp>
Paul's avatar
Paul committed
10
11
12
13

#include <miopen/miopen.h>

#include "test.hpp"
Paul's avatar
Paul committed
14
#include "verify.hpp"
Paul's avatar
Paul committed
15

Paul's avatar
Paul committed
16
template <class V>
Paul's avatar
Paul committed
17
migraph::argument run_cpu()
Paul's avatar
Paul committed
18
{
Paul's avatar
Paul committed
19
20
    V v;
    auto p = v.create_program();
Paul's avatar
Paul committed
21
    p.compile(migraph::cpu::cpu_target{});
Paul's avatar
Paul committed
22
    return p.eval(v.create_params());
Paul's avatar
Paul committed
23
24
}

Paul's avatar
Paul committed
25
template <class V>
Paul's avatar
Paul committed
26
migraph::argument run_gpu()
Paul's avatar
Paul committed
27
{
Paul's avatar
Paul committed
28
29
    V v;
    auto p = v.create_program();
Paul's avatar
Paul committed
30
    p.compile(migraph::gpu::target{});
Paul's avatar
Paul committed
31
32

    auto m = v.create_params();
Paul's avatar
Paul committed
33
    for(auto&& e : m)
Paul's avatar
Paul committed
34
    {
Paul's avatar
Paul committed
35
        e.second = migraph::gpu::to_gpu(e.second);
Paul's avatar
Paul committed
36
37
    }

Paul's avatar
Paul committed
38
    m["output"] = migraph::gpu::to_gpu(migraph::generate_argument(p.get_parameter_shape("output")));
Paul's avatar
Paul committed
39

Paul's avatar
Paul committed
40
    return migraph::gpu::from_gpu(p.eval(m));
Paul's avatar
Paul committed
41
42
}

Paul's avatar
Paul committed
43
template <class V>
Paul's avatar
Paul committed
44
void verify_program()
Paul's avatar
Paul committed
45
{
Paul's avatar
Paul committed
46
47
    auto cpu_arg = run_cpu<V>();
    auto gpu_arg = run_gpu<V>();
Paul's avatar
Paul committed
48
    visit_all(cpu_arg, gpu_arg)([](auto cpu, auto gpu) { EXPECT(test::verify_range(cpu, gpu)); });
Paul's avatar
Paul committed
49
50
}

Paul's avatar
Paul committed
51
52
53
54
55
struct test_literals
{
    migraph::program create_program() const
    {
        migraph::program p;
Paul's avatar
Paul committed
56
57
58
59
        auto input = p.add_literal(
            generate_literal(migraph::shape{migraph::shape::float_type, {4, 3, 3, 3}}));
        auto weights = p.add_literal(
            generate_literal(migraph::shape{migraph::shape::float_type, {4, 3, 3, 3}}));
Paul's avatar
Paul committed
60
61
62
63
64
        auto conv = p.add_instruction(migraph::convolution{}, input, weights);
        p.add_instruction(migraph::activation{"relu"}, conv);
        return p;
    }

Paul's avatar
Paul committed
65
    migraph::program::parameter_map create_params() const { return {}; }
Paul's avatar
Paul committed
66
67
};

Paul's avatar
Paul committed
68
69
struct test_add
{
Paul's avatar
Paul committed
70
    migraph::program create_program() const
Paul's avatar
Paul committed
71
    {
Paul's avatar
Paul committed
72
73
        migraph::program p;
        migraph::shape s{migraph::shape::float_type, {3}};
Paul's avatar
Paul committed
74
75
        auto x = p.add_parameter("x", s);
        auto y = p.add_parameter("y", s);
Paul's avatar
Paul committed
76
        p.add_instruction(migraph::add{}, x, y);
Paul's avatar
Paul committed
77
78
79
        return p;
    }

Paul's avatar
Paul committed
80
    migraph::program::parameter_map create_params() const
Paul's avatar
Paul committed
81
    {
Paul's avatar
Paul committed
82
83
84
        migraph::program::parameter_map m;
        m["x"] = migraph::generate_argument({migraph::shape::float_type, {3}});
        m["y"] = migraph::generate_argument({migraph::shape::float_type, {3}});
Paul's avatar
Paul committed
85
86
87
88
89
90
        return m;
    }
};

struct test_add_broadcast
{
Paul's avatar
Paul committed
91
    migraph::program create_program() const
Paul's avatar
Paul committed
92
    {
Paul's avatar
Paul committed
93
94
95
96
97
98
        migraph::program p;
        migraph::shape s{migraph::shape::float_type, {3}};
        auto x  = p.add_parameter("x", {migraph::shape::float_type, {2, 2, 3}});
        auto y  = p.add_parameter("y", {migraph::shape::float_type, {2, 2}});
        auto by = p.add_instruction(migraph::broadcast{0}, x, y);
        p.add_instruction(migraph::add{}, x, by);
Paul's avatar
Paul committed
99
100
101
        return p;
    }

Paul's avatar
Paul committed
102
    migraph::program::parameter_map create_params() const
Paul's avatar
Paul committed
103
    {
Paul's avatar
Paul committed
104
105
106
        migraph::program::parameter_map m;
        m["x"] = migraph::generate_argument({migraph::shape::float_type, {2, 2, 3}});
        m["y"] = migraph::generate_argument({migraph::shape::float_type, {2, 2}});
Paul's avatar
Paul committed
107
108
109
110
        return m;
    }
};

Paul's avatar
Paul committed
111
struct test_conv_relu
Paul's avatar
Paul committed
112
{
Paul's avatar
Paul committed
113
    migraph::program create_program() const
Paul's avatar
Paul committed
114
    {
Paul's avatar
Paul committed
115
        migraph::program p;
Paul's avatar
Paul committed
116
117
118
119
        auto input = p.add_parameter("x", migraph::shape{migraph::shape::float_type, {4, 3, 3, 3}});
        auto weights =
            p.add_parameter("w", migraph::shape{migraph::shape::float_type, {4, 3, 3, 3}});
        auto conv = p.add_instruction(migraph::convolution{}, input, weights);
Paul's avatar
Paul committed
120
        p.add_instruction(migraph::activation{"relu"}, conv);
Paul's avatar
Paul committed
121
122
123
        return p;
    }

Paul's avatar
Paul committed
124
    migraph::program::parameter_map create_params() const
Paul's avatar
Paul committed
125
    {
Paul's avatar
Paul committed
126
127
128
        migraph::program::parameter_map m;
        m["x"] = migraph::generate_argument({migraph::shape::float_type, {4, 3, 3, 3}});
        m["w"] = migraph::generate_argument({migraph::shape::float_type, {4, 3, 3, 3}});
Paul's avatar
Paul committed
129
130
131
132
        return m;
    }
};

Paul's avatar
Paul committed
133
134
struct test_conv_pooling
{
Paul's avatar
Paul committed
135
    migraph::program create_program() const
Paul's avatar
Paul committed
136
    {
Paul's avatar
Paul committed
137
        migraph::program p;
Paul's avatar
Paul committed
138
139
140
141
        auto input =
            p.add_parameter("x", migraph::shape{migraph::shape::float_type, {4, 3, 32, 32}});
        auto weights =
            p.add_parameter("w", migraph::shape{migraph::shape::float_type, {4, 3, 3, 3}});
Paul's avatar
Paul committed
142
143
144
        auto conv    = p.add_instruction(migraph::convolution{}, input, weights);
        auto pooling = p.add_instruction(migraph::pooling{"max"}, conv);
        p.add_instruction(migraph::activation{"relu"}, pooling);
Paul's avatar
Paul committed
145
146
147
        return p;
    }

Paul's avatar
Paul committed
148
    migraph::program::parameter_map create_params() const
Paul's avatar
Paul committed
149
    {
Paul's avatar
Paul committed
150
151
152
        migraph::program::parameter_map m;
        m["x"] = migraph::generate_argument({migraph::shape::float_type, {4, 3, 32, 32}});
        m["w"] = migraph::generate_argument({migraph::shape::float_type, {4, 3, 3, 3}});
Paul's avatar
Paul committed
153
154
155
156
        return m;
    }
};

Paul's avatar
Paul committed
157
158
struct test_gemm
{
Paul's avatar
Paul committed
159
    migraph::program create_program() const
Paul's avatar
Paul committed
160
    {
Paul's avatar
Paul committed
161
162
163
164
        migraph::program p;
        auto a = p.add_parameter("a", migraph::shape{migraph::shape::float_type, {4, 5}});
        auto b = p.add_parameter("b", migraph::shape{migraph::shape::float_type, {5, 3}});
        p.add_instruction(migraph::gemm{}, a, b);
Paul's avatar
Paul committed
165
166
167
        return p;
    }

Paul's avatar
Paul committed
168
    migraph::program::parameter_map create_params() const
Paul's avatar
Paul committed
169
    {
Paul's avatar
Paul committed
170
171
172
        migraph::program::parameter_map m;
        m["a"] = migraph::generate_argument({migraph::shape::float_type, {4, 5}});
        m["b"] = migraph::generate_argument({migraph::shape::float_type, {5, 3}});
Paul's avatar
Paul committed
173
174
175
176
        return m;
    }
};

177
178
179
180
181
struct test_contiguous
{
    migraph::program create_program() const
    {
        migraph::program p;
182
        migraph::shape s{migraph::shape::float_type, {4, 4, 4, 3}, {48, 4, 1, 16}};
183
184
185
186
187
188
189
190
        auto x = p.add_parameter("x", s);
        p.add_instruction(migraph::contiguous{}, x);
        return p;
    }

    migraph::program::parameter_map create_params() const
    {
        migraph::program::parameter_map m;
191
192
        m["x"] =
            migraph::generate_argument({migraph::shape::float_type, {4, 4, 4, 3}, {48, 4, 1, 16}});
193
194
195
196
        return m;
    }
};

197
struct test_transpose
198
{
199
200
201
202
203
204
205
206
207
208
    migraph::program create_program() const
    {
        migraph::program p;
        migraph::shape s{migraph::shape::float_type, {4, 3, 4, 4}};
        auto x                    = p.add_parameter("x", s);
        std::vector<int64_t> perm = {0, 2, 3, 1};
        auto l                    = p.add_instruction(migraph::transpose{perm}, x);
        p.add_instruction(migraph::contiguous{}, l);
        return p;
    }
209

210
211
212
213
214
215
216
    migraph::program::parameter_map create_params() const
    {
        migraph::program::parameter_map m;
        m["x"] = migraph::generate_argument({migraph::shape::float_type, {4, 3, 4, 4}});
        return m;
    }
};
217

wsttiger's avatar
wsttiger committed
218
219
220
221
222
223
224
225
226
227
228
229
struct test_batchnorm_inference
{
    const size_t width    = 3;
    const size_t height   = 3;
    const size_t channels = 3;
    const size_t batches  = 4;

    migraph::program create_program() const
    {
        migraph::program p;

        migraph::shape s{migraph::shape::float_type, {batches, channels, height, width}};
wsttiger's avatar
wsttiger committed
230
        migraph::shape vars{migraph::shape::float_type, {1, channels, 1, 1}};
wsttiger's avatar
wsttiger committed
231
232
233
234
235
236
237
238
239
240
241
242
243
244
245
246
247
248
249
250
251
252
253
254
255
256
257
258
259
260
261
262
        auto x        = p.add_parameter("x", s);
        auto mean     = p.add_parameter("mean", vars);
        auto variance = p.add_parameter("variance", vars);
        auto scale    = p.add_parameter("scale", vars);
        auto bias     = p.add_parameter("bias", vars);
        p.add_instruction(migraph::batch_norm_inference{}, x, mean, variance, scale, bias);
        return p;
    }

    migraph::program::parameter_map create_params() const
    {
        migraph::program::parameter_map m;
        migraph::shape s{migraph::shape::float_type, {batches, channels, height, width}};
        migraph::shape vars{migraph::shape::float_type, {channels}};
        m["x"]        = migraph::generate_argument(s);
        m["mean"]     = migraph::generate_argument(vars);
        m["variance"] = migraph::generate_argument(vars);
        m["scale"]    = migraph::generate_argument(vars);
        m["bias"]     = migraph::generate_argument(vars);
        return m;
    }
};

void batch_norm_inference_test()
{
    migraph::program p;
    const size_t width = 2, height = 2, channels = 4, batches = 2;
    const float x_val = 8.0f, mean_val = 2.0f, variance_val = 4.0f, scale_val = 2.0f,
                bias_val   = 1.0f;
    const float output_val = scale_val * (x_val - mean_val) / (std::sqrt(variance_val)) + bias_val;

    migraph::shape s{migraph::shape::float_type, {batches, channels, height, width}};
wsttiger's avatar
wsttiger committed
263
    migraph::shape vars{migraph::shape::float_type, {1, channels, 1, 1}};
wsttiger's avatar
wsttiger committed
264
265
266
267
268
269
270
271
272
273
274
275
276
277
278
279
280
281
282
283
    std::vector<float> x_data(width * height * channels * batches);
    std::vector<float> scale_data(channels);
    std::vector<float> bias_data(channels);
    std::vector<float> mean_data(channels);
    std::vector<float> variance_data(channels);

    std::fill(x_data.begin(), x_data.end(), x_val);
    std::fill(mean_data.begin(), mean_data.end(), mean_val);
    std::fill(variance_data.begin(), variance_data.end(), variance_val);
    std::fill(scale_data.begin(), scale_data.end(), scale_val);
    std::fill(bias_data.begin(), bias_data.end(), bias_val);

    auto x        = p.add_literal(migraph::literal{s, x_data});
    auto scale    = p.add_literal(migraph::literal{vars, scale_data});
    auto bias     = p.add_literal(migraph::literal{vars, bias_data});
    auto mean     = p.add_literal(migraph::literal{vars, mean_data});
    auto variance = p.add_literal(migraph::literal{vars, variance_data});

    p.add_instruction(migraph::batch_norm_inference{}, x, mean, variance, scale, bias);
    p.compile(migraph::gpu::target{});
284
285
286
287

    migraph::program::parameter_map m;
    m["output"] = migraph::gpu::to_gpu(migraph::generate_argument(p.get_parameter_shape("output")));
    auto result = migraph::gpu::from_gpu(p.eval(m));
wsttiger's avatar
wsttiger committed
288
289
290
291
292
293
294
295
296

    std::vector<float> result_vector(width * height * channels * batches);
    std::vector<float> gold(width * height * channels * batches);
    std::fill(gold.begin(), gold.end(), output_val);
    result.visit([&](auto output) { result_vector.assign(output.begin(), output.end()); });

    EXPECT(test::verify_range(result_vector, gold));
}

Paul's avatar
Paul committed
297
298
int main()
{
Paul's avatar
Paul committed
299
    verify_program<test_add>();
Paul's avatar
Paul committed
300
    verify_program<test_add_broadcast>();
Paul's avatar
Paul committed
301
302
303
    verify_program<test_conv_relu>();
    verify_program<test_conv_pooling>();
    verify_program<test_gemm>();
304
305
    verify_program<test_contiguous>();
    verify_program<test_transpose>();
306
307
    verify_program<test_batchnorm_inference>();
    batch_norm_inference_test();
Paul's avatar
Paul committed
308
}