miopen.cpp 21 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>
11
#include <migraph/verify_args.hpp>
wsttiger's avatar
wsttiger committed
12
#include <migraph/instruction.hpp>
Paul's avatar
Paul committed
13
14
15

#include <miopen/miopen.h>

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

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

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

Paul's avatar
Paul committed
26
27
// An improved async, that doesn't block
template <class Function>
Paul's avatar
Paul committed
28
29
std::future<typename std::result_of<Function()>::type> detach_async(Function&& f,
                                                                    bool parallel = true)
Paul's avatar
Paul committed
30
{
Paul's avatar
Paul committed
31
32
33
34
35
36
37
38
    if(parallel)
    {
        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);
    }
39
    return std::async(std::launch::deferred, std::forward<Function>(f));
Paul's avatar
Paul committed
40
41
}

Paul's avatar
Paul committed
42
43
struct auto_print
{
Paul's avatar
Paul committed
44
45
46
47
48
49
50
51
52
53
54
55
56
57
58
59
60
61
62
    static void set_terminate_handler(const std::string& name)
    {
        static std::string pname;
        pname = name;
        std::set_terminate(+[] {
            std::cout << "FAILED: " << pname << std::endl;
            try
            {
                std::rethrow_exception(std::current_exception());
            }
            catch(const std::exception& e)
            {
                std::cout << "    what(): " << e.what() << std::endl;
            }
            std::cout << std::endl;
            for(auto&& handle : auto_print::handlers)
                handle();
        });
    }
Paul's avatar
Paul committed
63
    static std::array<std::function<void()>, 2> handlers;
Paul's avatar
Paul committed
64
    int index;
Paul's avatar
Paul committed
65
    template <class T>
Paul's avatar
Paul committed
66
    auto_print(T& x, int i) : index(i)
Paul's avatar
Paul committed
67
    {
Paul's avatar
Paul committed
68
        handlers[index] = [&x] { std::cout << x << std::endl; };
Paul's avatar
Paul committed
69
    }
Paul's avatar
Paul committed
70

Paul's avatar
Paul committed
71
    ~auto_print()
Paul's avatar
Paul committed
72
    {
Paul's avatar
Paul committed
73
        handlers[index] = [] {};
Paul's avatar
Paul committed
74
75
    }
};
Paul's avatar
Paul committed
76
std::array<std::function<void()>, 2> auto_print::handlers = {};
Paul's avatar
Paul committed
77

Paul's avatar
Paul committed
78
template <class T>
Paul's avatar
Latest  
Paul committed
79
80
81
82
83
auto get_hash(const T& x)
{
    return std::hash<T>{}(x);
}

Paul's avatar
Paul committed
84
void compile_check(migraph::program& p, const migraph::target& t)
Paul's avatar
Paul committed
85
86
{
    auto name = t.name();
Paul's avatar
Paul committed
87
    auto s    = p.get_shape();
Paul's avatar
Paul committed
88
    std::stringstream ss;
Paul's avatar
Paul committed
89
    p.compile(t, migraph::tracer{ss});
Paul's avatar
Paul committed
90
    if(p.get_shape() != s)
Paul's avatar
Paul committed
91
92
93
94
95
96
    {
        std::cout << ss.str() << std::endl;
        throw std::runtime_error("Compiling program with " + name + " alters its shape");
    }
}

Paul's avatar
Paul committed
97
template <class V>
Paul's avatar
Paul committed
98
migraph::argument run_cpu(migraph::program& p)
Paul's avatar
Paul committed
99
{
Paul's avatar
Paul committed
100
    V v;
Paul's avatar
Paul committed
101
    p = v.create_program();
Paul's avatar
Paul committed
102
    auto_print pp{p, 0};
Paul's avatar
Paul committed
103
    compile_check(p, migraph::cpu::cpu_target{});
Paul's avatar
Paul committed
104
    migraph::program::parameter_map m;
Paul's avatar
Paul committed
105
    for(auto&& x : p.get_parameter_shapes())
Paul's avatar
Paul committed
106
    {
Paul's avatar
Latest  
Paul committed
107
        m[x.first] = migraph::generate_argument(x.second, get_hash(x.first));
Paul's avatar
Paul committed
108
    }
Paul's avatar
Paul committed
109
    return p.eval(m);
Paul's avatar
Paul committed
110
111
}

