miopen.cpp 19.7 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
    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);
    }
38
    return std::async(std::launch::deferred, std::forward<Function>(f));
Paul's avatar
Paul committed
39
40
}

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

struct test_gemm_transposeab
{
    migraph::program create_program() const
    {
        migraph::program p;
Paul's avatar
Paul committed
393
394
        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
395
396
397
        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);
398
399
400
401
        return p;
    }
};

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

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

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

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

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

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

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

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