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

Paul's avatar
Paul committed
44
45
struct auto_print
{
Paul's avatar
Paul committed
46
47
48
49
50
51
52
53
54
55
56
57
58
59
60
61
62
63
64
    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
65
    static std::array<std::function<void()>, 2> handlers;
Paul's avatar
Paul committed
66
    int index;
Paul's avatar
Paul committed
67
    template <class T>
Paul's avatar
Paul committed
68
    auto_print(T& x, int i) : index(i)
Paul's avatar
Paul committed
69
    {
Paul's avatar
Paul committed
70
        handlers[index] = [&x] { std::cout << x << std::endl; };
Paul's avatar
Paul committed
71
    }
Paul's avatar
Paul committed
72

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

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

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

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

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

Paul's avatar
Paul committed
130
131
132
133
template <class V>
void verify_program()
{
    auto_print::set_terminate_handler(migraph::get_type_name<V>());
Paul's avatar
Paul committed
134
135
136
137
    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
138
    bool passed    = verify_args(migraph::get_type_name<V>(), cpu_arg_f.get(), gpu_arg);
Paul's avatar
Paul committed
139
140
141
142
143
144
145
146
147
    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
148
    std::set_terminate(nullptr);
Paul's avatar
Paul committed
149
150
}

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

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

struct test_add_broadcast
{
Paul's avatar
Paul committed
181
    migraph::program create_program() const
Paul's avatar
Paul committed
182
    {
Paul's avatar
Paul committed
183
184
185
186
187
188
        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
189
190
191
192
        return p;
    }
};

Paul's avatar
Paul committed
193
194
195
196
197
198
199
200
201
202
203
204
205
206
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}});
        auto by = p.add_instruction(migraph::broadcast{1}, x, y);
        p.add_instruction(migraph::add{}, x, by);
        return p;
    }
};

Paul's avatar
Latest  
Paul committed
207
208
209
210
211
212
213
214
215
216
217
218
219
220
221
222
223
224
225
226
227
228
229
230
231
232
233
234
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}});
        auto by = p.add_instruction(migraph::broadcast{1}, x, y);
        p.add_instruction(migraph::add{}, x, by);
        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}});
        auto by = p.add_instruction(migraph::broadcast{1}, x, y);
        p.add_instruction(migraph::add{}, x, by);
        return p;
    }
};

Paul's avatar
Paul committed
235
236
237
238
239
240
241
242
243
244
245
246
247
248
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}});
        auto by = p.add_instruction(migraph::broadcast{1}, x, y);
        p.add_instruction(migraph::add{}, x, by);
        return p;
    }
};

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

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

Paul's avatar
Paul committed
271
272
273
274
275
276
277
278
279
280
281
282
283
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}});
        p.add_instruction(migraph::convolution{}, input, weights);
        return p;
    }
};

Paul's avatar
Paul committed
284
285
286
287
288
289
290
291
292
293
294
295
296
297
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}});
        p.add_instruction(migraph::convolution{{0, 0}, {1, 1}, {1, 1}}, input, weights);
        return p;
    }
};

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

Paul's avatar
Paul committed
312
313
314
315
316
struct test_add_relu
{
    migraph::program create_program() const
    {
        migraph::program p;
Paul's avatar
Paul committed
317
318
        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}});
Paul's avatar
Paul committed
319
320
321
322
323
324
        auto add = p.add_instruction(migraph::add{}, x, y);
        p.add_instruction(migraph::activation{"relu"}, add);
        return p;
    }
};

Paul's avatar
Paul committed
325
326
struct test_conv_pooling
{
Paul's avatar
Paul committed
327
    migraph::program create_program() const
Paul's avatar
Paul committed
328
    {
Paul's avatar
Paul committed
329
        migraph::program p;
Paul's avatar
Paul committed
330
331
332
333
        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
334
335
336
        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
337
338
339
340
        return p;
    }
};

Paul's avatar
Paul committed
341
342
struct test_gemm
{
Paul's avatar
Paul committed
343
    migraph::program create_program() const
Paul's avatar
Paul committed
344
    {
Paul's avatar
Paul committed
345
346
347
348
        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
349
350
351
352
        return p;
    }
};

Paul's avatar
Paul committed
353
354
355
356
357
358
359
360
361
362
363
364
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;
    }
};

365
366
367
368
369
struct test_gemm_transposeb
{
    migraph::program create_program() const
    {
        migraph::program p;
Paul's avatar
Paul committed
370
371
        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}});
372
373
374
375
376
377
378
379
380
381
382
        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
383
384
        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}});
385
386
387
388
389
390
391
392
393
394
395
        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
396
397
        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}});
398
399
400
401
402
403
404
        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;
    }
};

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

418
struct test_transpose
419
{
420
421
422
423
424
425
426
427
428
429
430
    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;
    }
};
431

Paul's avatar
Paul committed
432
433
434
435
436
437
438
439
440
441
442
443
444
445
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
446
447
448
449
450
        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)));
        p.add_instruction(migraph::batch_norm_inference{}, x, scale, bias, mean, variance);