Paul's avatar
Paul committed
112
template <class V>
Paul's avatar
Paul committed
113
migraph::argument run_gpu(migraph::program& p)
Paul's avatar
Paul committed
114
{
Paul's avatar
Paul committed
115
    V v;
Paul's avatar
Paul committed
116
    p = v.create_program();
Paul's avatar
Paul committed
117
    auto_print pp{p, 1};
Paul's avatar
Paul committed
118
    compile_check(p, migraph::gpu::target{});
Paul's avatar
Paul committed
119
    migraph::program::parameter_map m;
Paul's avatar
Paul committed
120
    for(auto&& x : p.get_parameter_shapes())
Paul's avatar
Paul committed
121
    {
Paul's avatar
Latest  
Paul committed
122
        m[x.first] = migraph::gpu::to_gpu(migraph::generate_argument(x.second, get_hash(x.first)));
Paul's avatar
Paul committed
123
    }
Paul's avatar
Paul committed
124
    EXPECT(bool{m.find("output") != m.end()});
Paul's avatar
Paul committed
125
    return migraph::gpu::from_gpu(p.eval(m));
Paul's avatar
Paul committed
126
127
}

Paul's avatar
Paul committed
128
129
130
131
template <class V>
void verify_program()
{
    auto_print::set_terminate_handler(migraph::get_type_name<V>());
Paul's avatar
Paul committed
132
133
134
135
    migraph::program cpu_prog;
    migraph::program gpu_prog;
    auto cpu_arg_f = detach_async([&] { return run_cpu<V>(cpu_prog); });
    auto gpu_arg   = run_gpu<V>(gpu_prog);
Paul's avatar
Paul committed
136
    bool passed    = verify_args(migraph::get_type_name<V>(), cpu_arg_f.get(), gpu_arg);
Paul's avatar
Paul committed
137
138
139
140
141
142
143
144
145
    if(not passed)
    {
        V v;
        auto p = v.create_program();
        std::cout << p << std::endl;
        std::cout << "cpu:\n" << cpu_prog << std::endl;
        std::cout << "gpu:\n" << gpu_prog << std::endl;
        std::cout << std::endl;
    }
Paul's avatar
Paul committed
146
    std::set_terminate(nullptr);
Paul's avatar
Paul committed
147
148
}

Paul's avatar
Paul committed
149
150
151
152
153
struct test_literals
{
    migraph::program create_program() const
    {
        migraph::program p;
Paul's avatar
Paul committed
154
155
156
157
        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}}));
wsttiger's avatar
wsttiger committed
158
159
        auto conv = p.add_instruction(migraph::op::convolution{}, input, weights);
        p.add_instruction(migraph::op::activation{"relu"}, conv);
Paul's avatar
Paul committed
160
161
162
163
        return p;
    }
};

Paul's avatar
Paul committed
164
165
struct test_add
{
Paul's avatar
Paul committed
166
    migraph::program create_program() const
Paul's avatar
Paul committed
167
    {
Paul's avatar
Paul committed
168
169
        migraph::program p;
        migraph::shape s{migraph::shape::float_type, {3}};
Paul's avatar
Paul committed
170
171
        auto x = p.add_parameter("x", s);
        auto y = p.add_parameter("y", s);
wsttiger's avatar
wsttiger committed
172
        p.add_instruction(migraph::op::add{}, x, y);
Paul's avatar
Paul committed
173
174
175
176
177
178
        return p;
    }
};

struct test_add_broadcast
{
Paul's avatar
Paul committed
179
    migraph::program create_program() const
Paul's avatar
Paul committed
180
    {
Paul's avatar
Paul committed
181
182
183
184
        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}});
wsttiger's avatar
wsttiger committed
185
        auto by = p.add_instruction(migraph::op::broadcast{0, x->get_shape()}, y);
wsttiger's avatar
wsttiger committed
186
        p.add_instruction(migraph::op::add{}, x, by);
Paul's avatar
Paul committed
187
188
189
190
        return p;
    }
};

