miopen.cpp 10 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
#include <migraph/type_name.hpp>
Paul's avatar
Paul committed
11
#include <migraph/verify.hpp>
Paul's avatar
Paul committed
12
13
14

#include <miopen/miopen.h>

Paul's avatar
Paul committed
15
16
17
#include <future>
#include <thread>

Paul's avatar
Paul committed
18
19
#include "test.hpp"

Paul's avatar
Paul committed
20
21
22
23
#ifdef __clang__
#pragma clang diagnostic push
#pragma clang diagnostic ignored "-Wglobal-constructors"
#endif
Paul's avatar
Paul committed
24

Paul's avatar
Paul committed
25
26
27
28
29
30
31
32
33
34
35
// An improved async, that doesn't block
template <class Function>
std::future<typename std::result_of<Function()>::type> detach_async(Function&& f)
{
    using result_type = typename std::result_of<Function()>::type;
    std::packaged_task<result_type()> task(std::forward<Function>(f));
    auto fut = task.get_future();
    std::thread(std::move(task)).detach();
    return std::move(fut);
}

Paul's avatar
Paul committed
36
37
struct auto_print
{
Paul's avatar
Paul committed
38
    static std::array<std::function<void()>, 2> handlers;
Paul's avatar
Paul committed
39
    int index;
Paul's avatar
Paul committed
40
    template <class T>
Paul's avatar
Paul committed
41
    auto_print(T& x, int i) : index(i)
Paul's avatar
Paul committed
42
    {
Paul's avatar
Paul committed
43
        handlers[index] = [&x] { std::cout << x << std::endl; };
Paul's avatar
Paul committed
44
    }
Paul's avatar
Paul committed
45

Paul's avatar
Paul committed
46
    ~auto_print()
Paul's avatar
Paul committed
47
    {
Paul's avatar
Paul committed
48
        handlers[index] = [] {};
Paul's avatar
Paul committed
49
50
    }
};
Paul's avatar
Paul committed
51
std::array<std::function<void()>, 2> auto_print::handlers = {};
Paul's avatar
Paul committed
52

Paul's avatar
Paul committed
53
template <class V>
Paul's avatar
Paul committed
54
migraph::argument run_cpu()
Paul's avatar
Paul committed
55
{
Paul's avatar
Paul committed
56
57
    V v;
    auto p = v.create_program();
Paul's avatar
Paul committed
58
    auto_print pp{p, 0};
Paul's avatar
Paul committed
59
    p.compile(migraph::cpu::cpu_target{});
Paul's avatar
Paul committed
60
    migraph::program::parameter_map m;
Paul's avatar
Paul committed
61
    for(auto&& x : p.get_parameter_shapes())
Paul's avatar
Paul committed
62
63
64
    {
        m[x.first] = migraph::generate_argument(x.second);
    }
Paul's avatar
Paul committed
65
    return p.eval(m);
Paul's avatar
Paul committed
66
67
}

Paul's avatar
Paul committed
68
template <class V>
Paul's avatar
Paul committed
69
migraph::argument run_gpu()
Paul's avatar
Paul committed
70
{
Paul's avatar
Paul committed
71
72
    V v;
    auto p = v.create_program();
Paul's avatar
Paul committed
73
    auto_print pp{p, 1};
Paul's avatar
Paul committed
74
    p.compile(migraph::gpu::target{});
Paul's avatar
Paul committed
75

Paul's avatar
Paul committed
76
    migraph::program::parameter_map m;
Paul's avatar
Paul committed
77
    for(auto&& x : p.get_parameter_shapes())
Paul's avatar
Paul committed
78
    {
Paul's avatar
Paul committed
79
        m[x.first] = migraph::gpu::to_gpu(migraph::generate_argument(x.second));
Paul's avatar
Paul committed
80
81
    }

Paul's avatar
Paul committed
82
    return migraph::gpu::from_gpu(p.eval(m));
Paul's avatar
Paul committed
83
84
}

