miopen.cpp 23.8 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
    // std::cout << migraph::get_type_name<V>() << std::endl;
Paul's avatar
Paul committed
133
134
135
136
    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
137
    bool passed    = verify_args(migraph::get_type_name<V>(), cpu_arg_f.get(), gpu_arg);
Paul's avatar
Paul committed
138
139
140
141
142
143
144
145
146
    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
147
    std::set_terminate(nullptr);
Paul's avatar
Paul committed
148
149
}

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

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

Khalique's avatar
Khalique committed
178
179
180
181
182
183
184
185
186
187
188
189
190
191
192
193
194
195
196
197
198
199
200
201
202
203
204
struct test_mul
{
    migraph::program create_program() const
    {
        migraph::program p;
        migraph::shape s{migraph::shape::float_type, {3}};
        auto x = p.add_parameter("x", s);
        auto y = p.add_parameter("y", s);
        p.add_instruction(migraph::op::mul{}, x, y);
        return p;
    }
};

struct test_scale
{
    migraph::program create_program() const
    {
        migraph::program p;
        migraph::shape s{migraph::shape::float_type, {3}};
        auto x = p.add_parameter("x", s);
        auto y = p.add_parameter("y", migraph::shape::float_type);
        auto scale = p.add_instruction(migraph::op::scalar{s}, y);
        p.add_instruction(migraph::op::mul{}, x, scale);
        return p;
    }
};

Paul's avatar
Paul committed
205
206
207
208
209
210
struct test_triadd
{
    migraph::program create_program() const
    {
        migraph::program p;
        migraph::shape s{migraph::shape::float_type, {3}};
Paul's avatar
Paul committed
211
212
213
        auto x   = p.add_parameter("x", s);
        auto y   = p.add_parameter("y", s);
        auto z   = p.add_parameter("z", s);
Paul's avatar
Paul committed
214
215
216
217
218
219
220
221
222
223
224
225
226
        auto sum = p.add_instruction(migraph::op::add{}, x, y);
        p.add_instruction(migraph::op::add{}, sum, z);
        return p;
    }
};

struct test_triadd2
{
    migraph::program create_program() const
    {
        migraph::program p;
        migraph::shape s{migraph::shape::float_type, {2, 3}};
        migraph::shape b{migraph::shape::float_type, {3}};
Paul's avatar
Paul committed
227
228
229
230
        auto x   = p.add_parameter("x", s);
        auto y   = p.add_parameter("y", s);
        auto z   = p.add_parameter("z", b);
        auto zb  = p.add_instruction(migraph::op::broadcast{1, s}, z);
Paul's avatar
Paul committed
231
232
233
234
235
236
        auto sum = p.add_instruction(migraph::op::add{}, x, y);
        p.add_instruction(migraph::op::add{}, sum, zb);
        return p;
    }
};

Paul's avatar
Paul committed
237
238
struct test_add_broadcast
{
Paul's avatar
Paul committed
239
    migraph::program create_program() const
Paul's avatar
Paul committed
240
    {
Paul's avatar
Paul committed
241
242
243
244
        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
245
        auto by = p.add_instruction(migraph::op::broadcast{0, x->get_shape()}, y);
wsttiger's avatar
wsttiger committed
246
        p.add_instruction(migraph::op::add{}, x, by);
Paul's avatar
Paul committed
247
248
249
250
        return p;
    }
};

Paul's avatar
Paul committed
251
252
253
254
255
256
257
258
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
259
        auto by = p.add_instruction(migraph::op::broadcast{1, x->get_shape()}, y);
wsttiger's avatar
wsttiger committed
260
        p.add_instruction(migraph::op::add{}, x, by);
Paul's avatar
Paul committed
261
262
263
264
        return p;
    }
};

Paul's avatar
Latest  
Paul committed
265
266
267
268
269
270
271
272
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
273
        auto by = p.add_instruction(migraph::op::broadcast{1, x->get_shape()}, y);
wsttiger's avatar
wsttiger committed
274
        p.add_instruction(migraph::op::add{}, x, by);
Paul's avatar
Latest  
Paul committed
275
276
277
278
279
280
281
282
283
284
285
286
        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
287
        auto by = p.add_instruction(migraph::op::broadcast{1, x->get_shape()}, y);
wsttiger's avatar
wsttiger committed
288
        p.add_instruction(migraph::op::add{}, x, by);
Paul's avatar
Latest  
Paul committed
289
290
291
292
        return p;
    }
};