Paul's avatar
Paul committed
191
192
193
194
195
196
197
198
struct test_add_broadcast2
{
    migraph::program create_program() const
    {
        migraph::program p;
        migraph::shape s{migraph::shape::float_type, {3}};
        auto x  = p.add_parameter("x", {migraph::shape::float_type, {2, 3, 4}});
        auto y  = p.add_parameter("y", {migraph::shape::float_type, {3}});
wsttiger's avatar
wsttiger committed
199
        auto by = p.add_instruction(migraph::op::broadcast{1, x->get_shape()}, y);
wsttiger's avatar
wsttiger committed
200
        p.add_instruction(migraph::op::add{}, x, by);
Paul's avatar
Paul committed
201
202
203
204
        return p;
    }
};

Paul's avatar
Latest  
Paul committed
205
206
207
208
209
210
211
212
struct test_add_broadcast3
{
    migraph::program create_program() const
    {
        migraph::program p;
        migraph::shape s{migraph::shape::float_type, {3}};
        auto x  = p.add_parameter("x", {migraph::shape::float_type, {2, 4, 5}});
        auto y  = p.add_parameter("y", {migraph::shape::float_type, {4}});
wsttiger's avatar
wsttiger committed
213
        auto by = p.add_instruction(migraph::op::broadcast{1, x->get_shape()}, y);
wsttiger's avatar
wsttiger committed
214
        p.add_instruction(migraph::op::add{}, x, by);
Paul's avatar
Latest  
Paul committed
215
216
217
218
219
220
221
222
223
224
225
226
        return p;
    }
};

struct test_add_broadcast4
{
    migraph::program create_program() const
    {
        migraph::program p;
        migraph::shape s{migraph::shape::float_type, {3}};
        auto x  = p.add_parameter("x", {migraph::shape::float_type, {2, 3, 5}});
        auto y  = p.add_parameter("y", {migraph::shape::float_type, {3}});
wsttiger's avatar
wsttiger committed
227
        auto by = p.add_instruction(migraph::op::broadcast{1, x->get_shape()}, y);
wsttiger's avatar
wsttiger committed
228
        p.add_instruction(migraph::op::add{}, x, by);
Paul's avatar
Latest  
Paul committed
229
230
231
232
        return p;
    }
};

Paul's avatar
Paul committed
233
234
235
236
237
238
239
240
struct test_add_broadcast5
{
    migraph::program create_program() const
    {
        migraph::program p;
        migraph::shape s{migraph::shape::float_type, {3}};
        auto x  = p.add_parameter("x", {migraph::shape::float_type, {2, 4, 8}});
        auto y  = p.add_parameter("y", {migraph::shape::float_type, {4}});
wsttiger's avatar
wsttiger committed
241
        auto by = p.add_instruction(migraph::op::broadcast{1, x->get_shape()}, y);
wsttiger's avatar
wsttiger committed
242
        p.add_instruction(migraph::op::add{}, x, by);
Paul's avatar
Paul committed
243
244
245
246
        return p;
    }
};

Paul's avatar
Paul committed
247
248
249
250
251
struct test_softmax
{
    migraph::program create_program() const
    {
        migraph::program p;
Paul's avatar
Paul committed
252
        auto x = p.add_parameter("x", migraph::shape{migraph::shape::float_type, {5, 3, 4, 2}});
wsttiger's avatar
wsttiger committed
253
        p.add_instruction(migraph::op::softmax{}, x);
Paul's avatar
Paul committed
254
255
256
257
258
259
260
261
262
        return p;
    }
};

struct test_softmax2
{
    migraph::program create_program() const
    {
        migraph::program p;
Paul's avatar
Paul committed
263
        auto x = p.add_parameter("x", migraph::shape{migraph::shape::float_type, {1, 1000, 1, 1}});
wsttiger's avatar
wsttiger committed
264
        p.add_instruction(migraph::op::softmax{}, x);
Paul's avatar
Paul committed
265
266
267
268
        return p;
    }
};

Paul's avatar
Paul committed
269
270
271
272
273
274
275
276
struct test_conv
{
    migraph::program create_program() const
    {
        migraph::program p;
        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}});
wsttiger's avatar
wsttiger committed
277
        p.add_instruction(migraph::op::convolution{}, input, weights);
Paul's avatar
Paul committed
278
279
280
281
        return p;
    }
};