Paul's avatar
Paul committed
85
template <class V>
Paul's avatar
Paul committed
86
void verify_program()
Paul's avatar
Paul committed
87
{
Paul's avatar
Paul committed
88
    std::set_terminate(+[] {
Paul's avatar
Paul committed
89
        std::cout << "FAILED: " << migraph::get_type_name<V>() << std::endl;
Paul's avatar
Paul committed
90
91
92
93
94
95
        try
        {
            std::rethrow_exception(std::current_exception());
        }
        catch(const std::exception& e)
        {
Paul's avatar
Paul committed
96
            std::cout << "    what(): " << e.what() << std::endl;
Paul's avatar
Paul committed
97
        }
Paul's avatar
Paul committed
98
99
        std::cout << std::endl;
        for(auto&& handle : auto_print::handlers)
Paul's avatar
Paul committed
100
            handle();
Paul's avatar
Paul committed
101
    });
Paul's avatar
Paul committed
102
    auto cpu_arg_f = detach_async([] { return run_cpu<V>(); });
Paul's avatar
Paul committed
103
    auto gpu_arg   = run_gpu<V>();
Paul's avatar
Paul committed
104
    visit_all(cpu_arg_f.get(), gpu_arg)([](auto cpu, auto gpu) {
Paul's avatar
Paul committed
105
        if(not migraph::verify_range(cpu, gpu))
Paul's avatar
Paul committed
106
        {
107
            // TODO: Check for nans
Paul's avatar
Paul committed
108
109
110
            std::cout << "FAILED: " << migraph::get_type_name<V>() << std::endl;
        }
    });
Paul's avatar
Paul committed
111
    std::set_terminate(nullptr);
Paul's avatar
Paul committed
112
113
}

Paul's avatar
Paul committed
114
115
116
117
118
struct test_literals
{
    migraph::program create_program() const
    {
        migraph::program p;
Paul's avatar
Paul committed
119
120
121
122
        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
123
124
125
126
127
128
        auto conv = p.add_instruction(migraph::convolution{}, input, weights);
        p.add_instruction(migraph::activation{"relu"}, conv);
        return p;
    }
};

Paul's avatar
Paul committed
129
130
struct test_add
{
Paul's avatar
Paul committed
131
    migraph::program create_program() const
Paul's avatar
Paul committed
132
    {
Paul's avatar
Paul committed
133
134
        migraph::program p;
        migraph::shape s{migraph::shape::float_type, {3}};
Paul's avatar
Paul committed
135
136
        auto x = p.add_parameter("x", s);
        auto y = p.add_parameter("y", s);
Paul's avatar
Paul committed
137
        p.add_instruction(migraph::add{}, x, y);
Paul's avatar
Paul committed
138
139
140
141
142
143
        return p;
    }
};

struct test_add_broadcast
{
Paul's avatar
Paul committed
144
    migraph::program create_program() const
Paul's avatar
Paul committed
145
    {
Paul's avatar
Paul committed
146
147
148
149
150
151
        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
152
153
154
155
        return p;
    }
};

Paul's avatar
Paul committed
156
struct test_conv_relu
Paul's avatar
Paul committed
157
{
Paul's avatar
Paul committed
158
    migraph::program create_program() const
Paul's avatar
Paul committed
159
    {
Paul's avatar
Paul committed
160
        migraph::program p;
Paul's avatar
Paul committed
161
162
163
164
        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
165
        p.add_instruction(migraph::activation{"relu"}, conv);
Paul's avatar
Paul committed
166
167
168
169
        return p;
    }
};

Paul's avatar
Paul committed
170
171
struct test_conv_pooling
{
Paul's avatar
Paul committed
172
    migraph::program create_program() const
Paul's avatar
Paul committed
173
    {
Paul's avatar
Paul committed
174
        migraph::program p;
Paul's avatar
Paul committed
175
176
177
178
        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
179
180
181
        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
182
183
184
185
        return p;
    }
};

Paul's avatar
Paul committed
186
187
struct test_gemm
{
Paul's avatar
Paul committed
188
    migraph::program create_program() const
Paul's avatar
Paul committed
189
    {
Paul's avatar
Paul committed
190
191
192
193
        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
194
195
196
197
        return p;
    }
};