Paul's avatar
Paul committed
451
452
453
454
        return p;
    }
};

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

Paul's avatar
Paul committed
478
479
480
481
482
483
484
485
486
487
488
489
490
491
492
493
494
495
496
497
498
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);
        auto conv     = p.add_instruction(migraph::convolution{{3, 3}, {2, 2}, {1, 1}}, x, w);
        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)));
        p.add_instruction(migraph::batch_norm_inference{}, conv, scale, bias, mean, variance);
        return p;
    }
};

Paul's avatar
Paul committed
499
500
501
502
503
504
505
506
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
507
        migraph::shape vars{migraph::shape::float_type, {64}};
Paul's avatar
Paul committed
508
509
510
        auto x        = p.add_parameter("x", xs);
        auto w        = p.add_parameter("w", ws);
        auto conv     = p.add_instruction(migraph::convolution{{3, 3}, {2, 2}, {1, 1}}, x, w);
Paul's avatar
Paul committed
511
512
513
514
        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)));
Paul's avatar
Paul committed
515
        auto bn =
Paul's avatar
Paul committed
516
            p.add_instruction(migraph::batch_norm_inference{}, conv, scale, bias, mean, variance);
Paul's avatar
Paul committed
517
518
519
520
521
522
        auto relu = p.add_instruction(migraph::activation{"relu"}, bn);
        p.add_instruction(migraph::pooling{"average", {1, 1}, {2, 2}, {3, 3}}, relu);
        return p;
    }
};

Paul's avatar
Paul committed
523
524
struct test_conv_bn_relu_pooling2
{
Paul's avatar
Paul committed
525
526
    static migraph::instruction_ref
    add_bn(migraph::program& p, migraph::instruction_ref x, std::size_t channels)
Paul's avatar
Paul committed
527
528
    {
        migraph::shape vars{migraph::shape::float_type, {channels}};
Paul's avatar
Paul committed
529
530
531
532
        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)));
Paul's avatar
Paul committed
533
534
535
536
537
538
539
540
541
542
        return p.add_instruction(migraph::batch_norm_inference{}, x, scale, bias, mean, variance);
    }
    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
543
544
545
546
547
548
549
550
551
552
        auto x1    = p.add_parameter("x1", xs1);
        auto w1    = p.add_parameter("w1", ws1);
        auto conv1 = p.add_instruction(migraph::convolution{{0, 0}, {1, 1}, {1, 1}}, x1, w1);
        auto bn1   = add_bn(p, conv1, 2048);
        auto x2    = p.add_parameter("x2", xs2);
        auto w2    = p.add_parameter("w2", ws2);
        auto conv2 = p.add_instruction(migraph::convolution{{0, 0}, {2, 2}, {1, 1}}, x2, w2);
        auto bn2   = add_bn(p, conv2, 2048);
        auto add   = p.add_instruction(migraph::add{}, bn1, bn2);
        auto relu  = p.add_instruction(migraph::activation{"relu"}, add);
Paul's avatar
Paul committed
553
554
555
556
557
        p.add_instruction(migraph::pooling{"average", {1, 1}, {2, 2}, {3, 3}}, relu);
        return p;
    }
};

Paul's avatar
Paul committed
558
559
int main()
{
Paul's avatar
Paul committed
560
    verify_program<test_add>();
Paul's avatar
Paul committed
561
    verify_program<test_add_broadcast>();
Paul's avatar
Paul committed
562
    verify_program<test_add_broadcast2>();
Paul's avatar
Latest  
Paul committed
563
564
    verify_program<test_add_broadcast3>();
    verify_program<test_add_broadcast4>();
Paul's avatar
Paul committed
565
    verify_program<test_add_broadcast5>();
Paul's avatar
Paul committed
566
    verify_program<test_softmax>();
Paul's avatar
Paul committed
567
    verify_program<test_softmax2>();
Paul's avatar
Paul committed
568
    verify_program<test_conv>();
Paul's avatar
Paul committed
569
    verify_program<test_conv2>();
Paul's avatar
Paul committed
570
    verify_program<test_conv_relu>();
Paul's avatar
Paul committed
571
    verify_program<test_add_relu>();
Paul's avatar
Paul committed
572
573
    verify_program<test_conv_pooling>();
    verify_program<test_gemm>();
Paul's avatar
Paul committed
574
    // verify_program<test_gemm_ld>();
575
576
577
    verify_program<test_gemm_transposeb>();
    verify_program<test_gemm_transposea>();
    verify_program<test_gemm_transposeab>();
578
579
    verify_program<test_contiguous>();
    verify_program<test_transpose>();
580
    verify_program<test_batchnorm_inference>();
Paul's avatar
Paul committed
581
    verify_program<test_batchnorm_inference_2>();
Paul's avatar
Paul committed
582
    verify_program<test_conv_bn>();
Paul's avatar
Paul committed
583
    verify_program<test_conv_bn_relu_pooling>();
Paul's avatar
Paul committed
584
    verify_program<test_conv_bn_relu_pooling2>();
Paul's avatar
Paul committed
585
}