Paul's avatar
Paul committed
282
283
284
285
286
287
288
289
290
struct test_conv2
{
    migraph::program create_program() const
    {
        migraph::program p;
        auto input =
            p.add_parameter("x", migraph::shape{migraph::shape::float_type, {1, 512, 28, 28}});
        auto weights =
            p.add_parameter("w", migraph::shape{migraph::shape::float_type, {256, 512, 1, 1}});
wsttiger's avatar
wsttiger committed
291
        p.add_instruction(migraph::op::convolution{{0, 0}, {1, 1}, {1, 1}}, input, weights);
Paul's avatar
Paul committed
292
293
294
295
        return p;
    }
};

Paul's avatar
Paul committed
296
struct test_conv_relu
Paul's avatar
Paul committed
297
{
Paul's avatar
Paul committed
298
    migraph::program create_program() const
Paul's avatar
Paul committed
299
    {
Paul's avatar
Paul committed
300
        migraph::program p;
Paul's avatar
Paul committed
301
302
303
        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}});
wsttiger's avatar
wsttiger committed
304
305
        auto conv = p.add_instruction(migraph::op::convolution{}, input, weights);
        p.add_instruction(migraph::op::activation{"relu"}, conv);
Paul's avatar
Paul committed
306
307
308
309
        return p;
    }
};

Paul's avatar
Paul committed
310
311
312
313
314
struct test_add_relu
{
    migraph::program create_program() const
    {
        migraph::program p;
Paul's avatar
Paul committed
315
316
        auto x   = p.add_parameter("x", migraph::shape{migraph::shape::float_type, {4, 3, 3, 3}});
        auto y   = p.add_parameter("y", migraph::shape{migraph::shape::float_type, {4, 3, 3, 3}});
wsttiger's avatar
wsttiger committed
317
318
        auto add = p.add_instruction(migraph::op::add{}, x, y);
        p.add_instruction(migraph::op::activation{"relu"}, add);
Paul's avatar
Paul committed
319
320
321
322
        return p;
    }
};

Paul's avatar
Paul committed
323
324
struct test_conv_pooling
{
Paul's avatar
Paul committed
325
    migraph::program create_program() const
Paul's avatar
Paul committed
326
    {
Paul's avatar
Paul committed
327
        migraph::program p;
Paul's avatar
Paul committed
328
329
330
331
        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}});
wsttiger's avatar
wsttiger committed
332
333
334
        auto conv    = p.add_instruction(migraph::op::convolution{}, input, weights);
        auto pooling = p.add_instruction(migraph::op::pooling{"max"}, conv);
        p.add_instruction(migraph::op::activation{"relu"}, pooling);
Paul's avatar
Paul committed
335
336
337
338
        return p;
    }
};

Paul's avatar
Paul committed
339
340
struct test_gemm
{
Paul's avatar
Paul committed
341
    migraph::program create_program() const
Paul's avatar
Paul committed
342
    {
Paul's avatar
Paul committed
343
344
345
        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}});
wsttiger's avatar
wsttiger committed
346
        p.add_instruction(migraph::op::gemm{}, a, b);
Paul's avatar
Paul committed
347
348
349
350
        return p;
    }
};

Paul's avatar
Paul committed
351
352
353
354
355
356
357
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}});
wsttiger's avatar
wsttiger committed
358
        p.add_instruction(migraph::op::gemm{}, a, b);
Paul's avatar
Paul committed
359
360
361
362
        return p;
    }
};

363
364
365
366
367
struct test_gemm_transposeb
{
    migraph::program create_program() const
    {
        migraph::program p;
Paul's avatar
Paul committed
368
369
        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}});
wsttiger's avatar
wsttiger committed
370
371
        auto bt = p.add_instruction(migraph::op::transpose{{1, 0}}, b);
        p.add_instruction(migraph::op::gemm{}, a, bt);
372
373
374
375
376
377
378
379
380
        return p;
    }
};

struct test_gemm_transposea
{
    migraph::program create_program() const
    {
        migraph::program p;
Paul's avatar
Paul committed
381
382
        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}});
wsttiger's avatar
wsttiger committed
383
384
        auto at = p.add_instruction(migraph::op::transpose{{1, 0}}, a);
        p.add_instruction(migraph::op::gemm{}, at, b);