Paul's avatar
Paul committed
293
294
295
296
297
298
299
300
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
301
        auto by = p.add_instruction(migraph::op::broadcast{1, x->get_shape()}, y);
wsttiger's avatar
wsttiger committed
302
        p.add_instruction(migraph::op::add{}, x, by);
Paul's avatar
Paul committed
303
304
305
306
        return p;
    }
};

Paul's avatar
Paul committed
307
308
309
310
311
312
struct test_triadd_broadcast
{
    migraph::program create_program() const
    {
        migraph::program p;
        migraph::shape s{migraph::shape::float_type, {3}};
Paul's avatar
Paul committed
313
314
315
316
        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 z   = p.add_parameter("z", {migraph::shape::float_type, {2, 2, 3}});
        auto by  = p.add_instruction(migraph::op::broadcast{0, x->get_shape()}, y);
Paul's avatar
Paul committed
317
318
319
320
321
322
        auto sum = p.add_instruction(migraph::op::add{}, x, by);
        p.add_instruction(migraph::op::add{}, sum, z);
        return p;
    }
};

Paul's avatar
Paul committed
323
324
325
326
327
struct test_softmax
{
    migraph::program create_program() const
    {
        migraph::program p;
Paul's avatar
Paul committed
328
        auto x = p.add_parameter("x", migraph::shape{migraph::shape::float_type, {5, 3, 4, 2}});
wsttiger's avatar
wsttiger committed
329
        p.add_instruction(migraph::op::softmax{}, x);
Paul's avatar
Paul committed
330
331
332
333
334
335
336
337
338
        return p;
    }
};

struct test_softmax2
{
    migraph::program create_program() const
    {
        migraph::program p;
Paul's avatar
Paul committed
339
        auto x = p.add_parameter("x", migraph::shape{migraph::shape::float_type, {1, 1000, 1, 1}});
wsttiger's avatar
wsttiger committed
340
        p.add_instruction(migraph::op::softmax{}, x);
Paul's avatar
Paul committed
341
342
343
344
        return p;
    }
};

Paul's avatar
Paul committed
345
346
347
348
349
350
351
352
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
353
        p.add_instruction(migraph::op::convolution{}, input, weights);
Paul's avatar
Paul committed
354
355
356
357
        return p;
    }
};

Paul's avatar
Paul committed
358
359
360
361
362
363
364
365
366
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
367
        p.add_instruction(migraph::op::convolution{{0, 0}, {1, 1}, {1, 1}}, input, weights);
Paul's avatar
Paul committed
368
369
370
371
        return p;
    }
};

Paul's avatar
Paul committed
372
struct test_conv_relu
Paul's avatar
Paul committed
373
{
Paul's avatar
Paul committed
374
    migraph::program create_program() const
Paul's avatar
Paul committed
375
    {
Paul's avatar
Paul committed
376
        migraph::program p;
Paul's avatar
Paul committed
377
378
379
        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
380
381
        auto conv = p.add_instruction(migraph::op::convolution{}, input, weights);
        p.add_instruction(migraph::op::activation{"relu"}, conv);
Paul's avatar
Paul committed
382
383
384
385
        return p;
    }
};

Paul's avatar
Paul committed
386
387
388
389
390
struct test_add_relu
{
    migraph::program create_program() const
    {
        migraph::program p;
Paul's avatar
Paul committed
391
392
        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
393
394
        auto add = p.add_instruction(migraph::op::add{}, x, y);
        p.add_instruction(migraph::op::activation{"relu"}, add);
Paul's avatar
Paul committed
395
396
397
398
        return p;
    }
};

399
400
401
402
403
404
405
406
407
408
409
struct test_leaky_relu
{
    migraph::program create_program() const
    {
        migraph::program p;
        auto x = p.add_parameter("x", migraph::shape{migraph::shape::float_type, {4, 3, 3, 3}});
        p.add_instruction(migraph::op::leaky_relu{0.01}, x);
        return p;
    }
};

Paul's avatar
Paul committed
410
411
struct test_conv_pooling
{
Paul's avatar
Paul committed
412
    migraph::program create_program() const
Paul's avatar
Paul committed
413
    {
Paul's avatar
Paul committed
414
        migraph::program p;
Paul's avatar
Paul committed
415
416
417
418
        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
419
420
421
        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
422
423
424
425
        return p;
    }
};

Paul's avatar
Paul committed
426
427
struct test_gemm
{
Paul's avatar
Paul committed
428
    migraph::program create_program() const
Paul's avatar
Paul committed
429
    {
Paul's avatar
Paul committed
430
431
432
        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}});