Paul's avatar
Paul committed
198
199
200
201
202
203
204
205
206
207
208
209
struct test_gemm_ld
{
    migraph::program create_program() const
    {
        migraph::program p;
        auto a = p.add_parameter("a", migraph::shape{migraph::shape::float_type, {4, 5}, {10, 1}});
        auto b = p.add_parameter("b", migraph::shape{migraph::shape::float_type, {5, 3}, {20, 1}});
        p.add_instruction(migraph::gemm{}, a, b);
        return p;
    }
};

210
211
212
213
214
struct test_gemm_transposeb
{
    migraph::program create_program() const
    {
        migraph::program p;
Paul's avatar
Paul committed
215
216
        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, {3, 5}});
217
218
219
220
221
222
223
224
225
226
227
        auto bt = p.add_instruction(migraph::transpose{{1, 0}}, b);
        p.add_instruction(migraph::gemm{}, a, bt);
        return p;
    }
};

struct test_gemm_transposea
{
    migraph::program create_program() const
    {
        migraph::program p;
Paul's avatar
Paul committed
228
229
        auto a  = p.add_parameter("a", migraph::shape{migraph::shape::float_type, {5, 4}});
        auto b  = p.add_parameter("b", migraph::shape{migraph::shape::float_type, {5, 3}});
230
231
232
233
234
235
236
237
238
239
240
        auto at = p.add_instruction(migraph::transpose{{1, 0}}, a);
        p.add_instruction(migraph::gemm{}, at, b);
        return p;
    }
};

struct test_gemm_transposeab
{
    migraph::program create_program() const
    {
        migraph::program p;
Paul's avatar
Paul committed
241
242
        auto a  = p.add_parameter("a", migraph::shape{migraph::shape::float_type, {5, 4}});
        auto b  = p.add_parameter("b", migraph::shape{migraph::shape::float_type, {3, 5}});
243
244
245
246
247
248
249
        auto at = p.add_instruction(migraph::transpose{{1, 0}}, a);
        auto bt = p.add_instruction(migraph::transpose{{1, 0}}, b);
        p.add_instruction(migraph::gemm{}, at, bt);
        return p;
    }
};

250
251
252
253
254
struct test_contiguous
{
    migraph::program create_program() const
    {
        migraph::program p;
255
        migraph::shape s{migraph::shape::float_type, {4, 4, 4, 3}, {48, 4, 1, 16}};
256
257
258
259
260
261
        auto x = p.add_parameter("x", s);
        p.add_instruction(migraph::contiguous{}, x);
        return p;
    }
};

262
struct test_transpose
263
{
264
265
266
267
268
269
270
271
272
273
274
    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;
    }
};
275

Paul's avatar
Paul committed
276
277
278
279
280
281
282
283
284
285
286
287
288
289
290
291
292
293
294
295
296
297
298
struct test_batchnorm_inference_2
{
    const size_t width    = 14;
    const size_t height   = 14;
    const size_t channels = 256;
    const size_t batches  = 1;

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

        migraph::shape s{migraph::shape::float_type, {batches, channels, height, width}};
        migraph::shape vars{migraph::shape::float_type, {channels}};
        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;
    }
};

wsttiger's avatar
wsttiger committed
299
300
301
302
303
304
305
306
307
308
309
310
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}};
311
        migraph::shape vars{migraph::shape::float_type, {channels}};
wsttiger's avatar
wsttiger committed
312
313
314
315
316
317
318
319
320
321
        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;
    }
};

Paul's avatar
Paul committed
322
323
int main()
{
Paul's avatar
Paul committed
324
    verify_program<test_add>();
Paul's avatar
Paul committed
325
    verify_program<test_add_broadcast>();
Paul's avatar
Paul committed
326
327
328
    verify_program<test_conv_relu>();
    verify_program<test_conv_pooling>();
    verify_program<test_gemm>();
Paul's avatar
Paul committed
329
    // verify_program<test_gemm_ld>();
330
331
332
    verify_program<test_gemm_transposeb>();
    verify_program<test_gemm_transposea>();
    verify_program<test_gemm_transposeab>();
333
334
    verify_program<test_contiguous>();
    verify_program<test_transpose>();
335
    verify_program<test_batchnorm_inference>();
Paul's avatar
Paul committed
336
    verify_program<test_batchnorm_inference_2>();
Paul's avatar
Paul committed
337
}