385
386
387
388
389
390
391
392
393
        return p;
    }
};

struct test_gemm_transposeab
{
    migraph::program create_program() const
    {
        migraph::program p;
Paul's avatar
Paul committed
394
395
        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}});
wsttiger's avatar
wsttiger committed
396
397
398
        auto at = p.add_instruction(migraph::op::transpose{{1, 0}}, a);
        auto bt = p.add_instruction(migraph::op::transpose{{1, 0}}, b);
        p.add_instruction(migraph::op::gemm{}, at, bt);
399
400
401
402
        return p;
    }
};

403
404
405
406
407
struct test_contiguous
{
    migraph::program create_program() const
    {
        migraph::program p;
408
        migraph::shape s{migraph::shape::float_type, {4, 4, 4, 3}, {48, 4, 1, 16}};
409
        auto x = p.add_parameter("x", s);
wsttiger's avatar
wsttiger committed
410
        p.add_instruction(migraph::op::contiguous{}, x);
Paul's avatar
Paul committed
411
        EXPECT(p.get_shape().standard());
412
413
414
415
        return p;
    }
};

416
struct test_transpose
417
{
418
419
420
421
422
423
    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};
wsttiger's avatar
wsttiger committed
424
425
        auto l                    = p.add_instruction(migraph::op::transpose{perm}, x);
        p.add_instruction(migraph::op::contiguous{}, l);
426
427
428
        return p;
    }
};
429

Paul's avatar
Paul committed
430
431
432
433
434
435
436
437
438
439
440
441
442
443
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);
Paul's avatar
Paul committed
444
445
446
447
        auto scale    = p.add_literal(migraph::abs(migraph::generate_literal(vars, 1)));
        auto bias     = p.add_literal(migraph::abs(migraph::generate_literal(vars, 2)));
        auto mean     = p.add_literal(migraph::abs(migraph::generate_literal(vars, 3)));
        auto variance = p.add_literal(migraph::abs(migraph::generate_literal(vars, 4)));
wsttiger's avatar
wsttiger committed
448
        p.add_instruction(migraph::op::batch_norm_inference{}, x, scale, bias, mean, variance);
Paul's avatar
Paul committed
449
450
451
452
        return p;
    }
};

wsttiger's avatar
wsttiger committed
453
454
455
456
457
458
459
460
461
462
463
464
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}};
465
        migraph::shape vars{migraph::shape::float_type, {channels}};
wsttiger's avatar
wsttiger committed
466
        auto x        = p.add_parameter("x", s);
Paul's avatar
Paul committed
467
468
469
470
        auto scale    = p.add_literal(migraph::abs(migraph::generate_literal(vars, 1)));
        auto bias     = p.add_literal(migraph::abs(migraph::generate_literal(vars, 2)));
        auto mean     = p.add_literal(migraph::abs(migraph::generate_literal(vars, 3)));
        auto variance = p.add_literal(migraph::abs(migraph::generate_literal(vars, 4)));
wsttiger's avatar
wsttiger committed
471
        p.add_instruction(migraph::op::batch_norm_inference{}, x, scale, bias, mean, variance);
wsttiger's avatar
wsttiger committed
472
473
474
475
        return p;
    }
};

Paul's avatar
Paul committed
476
477
478
479
480
481
482
483
484
485
486
struct test_conv_bn
{
    migraph::program create_program() const
    {
        migraph::program p;

        migraph::shape xs{migraph::shape::float_type, {1, 3, 224, 224}};
        migraph::shape ws{migraph::shape::float_type, {64, 3, 7, 7}};
        migraph::shape vars{migraph::shape::float_type, {64}};
        auto x        = p.add_parameter("x", xs);
        auto w        = p.add_parameter("w", ws);
wsttiger's avatar
wsttiger committed
487
        auto conv     = p.add_instruction(migraph::op::convolution{{3, 3}, {2, 2}, {1, 1}}, x, w);
Paul's avatar
Paul committed
488
489
490
491
        auto scale    = p.add_literal(migraph::abs(migraph::generate_literal(vars, 1)));
        auto bias     = p.add_literal(migraph::abs(migraph::generate_literal(vars, 2)));
        auto mean     = p.add_literal(migraph::abs(migraph::generate_literal(vars, 3)));
        auto variance = p.add_literal(migraph::abs(migraph::generate_literal(vars, 4)));
wsttiger's avatar
wsttiger committed
492
        p.add_instruction(migraph::op::batch_norm_inference{}, conv, scale, bias, mean, variance);
Paul's avatar
Paul committed
493
494
495
496
        return p;
    }
};