Shucai Xiao's avatar
Shucai Xiao committed
433
        p.add_instruction(migraph::op::dot{}, a, b);
Paul's avatar
Paul committed
434
435
436
437
        return p;
    }
};

Paul's avatar
Paul committed
438
439
440
441
442
443
444
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}});
Shucai Xiao's avatar
Shucai Xiao committed
445
        p.add_instruction(migraph::op::dot{}, a, b);
Paul's avatar
Paul committed
446
447
448
449
        return p;
    }
};

450
451
452
453
454
struct test_gemm_transposeb
{
    migraph::program create_program() const
    {
        migraph::program p;
Paul's avatar
Paul committed
455
456
        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
457
        auto bt = p.add_instruction(migraph::op::transpose{{1, 0}}, b);
Shucai Xiao's avatar
Shucai Xiao committed
458
        p.add_instruction(migraph::op::dot{}, a, bt);
459
460
461
462
463
464
465
466
467
        return p;
    }
};

struct test_gemm_transposea
{
    migraph::program create_program() const
    {
        migraph::program p;
Paul's avatar
Paul committed
468
469
        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
470
        auto at = p.add_instruction(migraph::op::transpose{{1, 0}}, a);
Shucai Xiao's avatar
Shucai Xiao committed
471
        p.add_instruction(migraph::op::dot{}, at, b);
472
473
474
475
476
477
478
479
480
        return p;
    }
};

struct test_gemm_transposeab
{
    migraph::program create_program() const
    {
        migraph::program p;
Paul's avatar
Paul committed
481
482
        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
483
484
        auto at = p.add_instruction(migraph::op::transpose{{1, 0}}, a);
        auto bt = p.add_instruction(migraph::op::transpose{{1, 0}}, b);
Shucai Xiao's avatar
Shucai Xiao committed
485
        p.add_instruction(migraph::op::dot{}, at, bt);
486
487
488
489
        return p;
    }
};

490
491
492
493
494
struct test_contiguous
{
    migraph::program create_program() const
    {
        migraph::program p;
495
        migraph::shape s{migraph::shape::float_type, {4, 4, 4, 3}, {48, 4, 1, 16}};
496
        auto x = p.add_parameter("x", s);
wsttiger's avatar
wsttiger committed
497
        p.add_instruction(migraph::op::contiguous{}, x);
Paul's avatar
Paul committed
498
        EXPECT(p.get_shape().standard());
499
500
501
502
        return p;
    }
};

503
struct test_transpose
504
{
505
506
507
508
509
510
    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
511
512
        auto l                    = p.add_instruction(migraph::op::transpose{perm}, x);
        p.add_instruction(migraph::op::contiguous{}, l);
513
514
515
        return p;
    }
};
516

Paul's avatar
Paul committed
517
518
519
520
521
522
523
524
525
526
527
528
529
530
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
531
532
533
534
        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
535
        p.add_instruction(migraph::op::batch_norm_inference{}, x, scale, bias, mean, variance);
Paul's avatar
Paul committed
536
537
538
539
        return p;
    }
};

wsttiger's avatar
wsttiger committed
540
541
542
543
544
545
546
547
548
549
550
551
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}};
552
        migraph::shape vars{migraph::shape::float_type, {channels}};
wsttiger's avatar
wsttiger committed
553
        auto x        = p.add_parameter("x", s);
Paul's avatar
Paul committed
554
555
556
557
        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
558
        p.add_instruction(migraph::op::batch_norm_inference{}, x, scale, bias, mean, variance);
wsttiger's avatar
wsttiger committed
559
560
561
562
        return p;
    }
};

Paul's avatar
Paul committed
563
564
565
566
567
568
569
570
571
572
573
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
574
        auto conv     = p.add_instruction(migraph::op::convolution{{3, 3}, {2, 2}, {1, 1}}, x, w);
Paul's avatar
Paul committed
575
576
577
578
        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
579
        p.add_instruction(migraph::op::batch_norm_inference{}, conv, scale, bias, mean, variance);
Paul's avatar
Paul committed
580
581
582
583
        return p;
    }
};

Paul's avatar
Paul committed
584
585
586
587
588
589
590
591
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
592
        migraph::shape vars{migraph::shape::float_type, {64}};
Paul's avatar
Paul committed
593
594
        auto x        = p.add_parameter("x", xs);
        auto w        = p.add_parameter("w", ws);