Paul's avatar
Paul committed
497
498
499
500
501
502
503
504
struct test_conv_bn_relu_pooling
{
    migraph::program create_program() const
    {
        migraph::program p;

        migraph::shape xs{migraph::shape::float_type, {1, 3, 224, 224}};
        migraph::shape ws{migraph::shape::float_type, {64, 3, 7, 7}};
Paul's avatar
Paul committed
505
        migraph::shape vars{migraph::shape::float_type, {64}};
Paul's avatar
Paul committed
506
507
        auto x        = p.add_parameter("x", xs);
        auto w        = p.add_parameter("w", ws);
wsttiger's avatar
wsttiger committed
508
        auto conv     = p.add_instruction(migraph::op::convolution{{3, 3}, {2, 2}, {1, 1}}, x, w);
Paul's avatar
Paul committed
509
510
511
512
        auto scale    = p.add_literal(migraph::abs(migraph::generate_literal(vars, 1)));
        auto bias     = p.add_literal(migraph::abs(migraph::generate_literal(vars, 2)));
        auto mean     = p.add_literal(migraph::abs(migraph::generate_literal(vars, 3)));
        auto variance = p.add_literal(migraph::abs(migraph::generate_literal(vars, 4)));
wsttiger's avatar
wsttiger committed
513
514
        auto bn       = p.add_instruction(
            migraph::op::batch_norm_inference{}, conv, scale, bias, mean, variance);
wsttiger's avatar
wsttiger committed
515
516
        auto relu = p.add_instruction(migraph::op::activation{"relu"}, bn);
        p.add_instruction(migraph::op::pooling{"average", {1, 1}, {2, 2}, {3, 3}}, relu);
Paul's avatar
Paul committed
517
518
519
520
        return p;
    }
};

521
522
523
524
525
526
527
528
529
530
531
532
533
534
535
536
537
538
539
540
541
542
543
544
545
546
547
548
549
550
551
552
553
554
struct test_concat
{
    migraph::program create_program() const
    {
        migraph::program p;
        std::size_t axis       = 1;
        migraph::shape s0{migraph::shape::int32_type, {2, 2}};
        migraph::shape s1{migraph::shape::int32_type, {2, 3}};
        migraph::shape s2{migraph::shape::int32_type, {2, 1}};
        auto l0 = p.add_parameter("x", s0);
        auto l1 = p.add_parameter("y", s1);
        auto l2 = p.add_parameter("z", s2);
        p.add_instruction(migraph::op::concat{axis}, l0, l1, l2);
        return p;
    }
};

struct test_concat2
{
    migraph::program create_program() const
    {
        migraph::program p;
        std::size_t axis       = 0;
        migraph::shape s0{migraph::shape::int32_type, {2, 2}};
        migraph::shape s1{migraph::shape::int32_type, {3, 2}};
        migraph::shape s2{migraph::shape::int32_type, {1, 2}};
        auto l0 = p.add_parameter("x", s0);
        auto l1 = p.add_parameter("y", s1);
        auto l2 = p.add_parameter("z", s2);
        p.add_instruction(migraph::op::concat{axis}, l0, l1, l2);
        return p;
    }
};

Paul's avatar
Paul committed
555
556
struct test_conv_bn_relu_pooling2
{
Paul's avatar
Paul committed
557
558
    static migraph::instruction_ref
    add_bn(migraph::program& p, migraph::instruction_ref x, std::size_t channels)
Paul's avatar
Paul committed
559
560
    {
        migraph::shape vars{migraph::shape::float_type, {channels}};
Paul's avatar
Paul committed
561
562
563
564
        auto scale    = p.add_literal(migraph::abs(migraph::generate_literal(vars, 1 + channels)));
        auto bias     = p.add_literal(migraph::abs(migraph::generate_literal(vars, 2 + channels)));
        auto mean     = p.add_literal(migraph::abs(migraph::generate_literal(vars, 3 + channels)));
        auto variance = p.add_literal(migraph::abs(migraph::generate_literal(vars, 4 + channels)));
wsttiger's avatar
wsttiger committed
565
566
        return p.add_instruction(
            migraph::op::batch_norm_inference{}, x, scale, bias, mean, variance);
Paul's avatar
Paul committed
567
568
569
570
571
572
573
574
575
    }
    migraph::program create_program() const
    {
        migraph::program p;

        migraph::shape xs1{migraph::shape::float_type, {1, 512, 7, 7}};
        migraph::shape xs2{migraph::shape::float_type, {1, 1024, 14, 14}};
        migraph::shape ws1{migraph::shape::float_type, {2048, 512, 1, 1}};
        migraph::shape ws2{migraph::shape::float_type, {2048, 1024, 1, 1}};
Paul's avatar
Paul committed
576
577
        auto x1    = p.add_parameter("x1", xs1);
        auto w1    = p.add_parameter("w1", ws1);
wsttiger's avatar
wsttiger committed
578
        auto conv1 = p.add_instruction(migraph::op::convolution{{0, 0}, {1, 1}, {1, 1}}, x1, w1);
Paul's avatar
Paul committed
579
580
581
        auto bn1   = add_bn(p, conv1, 2048);
        auto x2    = p.add_parameter("x2", xs2);
        auto w2    = p.add_parameter("w2", ws2);
wsttiger's avatar
wsttiger committed
582
        auto conv2 = p.add_instruction(migraph::op::convolution{{0, 0}, {2, 2}, {1, 1}}, x2, w2);
Paul's avatar
Paul committed
583
        auto bn2   = add_bn(p, conv2, 2048);
wsttiger's avatar
wsttiger committed
584
585
586
        auto add   = p.add_instruction(migraph::op::add{}, bn1, bn2);
        auto relu  = p.add_instruction(migraph::op::activation{"relu"}, add);
        p.add_instruction(migraph::op::pooling{"average", {1, 1}, {2, 2}, {3, 3}}, relu);
Paul's avatar
Paul committed
587
588
589
590
        return p;
    }
};

Paul's avatar
Paul committed
591
592
int main()
{
593
594
    verify_program<test_concat>();
    verify_program<test_concat2>();
Paul's avatar
Paul committed
595
    verify_program<test_add>();
Paul's avatar
Paul committed
596
    verify_program<test_add_broadcast>();
Paul's avatar
Paul committed
597
    verify_program<test_add_broadcast2>();
Paul's avatar
Latest  
Paul committed
598
599
    verify_program<test_add_broadcast3>();
    verify_program<test_add_broadcast4>();
Paul's avatar
Paul committed
600
    verify_program<test_add_broadcast5>();
Paul's avatar
Paul committed
601
    verify_program<test_softmax>();
Paul's avatar
Paul committed
602
    verify_program<test_softmax2>();
Paul's avatar
Paul committed
603
    verify_program<test_conv>();
Paul's avatar
Paul committed
604
    verify_program<test_conv2>();
Paul's avatar
Paul committed
605
    verify_program<test_conv_relu>();
Paul's avatar
Paul committed
606
    verify_program<test_add_relu>();
Paul's avatar
Paul committed
607
608
    verify_program<test_conv_pooling>();
    verify_program<test_gemm>();
Paul's avatar
Paul committed
609
    // verify_program<test_gemm_ld>();
610
611
612
    verify_program<test_gemm_transposeb>();
    verify_program<test_gemm_transposea>();
    verify_program<test_gemm_transposeab>();
613
614
    verify_program<test_contiguous>();
    verify_program<test_transpose>();
615
    verify_program<test_batchnorm_inference>();
Paul's avatar
Paul committed
616
    verify_program<test_batchnorm_inference_2>();
Paul's avatar
Paul committed
617
    verify_program<test_conv_bn>();
Paul's avatar
Paul committed
618
    verify_program<test_conv_bn_relu_pooling>();
Paul's avatar
Paul committed
619
    verify_program<test_conv_bn_relu_pooling2>();
Paul's avatar
Paul committed
620
}