wsttiger's avatar
wsttiger committed
595
        auto conv     = p.add_instruction(migraph::op::convolution{{3, 3}, {2, 2}, {1, 1}}, x, w);
Paul's avatar
Paul committed
596
597
598
599
        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
600
601
        auto bn       = p.add_instruction(
            migraph::op::batch_norm_inference{}, conv, scale, bias, mean, variance);
wsttiger's avatar
wsttiger committed
602
603
        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
604
605
606
607
        return p;
    }
};

608
609
610
611
612
struct test_concat
{
    migraph::program create_program() const
    {
        migraph::program p;
wsttiger's avatar
wsttiger committed
613
        std::size_t axis = 1;
614
615
616
617
618
619
620
621
622
623
624
625
626
627
628
629
        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;
wsttiger's avatar
wsttiger committed
630
        std::size_t axis = 0;
631
632
633
634
635
636
637
638
639
640
641
        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
642
643
struct test_conv_bn_relu_pooling2
{
Paul's avatar
Paul committed
644
645
    static migraph::instruction_ref
    add_bn(migraph::program& p, migraph::instruction_ref x, std::size_t channels)
Paul's avatar
Paul committed
646
647
    {
        migraph::shape vars{migraph::shape::float_type, {channels}};
Paul's avatar
Paul committed
648
649
650
651
        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
652
653
        return p.add_instruction(
            migraph::op::batch_norm_inference{}, x, scale, bias, mean, variance);
Paul's avatar
Paul committed
654
655
656
657
658
659
660
661
662
    }
    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
663
664
        auto x1    = p.add_parameter("x1", xs1);
        auto w1    = p.add_parameter("w1", ws1);
wsttiger's avatar
wsttiger committed
665
        auto conv1 = p.add_instruction(migraph::op::convolution{{0, 0}, {1, 1}, {1, 1}}, x1, w1);
Paul's avatar
Paul committed
666
667
668
        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
669
        auto conv2 = p.add_instruction(migraph::op::convolution{{0, 0}, {2, 2}, {1, 1}}, x2, w2);
Paul's avatar
Paul committed
670
        auto bn2   = add_bn(p, conv2, 2048);
wsttiger's avatar
wsttiger committed
671
672
673
        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
674
675
676
677
        return p;
    }
};

Paul's avatar
Paul committed
678
679
int main()
{
680
681
    verify_program<test_concat>();
    verify_program<test_concat2>();
Paul's avatar
Paul committed
682
    verify_program<test_add>();
Khalique's avatar
Khalique committed
683
684
    verify_program<test_mul>();
    verify_program<test_scale>();
Paul's avatar
Paul committed
685
686
    verify_program<test_triadd>();
    verify_program<test_triadd2>();
Paul's avatar
Paul committed
687
    verify_program<test_add_broadcast>();
Paul's avatar
Paul committed
688
    verify_program<test_add_broadcast2>();
Paul's avatar
Latest  
Paul committed
689
690
    verify_program<test_add_broadcast3>();
    verify_program<test_add_broadcast4>();
Paul's avatar
Paul committed
691
    verify_program<test_add_broadcast5>();
Paul's avatar
Paul committed
692
    verify_program<test_triadd_broadcast>();
Paul's avatar
Paul committed
693
    verify_program<test_softmax>();
Paul's avatar
Paul committed
694
    verify_program<test_softmax2>();
Paul's avatar
Paul committed
695
    verify_program<test_conv>();
Paul's avatar
Paul committed
696
    verify_program<test_conv2>();
Paul's avatar
Paul committed
697
    verify_program<test_conv_relu>();
Paul's avatar
Paul committed
698
    verify_program<test_add_relu>();
699
    verify_program<test_leaky_relu>();
Paul's avatar
Paul committed
700
701
    verify_program<test_conv_pooling>();
    verify_program<test_gemm>();
702
    // verify_program<test_gemm_ld>();
703
704
705
    verify_program<test_gemm_transposeb>();
    verify_program<test_gemm_transposea>();
    verify_program<test_gemm_transposeab>();
706
707
    verify_program<test_contiguous>();
    verify_program<test_transpose>();
708
    verify_program<test_batchnorm_inference>();
Paul's avatar
Paul committed
709
    verify_program<test_batchnorm_inference_2>();
Paul's avatar
Paul committed
710
    verify_program<test_conv_bn>();
Paul's avatar
Paul committed
711
    verify_program<test_conv_bn_relu_pooling>();
Paul's avatar
Paul committed
712
    verify_program<test_conv_bn_relu_pooling2>();
Paul's avatar
Paul committed
713
}