miopen.cpp 33.2 KB
Newer Older
Paul's avatar
Paul committed
1

Paul's avatar
Paul committed
2
3
4
5
6
7
8
9
10
11
12
#include <migraphx/program.hpp>
#include <migraphx/operators.hpp>
#include <migraphx/generate.hpp>
#include <migraphx/cpu/target.hpp>
#include <migraphx/gpu/target.hpp>
#include <migraphx/gpu/miopen.hpp>
#include <migraphx/gpu/hip.hpp>
#include <migraphx/manage_ptr.hpp>
#include <migraphx/type_name.hpp>
#include <migraphx/verify_args.hpp>
#include <migraphx/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(migraphx::program& p, const migraphx::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, migraphx::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
migraphx::argument run_cpu(migraphx::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
104
    compile_check(p, migraphx::cpu::target{});
    migraphx::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
Paul committed
107
        m[x.first] = migraphx::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
migraphx::argument run_gpu(migraphx::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
119
    compile_check(p, migraphx::gpu::target{});
    migraphx::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
Paul committed
122
123
        m[x.first] =
            migraphx::gpu::to_gpu(migraphx::generate_argument(x.second, get_hash(x.first)));
Paul's avatar
Paul committed
124
    }
Paul's avatar
Paul committed
125
    EXPECT(bool{m.find("output") != m.end()});
Paul's avatar
Paul committed
126
    return migraphx::gpu::from_gpu(p.eval(m));
Paul's avatar
Paul committed
127
128
}

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

Paul's avatar
Paul committed
152
153
struct test_literals
{
Paul's avatar
Paul committed
154
    migraphx::program create_program() const
Paul's avatar
Paul committed
155
    {
Paul's avatar
Paul committed
156
        migraphx::program p;
Paul's avatar
Paul committed
157
        auto input = p.add_literal(
Paul's avatar
Paul committed
158
            generate_literal(migraphx::shape{migraphx::shape::float_type, {4, 3, 3, 3}}));
Paul's avatar
Paul committed
159
        auto weights = p.add_literal(
Paul's avatar
Paul committed
160
161
162
            generate_literal(migraphx::shape{migraphx::shape::float_type, {4, 3, 3, 3}}));
        auto conv = p.add_instruction(migraphx::op::convolution{}, input, weights);
        p.add_instruction(migraphx::op::relu{}, conv);
Paul's avatar
Paul committed
163
164
165
166
        return p;
    }
};

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

Paul's avatar
Paul committed
180
181
struct test_add_half
{
Paul's avatar
Paul committed
182
    migraphx::program create_program() const
Paul's avatar
Paul committed
183
    {
Paul's avatar
Paul committed
184
185
        migraphx::program p;
        migraphx::shape s{migraphx::shape::half_type, {3}};
Paul's avatar
Paul committed
186
187
        auto x = p.add_parameter("x", s);
        auto y = p.add_parameter("y", s);
Paul's avatar
Paul committed
188
        p.add_instruction(migraphx::op::add{}, x, y);
Paul's avatar
Paul committed
189
190
191
192
        return p;
    }
};

Khalique's avatar
Khalique committed
193
194
struct test_mul
{
Paul's avatar
Paul committed
195
    migraphx::program create_program() const
Khalique's avatar
Khalique committed
196
    {
Paul's avatar
Paul committed
197
198
        migraphx::program p;
        migraphx::shape s{migraphx::shape::float_type, {3}};
Khalique's avatar
Khalique committed
199
200
        auto x = p.add_parameter("x", s);
        auto y = p.add_parameter("y", s);
Paul's avatar
Paul committed
201
        p.add_instruction(migraphx::op::mul{}, x, y);
Khalique's avatar
Khalique committed
202
203
204
205
        return p;
    }
};

206
207
208
209
210
211
212
213
214
215
216
217
struct test_sin
{
    migraphx::program create_program() const
    {
        migraphx::program p;
        migraphx::shape s{migraphx::shape::float_type, {10}};
        auto x = p.add_parameter("x", s);
        p.add_instruction(migraphx::op::sin{}, x);
        return p;
    }
};

Shucai Xiao's avatar
Shucai Xiao committed
218
219
220
221
222
223
224
225
226
227
228
229
230
231
232
233
234
235
236
237
238
239
240
241
struct test_cos
{
    migraphx::program create_program() const
    {
        migraphx::program p;
        migraphx::shape s{migraphx::shape::double_type, {8}};
        auto x = p.add_parameter("x", s);
        p.add_instruction(migraphx::op::cos{}, x);
        return p;
    }
};

struct test_tan
{
    migraphx::program create_program() const
    {
        migraphx::program p;
        migraphx::shape s{migraphx::shape::float_type, {16}};
        auto x = p.add_parameter("x", s);
        p.add_instruction(migraphx::op::tan{}, x);
        return p;
    }
};

242
243
244
245
246
247
248
249
250
251
252
253
254
255
256
257
258
259
260
261
262
263
264
265
struct test_sinh
{
    migraphx::program create_program() const
    {
        migraphx::program p;
        migraphx::shape s{migraphx::shape::double_type, {16}};
        auto x = p.add_parameter("x", s);
        p.add_instruction(migraphx::op::sinh{}, x);
        return p;
    }
};

struct test_cosh
{
    migraphx::program create_program() const
    {
        migraphx::program p;
        migraphx::shape s{migraphx::shape::double_type, {16}};
        auto x = p.add_parameter("x", s);
        p.add_instruction(migraphx::op::cosh{}, x);
        return p;
    }
};

266
267
268
269
270
271
272
273
274
275
276
struct test_tanh
{
    migraphx::program create_program() const
    {
        migraphx::program p;
        auto x = p.add_parameter("x", migraphx::shape{migraphx::shape::float_type, {4, 3, 3, 3}});
        p.add_instruction(migraphx::op::tanh{}, x);
        return p;
    }
};

277
278
279
280
281
282
283
284
285
286
287
288
289
290
291
292
293
294
295
296
297
298
299
300
301
302
303
304
305
306
307
308
309
310
311
312
struct test_asin
{
    migraphx::program create_program() const
    {
        migraphx::program p;
        migraphx::shape s{migraphx::shape::double_type, {16}};
        auto x = p.add_parameter("x", s);
        p.add_instruction(migraphx::op::asin{}, x);
        return p;
    }
};

struct test_acos
{
    migraphx::program create_program() const
    {
        migraphx::program p;
        migraphx::shape s{migraphx::shape::double_type, {16}};
        auto x = p.add_parameter("x", s);
        p.add_instruction(migraphx::op::acos{}, x);
        return p;
    }
};

struct test_atan
{
    migraphx::program create_program() const
    {
        migraphx::program p;
        migraphx::shape s{migraphx::shape::double_type, {16}};
        auto x = p.add_parameter("x", s);
        p.add_instruction(migraphx::op::atan{}, x);
        return p;
    }
};

Khalique's avatar
Khalique committed
313
314
struct test_scale
{
Paul's avatar
Paul committed
315
    migraphx::program create_program() const
Khalique's avatar
Khalique committed
316
    {
Paul's avatar
Paul committed
317
318
        migraphx::program p;
        migraphx::shape s{migraphx::shape::float_type, {3}};
Khalique's avatar
Khalique committed
319
        auto x     = p.add_parameter("x", s);
Paul's avatar
Paul committed
320
321
322
        auto y     = p.add_parameter("y", migraphx::shape::float_type);
        auto scale = p.add_instruction(migraphx::op::scalar{s}, y);
        p.add_instruction(migraphx::op::mul{}, x, scale);
Khalique's avatar
Khalique committed
323
324
325
326
        return p;
    }
};

327
328
struct test_slice
{
Paul's avatar
Paul committed
329
    migraphx::program create_program() const
330
    {
Paul's avatar
Paul committed
331
332
        migraphx::program p;
        migraphx::shape s{migraphx::shape::int32_type, {2, 2, 4}};
333
        auto x      = p.add_parameter("x", s);
Paul's avatar
Paul committed
334
335
336
        auto y      = p.add_parameter("y", {migraphx::shape::int32_type, {2, 2, 2}});
        auto slice0 = p.add_instruction(migraphx::op::slice{{2}, {0}, {2}}, x);
        p.add_instruction(migraphx::op::add{}, y, slice0);
337
338
339
340
341

        return p;
    }
};

Paul's avatar
Paul committed
342
343
struct test_triadd
{
Paul's avatar
Paul committed
344
    migraphx::program create_program() const
Paul's avatar
Paul committed
345
    {
Paul's avatar
Paul committed
346
347
        migraphx::program p;
        migraphx::shape s{migraphx::shape::float_type, {3}};
Paul's avatar
Paul committed
348
349
350
        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
351
352
        auto sum = p.add_instruction(migraphx::op::add{}, x, y);
        p.add_instruction(migraphx::op::add{}, sum, z);
Paul's avatar
Paul committed
353
354
355
356
357
358
        return p;
    }
};

struct test_triadd2
{
Paul's avatar
Paul committed
359
    migraphx::program create_program() const
Paul's avatar
Paul committed
360
    {
Paul's avatar
Paul committed
361
362
363
        migraphx::program p;
        migraphx::shape s{migraphx::shape::float_type, {2, 3}};
        migraphx::shape b{migraphx::shape::float_type, {3}};
Paul's avatar
Paul committed
364
365
366
        auto x   = p.add_parameter("x", s);
        auto y   = p.add_parameter("y", s);
        auto z   = p.add_parameter("z", b);
Paul's avatar
Paul committed
367
368
369
        auto zb  = p.add_instruction(migraphx::op::broadcast{1, s}, z);
        auto sum = p.add_instruction(migraphx::op::add{}, x, y);
        p.add_instruction(migraphx::op::add{}, sum, zb);
Paul's avatar
Paul committed
370
371
372
373
        return p;
    }
};

Paul's avatar
Paul committed
374
375
struct test_add_broadcast
{
Paul's avatar
Paul committed
376
    migraphx::program create_program() const
Paul's avatar
Paul committed
377
    {
Paul's avatar
Paul committed
378
379
380
381
382
383
        migraphx::program p;
        migraphx::shape s{migraphx::shape::float_type, {3}};
        auto x  = p.add_parameter("x", {migraphx::shape::float_type, {2, 2, 3}});
        auto y  = p.add_parameter("y", {migraphx::shape::float_type, {2, 2}});
        auto by = p.add_instruction(migraphx::op::broadcast{0, x->get_shape()}, y);
        p.add_instruction(migraphx::op::add{}, x, by);
Paul's avatar
Paul committed
384
385
386
387
        return p;
    }
};

Paul's avatar
Paul committed
388
389
struct test_add_broadcast2
{
Paul's avatar
Paul committed
390
    migraphx::program create_program() const
Paul's avatar
Paul committed
391
    {
Paul's avatar
Paul committed
392
393
394
395
396
397
        migraphx::program p;
        migraphx::shape s{migraphx::shape::float_type, {3}};
        auto x  = p.add_parameter("x", {migraphx::shape::float_type, {2, 3, 4}});
        auto y  = p.add_parameter("y", {migraphx::shape::float_type, {3}});
        auto by = p.add_instruction(migraphx::op::broadcast{1, x->get_shape()}, y);
        p.add_instruction(migraphx::op::add{}, x, by);
Paul's avatar
Paul committed
398
399
400
401
        return p;
    }
};

Paul's avatar
Latest  
Paul committed
402
403
struct test_add_broadcast3
{
Paul's avatar
Paul committed
404
    migraphx::program create_program() const
Paul's avatar
Latest  
Paul committed
405
    {
Paul's avatar
Paul committed
406
407
408
409
410
411
        migraphx::program p;
        migraphx::shape s{migraphx::shape::float_type, {3}};
        auto x  = p.add_parameter("x", {migraphx::shape::float_type, {2, 4, 5}});
        auto y  = p.add_parameter("y", {migraphx::shape::float_type, {4}});
        auto by = p.add_instruction(migraphx::op::broadcast{1, x->get_shape()}, y);
        p.add_instruction(migraphx::op::add{}, x, by);
Paul's avatar
Latest  
Paul committed
412
413
414
415
416
417
        return p;
    }
};

struct test_add_broadcast4
{
Paul's avatar
Paul committed
418
    migraphx::program create_program() const
Paul's avatar
Latest  
Paul committed
419
    {
Paul's avatar
Paul committed
420
421
422
423
424
425
        migraphx::program p;
        migraphx::shape s{migraphx::shape::float_type, {3}};
        auto x  = p.add_parameter("x", {migraphx::shape::float_type, {2, 3, 5}});
        auto y  = p.add_parameter("y", {migraphx::shape::float_type, {3}});
        auto by = p.add_instruction(migraphx::op::broadcast{1, x->get_shape()}, y);
        p.add_instruction(migraphx::op::add{}, x, by);
Paul's avatar
Latest  
Paul committed
426
427
428
429
        return p;
    }
};

Paul's avatar
Paul committed
430
431
struct test_add_broadcast5
{
Paul's avatar
Paul committed
432
    migraphx::program create_program() const
Paul's avatar
Paul committed
433
    {
Paul's avatar
Paul committed
434
435
436
437
438
439
        migraphx::program p;
        migraphx::shape s{migraphx::shape::float_type, {3}};
        auto x  = p.add_parameter("x", {migraphx::shape::float_type, {2, 4, 8}});
        auto y  = p.add_parameter("y", {migraphx::shape::float_type, {4}});
        auto by = p.add_instruction(migraphx::op::broadcast{1, x->get_shape()}, y);
        p.add_instruction(migraphx::op::add{}, x, by);
Paul's avatar
Paul committed
440
441
442
443
        return p;
    }
};

Paul's avatar
Paul committed
444
445
struct test_triadd_broadcast
{
Paul's avatar
Paul committed
446
    migraphx::program create_program() const
Paul's avatar
Paul committed
447
    {
Paul's avatar
Paul committed
448
449
450
451
452
453
454
455
        migraphx::program p;
        migraphx::shape s{migraphx::shape::float_type, {3}};
        auto x   = p.add_parameter("x", {migraphx::shape::float_type, {2, 2, 3}});
        auto y   = p.add_parameter("y", {migraphx::shape::float_type, {2, 2}});
        auto z   = p.add_parameter("z", {migraphx::shape::float_type, {2, 2, 3}});
        auto by  = p.add_instruction(migraphx::op::broadcast{0, x->get_shape()}, y);
        auto sum = p.add_instruction(migraphx::op::add{}, x, by);
        p.add_instruction(migraphx::op::add{}, sum, z);
Paul's avatar
Paul committed
456
457
458
459
        return p;
    }
};

Paul's avatar
Paul committed
460
461
struct test_softmax
{
Paul's avatar
Paul committed
462
    migraphx::program create_program() const
Paul's avatar
Paul committed
463
    {
Paul's avatar
Paul committed
464
465
466
        migraphx::program p;
        auto x = p.add_parameter("x", migraphx::shape{migraphx::shape::float_type, {5, 3, 4, 2}});
        p.add_instruction(migraphx::op::softmax{}, x);
Paul's avatar
Paul committed
467
468
469
470
471
472
        return p;
    }
};

struct test_softmax2
{
Paul's avatar
Paul committed
473
    migraphx::program create_program() const
Paul's avatar
Paul committed
474
    {
Paul's avatar
Paul committed
475
        migraphx::program p;
Paul's avatar
Paul committed
476
477
        auto x =
            p.add_parameter("x", migraphx::shape{migraphx::shape::float_type, {1, 1000, 1, 1}});
Paul's avatar
Paul committed
478
        p.add_instruction(migraphx::op::softmax{}, x);
Paul's avatar
Paul committed
479
480
481
482
        return p;
    }
};

Paul's avatar
Paul committed
483
484
struct test_conv
{
Paul's avatar
Paul committed
485
    migraphx::program create_program() const
Paul's avatar
Paul committed
486
    {
Paul's avatar
Paul committed
487
        migraphx::program p;
Paul's avatar
Paul committed
488
489
        auto input =
            p.add_parameter("x", migraphx::shape{migraphx::shape::float_type, {4, 3, 3, 3}});
Paul's avatar
Paul committed
490
        auto weights =
Paul's avatar
Paul committed
491
492
            p.add_parameter("w", migraphx::shape{migraphx::shape::float_type, {4, 3, 3, 3}});
        p.add_instruction(migraphx::op::convolution{}, input, weights);
Paul's avatar
Paul committed
493
494
495
496
        return p;
    }
};

Paul's avatar
Paul committed
497
498
struct test_conv2
{
Paul's avatar
Paul committed
499
    migraphx::program create_program() const
Paul's avatar
Paul committed
500
    {
Paul's avatar
Paul committed
501
        migraphx::program p;
Paul's avatar
Paul committed
502
        auto input =
Paul's avatar
Paul committed
503
            p.add_parameter("x", migraphx::shape{migraphx::shape::float_type, {1, 512, 28, 28}});
Paul's avatar
Paul committed
504
        auto weights =
Paul's avatar
Paul committed
505
506
            p.add_parameter("w", migraphx::shape{migraphx::shape::float_type, {256, 512, 1, 1}});
        p.add_instruction(migraphx::op::convolution{{0, 0}, {1, 1}, {1, 1}}, input, weights);
Paul's avatar
Paul committed
507
508
509
510
        return p;
    }
};

Paul's avatar
Paul committed
511
struct test_conv_relu
Paul's avatar
Paul committed
512
{
Paul's avatar
Paul committed
513
    migraphx::program create_program() const
Paul's avatar
Paul committed
514
    {
Paul's avatar
Paul committed
515
        migraphx::program p;
Paul's avatar
Paul committed
516
517
        auto input =
            p.add_parameter("x", migraphx::shape{migraphx::shape::float_type, {4, 3, 3, 3}});
Paul's avatar
Paul committed
518
        auto weights =
Paul's avatar
Paul committed
519
520
521
            p.add_parameter("w", migraphx::shape{migraphx::shape::float_type, {4, 3, 3, 3}});
        auto conv = p.add_instruction(migraphx::op::convolution{}, input, weights);
        p.add_instruction(migraphx::op::relu{}, conv);
Paul's avatar
Paul committed
522
523
524
525
        return p;
    }
};

Paul's avatar
Paul committed
526
527
struct test_conv_relu_half
{
Paul's avatar
Paul committed
528
    migraphx::program create_program() const
Paul's avatar
Paul committed
529
    {
Paul's avatar
Paul committed
530
        migraphx::program p;
Paul's avatar
Paul committed
531
532
        auto input =
            p.add_parameter("x", migraphx::shape{migraphx::shape::half_type, {4, 3, 3, 3}});
Paul's avatar
Paul committed
533
        auto weights =
Paul's avatar
Paul committed
534
535
536
            p.add_parameter("w", migraphx::shape{migraphx::shape::half_type, {4, 3, 3, 3}});
        auto conv = p.add_instruction(migraphx::op::convolution{}, input, weights);
        p.add_instruction(migraphx::op::relu{}, conv);
Paul's avatar
Paul committed
537
538
539
540
        return p;
    }
};

Paul's avatar
Paul committed
541
542
struct test_add_relu
{
Paul's avatar
Paul committed
543
    migraphx::program create_program() const
Paul's avatar
Paul committed
544
    {
Paul's avatar
Paul committed
545
546
547
548
549
        migraphx::program p;
        auto x   = p.add_parameter("x", migraphx::shape{migraphx::shape::float_type, {4, 3, 3, 3}});
        auto y   = p.add_parameter("y", migraphx::shape{migraphx::shape::float_type, {4, 3, 3, 3}});
        auto add = p.add_instruction(migraphx::op::add{}, x, y);
        p.add_instruction(migraphx::op::relu{}, add);
Paul's avatar
Paul committed
550
551
552
553
        return p;
    }
};

Khalique's avatar
Khalique committed
554
555
556
557
558
struct test_sigmoid
{
    migraphx::program create_program() const
    {
        migraphx::program p;
Khalique's avatar
Khalique committed
559
        auto x = p.add_parameter("x", migraphx::shape{migraphx::shape::float_type, {4, 3, 3, 3}});
Khalique's avatar
Khalique committed
560
561
562
563
564
565
566
567
568
569
        p.add_instruction(migraphx::op::sigmoid{}, x);
        return p;
    }
};

struct test_abs
{
    migraphx::program create_program() const
    {
        migraphx::program p;
Khalique's avatar
Khalique committed
570
        auto x = p.add_parameter("x", migraphx::shape{migraphx::shape::float_type, {4, 3, 3, 3}});
Khalique's avatar
Khalique committed
571
572
573
574
575
        p.add_instruction(migraphx::op::abs{}, x);
        return p;
    }
};

576
577
struct test_leaky_relu
{
Paul's avatar
Paul committed
578
    migraphx::program create_program() const
579
    {
Paul's avatar
Paul committed
580
581
582
        migraphx::program p;
        auto x = p.add_parameter("x", migraphx::shape{migraphx::shape::float_type, {4, 3, 3, 3}});
        p.add_instruction(migraphx::op::leaky_relu{0.01}, x);
583
584
585
586
        return p;
    }
};

Khalique's avatar
Khalique committed
587
588
589
590
591
592
593
594
595
596
597
struct test_elu
{
    migraphx::program create_program() const
    {
        migraphx::program p;
        auto x = p.add_parameter("x", migraphx::shape{migraphx::shape::float_type, {4, 3, 3, 3}});
        p.add_instruction(migraphx::op::leaky_relu{1.0}, x);
        return p;
    }
};

Paul's avatar
Paul committed
598
599
struct test_conv_pooling
{
Paul's avatar
Paul committed
600
    migraphx::program create_program() const
Paul's avatar
Paul committed
601
    {
Paul's avatar
Paul committed
602
        migraphx::program p;
Paul's avatar
Paul committed
603
        auto input =
Paul's avatar
Paul committed
604
            p.add_parameter("x", migraphx::shape{migraphx::shape::float_type, {4, 3, 32, 32}});
Paul's avatar
Paul committed
605
        auto weights =
Paul's avatar
Paul committed
606
607
608
609
            p.add_parameter("w", migraphx::shape{migraphx::shape::float_type, {4, 3, 3, 3}});
        auto conv    = p.add_instruction(migraphx::op::convolution{}, input, weights);
        auto pooling = p.add_instruction(migraphx::op::pooling{"max"}, conv);
        p.add_instruction(migraphx::op::relu{}, pooling);
Paul's avatar
Paul committed
610
611
612
613
        return p;
    }
};

614
615
struct test_global_avg_pooling
{
Paul's avatar
Paul committed
616
    migraphx::program create_program() const
617
    {
Paul's avatar
Paul committed
618
        migraphx::program p;
619
        auto input =
Paul's avatar
Paul committed
620
621
            p.add_parameter("x", migraphx::shape{migraphx::shape::float_type, {1, 3, 16, 16}});
        auto op    = migraphx::op::pooling{"average"};
622
        auto lens  = input->get_shape().lens();
Khalique's avatar
Khalique committed
623
        op.lengths = {lens[2], lens[3]};
624
625
626
627
628
629
630
        p.add_instruction(op, input);
        return p;
    }
};

struct test_global_max_pooling
{
Paul's avatar
Paul committed
631
    migraphx::program create_program() const
632
    {
Paul's avatar
Paul committed
633
        migraphx::program p;
634
        auto input =
Paul's avatar
Paul committed
635
636
            p.add_parameter("x", migraphx::shape{migraphx::shape::float_type, {1, 3, 16, 16}});
        auto op    = migraphx::op::pooling{"max"};
637
        auto lens  = input->get_shape().lens();
Khalique's avatar
Khalique committed
638
        op.lengths = {lens[2], lens[3]};
639
640
641
642
643
        p.add_instruction(op, input);
        return p;
    }
};

Paul's avatar
Paul committed
644
645
struct test_gemm
{
Paul's avatar
Paul committed
646
    migraphx::program create_program() const
Paul's avatar
Paul committed
647
    {
Paul's avatar
Paul committed
648
649
650
651
        migraphx::program p;
        auto a = p.add_parameter("a", migraphx::shape{migraphx::shape::float_type, {4, 5}});
        auto b = p.add_parameter("b", migraphx::shape{migraphx::shape::float_type, {5, 3}});
        p.add_instruction(migraphx::op::dot{}, a, b);
Paul's avatar
Paul committed
652
653
654
655
        return p;
    }
};

Paul's avatar
Paul committed
656
657
struct test_gemm_half
{
Paul's avatar
Paul committed
658
    migraphx::program create_program() const
Paul's avatar
Paul committed
659
    {
Paul's avatar
Paul committed
660
661
662
663
        migraphx::program p;
        auto a = p.add_parameter("a", migraphx::shape{migraphx::shape::half_type, {4, 5}});
        auto b = p.add_parameter("b", migraphx::shape{migraphx::shape::half_type, {5, 3}});
        p.add_instruction(migraphx::op::dot{}, a, b);
Paul's avatar
Paul committed
664
665
666
667
        return p;
    }
};

Paul's avatar
Paul committed
668
669
struct test_gemm_ld
{
Paul's avatar
Paul committed
670
    migraphx::program create_program() const
Paul's avatar
Paul committed
671
    {
Paul's avatar
Paul committed
672
        migraphx::program p;
Paul's avatar
Paul committed
673
674
675
676
        auto a =
            p.add_parameter("a", migraphx::shape{migraphx::shape::float_type, {4, 5}, {10, 1}});
        auto b =
            p.add_parameter("b", migraphx::shape{migraphx::shape::float_type, {5, 3}, {20, 1}});
Paul's avatar
Paul committed
677
        p.add_instruction(migraphx::op::dot{}, a, b);
Paul's avatar
Paul committed
678
679
680
681
        return p;
    }
};

682
683
struct test_gemm_transposeb
{
Paul's avatar
Paul committed
684
    migraphx::program create_program() const
685
    {
Paul's avatar
Paul committed
686
687
688
689
690
        migraphx::program p;
        auto a  = p.add_parameter("a", migraphx::shape{migraphx::shape::float_type, {4, 5}});
        auto b  = p.add_parameter("b", migraphx::shape{migraphx::shape::float_type, {3, 5}});
        auto bt = p.add_instruction(migraphx::op::transpose{{1, 0}}, b);
        p.add_instruction(migraphx::op::dot{}, a, bt);
691
692
693
694
695
696
        return p;
    }
};

struct test_gemm_transposea
{
Paul's avatar
Paul committed
697
    migraphx::program create_program() const
698
    {
Paul's avatar
Paul committed
699
700
701
702
703
        migraphx::program p;
        auto a  = p.add_parameter("a", migraphx::shape{migraphx::shape::float_type, {5, 4}});
        auto b  = p.add_parameter("b", migraphx::shape{migraphx::shape::float_type, {5, 3}});
        auto at = p.add_instruction(migraphx::op::transpose{{1, 0}}, a);
        p.add_instruction(migraphx::op::dot{}, at, b);
704
705
706
707
708
709
        return p;
    }
};

struct test_gemm_transposeab
{
Paul's avatar
Paul committed
710
    migraphx::program create_program() const
711
    {
Paul's avatar
Paul committed
712
713
714
715
716
717
        migraphx::program p;
        auto a  = p.add_parameter("a", migraphx::shape{migraphx::shape::float_type, {5, 4}});
        auto b  = p.add_parameter("b", migraphx::shape{migraphx::shape::float_type, {3, 5}});
        auto at = p.add_instruction(migraphx::op::transpose{{1, 0}}, a);
        auto bt = p.add_instruction(migraphx::op::transpose{{1, 0}}, b);
        p.add_instruction(migraphx::op::dot{}, at, bt);
718
719
720
721
        return p;
    }
};

722
723
struct test_contiguous
{
Paul's avatar
Paul committed
724
    migraphx::program create_program() const
725
    {
Paul's avatar
Paul committed
726
727
        migraphx::program p;
        migraphx::shape s{migraphx::shape::float_type, {4, 4, 4, 3}, {48, 4, 1, 16}};
728
        auto x = p.add_parameter("x", s);
Paul's avatar
Paul committed
729
        p.add_instruction(migraphx::op::contiguous{}, x);
Paul's avatar
Paul committed
730
        EXPECT(p.get_shape().standard());
731
732
733
734
        return p;
    }
};

735
struct test_transpose
736
{
Paul's avatar
Paul committed
737
    migraphx::program create_program() const
738
    {
Paul's avatar
Paul committed
739
740
        migraphx::program p;
        migraphx::shape s{migraphx::shape::float_type, {4, 3, 4, 4}};
741
742
        auto x                    = p.add_parameter("x", s);
        std::vector<int64_t> perm = {0, 2, 3, 1};
Paul's avatar
Paul committed
743
744
        auto l                    = p.add_instruction(migraphx::op::transpose{perm}, x);
        p.add_instruction(migraphx::op::contiguous{}, l);
745
746
747
        return p;
    }
};
748

Paul's avatar
Paul committed
749
750
751
752
753
754
755
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;

Paul's avatar
Paul committed
756
    migraphx::program create_program() const
Paul's avatar
Paul committed
757
    {
Paul's avatar
Paul committed
758
        migraphx::program p;
Paul's avatar
Paul committed
759

Paul's avatar
Paul committed
760
761
        migraphx::shape s{migraphx::shape::float_type, {batches, channels, height, width}};
        migraphx::shape vars{migraphx::shape::float_type, {channels}};
Paul's avatar
Paul committed
762
        auto x        = p.add_parameter("x", s);
Paul's avatar
Paul committed
763
764
765
766
767
        auto scale    = p.add_literal(migraphx::abs(migraphx::generate_literal(vars, 1)));
        auto bias     = p.add_literal(migraphx::abs(migraphx::generate_literal(vars, 2)));
        auto mean     = p.add_literal(migraphx::abs(migraphx::generate_literal(vars, 3)));
        auto variance = p.add_literal(migraphx::abs(migraphx::generate_literal(vars, 4)));
        p.add_instruction(migraphx::op::batch_norm_inference{}, x, scale, bias, mean, variance);
Paul's avatar
Paul committed
768
769
770
771
        return p;
    }
};

wsttiger's avatar
wsttiger committed
772
773
774
775
776
777
778
struct test_batchnorm_inference
{
    const size_t width    = 3;
    const size_t height   = 3;
    const size_t channels = 3;
    const size_t batches  = 4;

Paul's avatar
Paul committed
779
    migraphx::program create_program() const
wsttiger's avatar
wsttiger committed
780
    {
Paul's avatar
Paul committed
781
        migraphx::program p;
wsttiger's avatar
wsttiger committed
782

Paul's avatar
Paul committed
783
784
        migraphx::shape s{migraphx::shape::float_type, {batches, channels, height, width}};
        migraphx::shape vars{migraphx::shape::float_type, {channels}};
wsttiger's avatar
wsttiger committed
785
        auto x        = p.add_parameter("x", s);
Paul's avatar
Paul committed
786
787
788
789
790
        auto scale    = p.add_literal(migraphx::abs(migraphx::generate_literal(vars, 1)));
        auto bias     = p.add_literal(migraphx::abs(migraphx::generate_literal(vars, 2)));
        auto mean     = p.add_literal(migraphx::abs(migraphx::generate_literal(vars, 3)));
        auto variance = p.add_literal(migraphx::abs(migraphx::generate_literal(vars, 4)));
        p.add_instruction(migraphx::op::batch_norm_inference{}, x, scale, bias, mean, variance);
wsttiger's avatar
wsttiger committed
791
792
793
794
        return p;
    }
};

Paul's avatar
Paul committed
795
796
struct test_conv_bn
{
Paul's avatar
Paul committed
797
    migraphx::program create_program() const
Paul's avatar
Paul committed
798
    {
Paul's avatar
Paul committed
799
        migraphx::program p;
Paul's avatar
Paul committed
800

Paul's avatar
Paul committed
801
802
803
        migraphx::shape xs{migraphx::shape::float_type, {1, 3, 224, 224}};
        migraphx::shape ws{migraphx::shape::float_type, {64, 3, 7, 7}};
        migraphx::shape vars{migraphx::shape::float_type, {64}};
Paul's avatar
Paul committed
804
805
        auto x        = p.add_parameter("x", xs);
        auto w        = p.add_parameter("w", ws);
Paul's avatar
Paul committed
806
807
808
809
810
811
        auto conv     = p.add_instruction(migraphx::op::convolution{{3, 3}, {2, 2}, {1, 1}}, x, w);
        auto scale    = p.add_literal(migraphx::abs(migraphx::generate_literal(vars, 1)));
        auto bias     = p.add_literal(migraphx::abs(migraphx::generate_literal(vars, 2)));
        auto mean     = p.add_literal(migraphx::abs(migraphx::generate_literal(vars, 3)));
        auto variance = p.add_literal(migraphx::abs(migraphx::generate_literal(vars, 4)));
        p.add_instruction(migraphx::op::batch_norm_inference{}, conv, scale, bias, mean, variance);
Paul's avatar
Paul committed
812
813
814
815
        return p;
    }
};

Paul's avatar
Paul committed
816
817
struct test_conv_bn_relu_pooling
{
Paul's avatar
Paul committed
818
    migraphx::program create_program() const
Paul's avatar
Paul committed
819
    {
Paul's avatar
Paul committed
820
        migraphx::program p;
Paul's avatar
Paul committed
821

Paul's avatar
Paul committed
822
823
824
        migraphx::shape xs{migraphx::shape::float_type, {1, 3, 224, 224}};
        migraphx::shape ws{migraphx::shape::float_type, {64, 3, 7, 7}};
        migraphx::shape vars{migraphx::shape::float_type, {64}};
Paul's avatar
Paul committed
825
826
        auto x        = p.add_parameter("x", xs);
        auto w        = p.add_parameter("w", ws);
Paul's avatar
Paul committed
827
828
829
830
831
        auto conv     = p.add_instruction(migraphx::op::convolution{{3, 3}, {2, 2}, {1, 1}}, x, w);
        auto scale    = p.add_literal(migraphx::abs(migraphx::generate_literal(vars, 1)));
        auto bias     = p.add_literal(migraphx::abs(migraphx::generate_literal(vars, 2)));
        auto mean     = p.add_literal(migraphx::abs(migraphx::generate_literal(vars, 3)));
        auto variance = p.add_literal(migraphx::abs(migraphx::generate_literal(vars, 4)));
wsttiger's avatar
wsttiger committed
832
        auto bn       = p.add_instruction(
Paul's avatar
Paul committed
833
834
835
            migraphx::op::batch_norm_inference{}, conv, scale, bias, mean, variance);
        auto relu = p.add_instruction(migraphx::op::relu{}, bn);
        p.add_instruction(migraphx::op::pooling{"average", {1, 1}, {2, 2}, {3, 3}}, relu);
Paul's avatar
Paul committed
836
837
838
839
        return p;
    }
};

840
841
struct test_concat
{
Paul's avatar
Paul committed
842
    migraphx::program create_program() const
843
    {
Paul's avatar
Paul committed
844
        migraphx::program p;
wsttiger's avatar
wsttiger committed
845
        std::size_t axis = 1;
Paul's avatar
Paul committed
846
847
848
        migraphx::shape s0{migraphx::shape::int32_type, {2, 2}};
        migraphx::shape s1{migraphx::shape::int32_type, {2, 3}};
        migraphx::shape s2{migraphx::shape::int32_type, {2, 1}};
849
850
851
        auto l0 = p.add_parameter("x", s0);
        auto l1 = p.add_parameter("y", s1);
        auto l2 = p.add_parameter("z", s2);
Paul's avatar
Paul committed
852
        p.add_instruction(migraphx::op::concat{axis}, l0, l1, l2);
853
854
855
856
857
858
        return p;
    }
};

struct test_concat2
{
Paul's avatar
Paul committed
859
    migraphx::program create_program() const
860
    {
Paul's avatar
Paul committed
861
        migraphx::program p;
wsttiger's avatar
wsttiger committed
862
        std::size_t axis = 0;
Paul's avatar
Paul committed
863
864
865
        migraphx::shape s0{migraphx::shape::int32_type, {2, 2}};
        migraphx::shape s1{migraphx::shape::int32_type, {3, 2}};
        migraphx::shape s2{migraphx::shape::int32_type, {1, 2}};
866
867
868
        auto l0 = p.add_parameter("x", s0);
        auto l1 = p.add_parameter("y", s1);
        auto l2 = p.add_parameter("z", s2);
Paul's avatar
Paul committed
869
        p.add_instruction(migraphx::op::concat{axis}, l0, l1, l2);
870
871
872
873
        return p;
    }
};

wsttiger's avatar
wsttiger committed
874
875
struct test_concat_relu
{
Paul's avatar
Paul committed
876
    migraphx::program create_program() const
wsttiger's avatar
wsttiger committed
877
    {
Paul's avatar
Paul committed
878
        migraphx::program p;
wsttiger's avatar
wsttiger committed
879
        std::size_t axis = 0;
Paul's avatar
Paul committed
880
881
882
        migraphx::shape s0{migraphx::shape::float_type, {2, 2}};
        migraphx::shape s1{migraphx::shape::float_type, {3, 2}};
        migraphx::shape s2{migraphx::shape::float_type, {1, 2}};
wsttiger's avatar
wsttiger committed
883
884
885
        auto l0 = p.add_parameter("x", s0);
        auto l1 = p.add_parameter("y", s1);
        auto l2 = p.add_parameter("z", s2);
Paul's avatar
Paul committed
886
887
888
889
890
        auto r0 = p.add_instruction(migraphx::op::relu{}, l0);
        auto r1 = p.add_instruction(migraphx::op::relu{}, l1);
        auto r2 = p.add_instruction(migraphx::op::relu{}, l2);
        auto c0 = p.add_instruction(migraphx::op::concat{axis}, r0, r1, r2);
        p.add_instruction(migraphx::op::relu{}, c0);
wsttiger's avatar
wsttiger committed
891
892
893
894
895
896
        return p;
    }
};

void manual_identity()
{
Paul's avatar
Paul committed
897
    migraphx::program p;
wsttiger's avatar
wsttiger committed
898
    std::vector<float> data0 = {0, 1, 2, 3};
Paul's avatar
Paul committed
899
900
901
902
903
    migraphx::shape s0{migraphx::shape::float_type, {2, 2}};
    auto l0 = p.add_literal(migraphx::literal{s0, data0});
    p.add_instruction(migraphx::op::identity{}, l0);
    p.compile(migraphx::gpu::target{});
    migraphx::program::parameter_map m;
wsttiger's avatar
wsttiger committed
904
905
    for(auto&& x : p.get_parameter_shapes())
    {
Paul's avatar
Paul committed
906
        m[x.first] = migraphx::gpu::to_gpu(migraphx::generate_argument(x.second));
wsttiger's avatar
wsttiger committed
907
    }
Paul's avatar
Paul committed
908
    auto result = migraphx::gpu::from_gpu(p.eval(m));
wsttiger's avatar
wsttiger committed
909
910
911
912
913
    std::cout << result << std::endl;
}

void manual_test_concat_relu()
{
Paul's avatar
Paul committed
914
    migraphx::program p;
wsttiger's avatar
wsttiger committed
915
    std::size_t axis         = 0;
wsttiger's avatar
wsttiger committed
916
917
918
    std::vector<float> data0 = {0, 1, 2, 3};
    std::vector<float> data1 = {4, 5, 6, 7, 8, 9};
    std::vector<float> data2 = {10, 11};
Paul's avatar
Paul committed
919
920
921
922
923
924
925
926
927
928
929
930
931
932
    migraphx::shape s0{migraphx::shape::float_type, {2, 2}};
    migraphx::shape s1{migraphx::shape::float_type, {3, 2}};
    migraphx::shape s2{migraphx::shape::float_type, {1, 2}};
    auto l0 = p.add_literal(migraphx::literal{s0, data0});
    auto l1 = p.add_literal(migraphx::literal{s1, data1});
    auto l2 = p.add_literal(migraphx::literal{s2, data2});
    auto r0 = p.add_instruction(migraphx::op::relu{}, l0);
    auto r1 = p.add_instruction(migraphx::op::relu{}, l1);
    auto r2 = p.add_instruction(migraphx::op::relu{}, l2);
    auto c0 = p.add_instruction(migraphx::op::concat{axis}, r0, r1, r2);
    p.add_instruction(migraphx::op::relu{}, c0);

    p.compile(migraphx::gpu::target{});
    migraphx::program::parameter_map m;
wsttiger's avatar
wsttiger committed
933
934
    for(auto&& x : p.get_parameter_shapes())
    {
Paul's avatar
Paul committed
935
        m[x.first] = migraphx::gpu::to_gpu(migraphx::generate_argument(x.second));
wsttiger's avatar
wsttiger committed
936
    }
Paul's avatar
Paul committed
937
    auto result = migraphx::gpu::from_gpu(p.eval(m));
wsttiger's avatar
wsttiger committed
938
939
940
    std::cout << result << std::endl;
}

Paul's avatar
Paul committed
941
942
struct test_conv_bn_relu_pooling2
{
Paul's avatar
Paul committed
943
944
    static migraphx::instruction_ref
    add_bn(migraphx::program& p, migraphx::instruction_ref x, std::size_t channels)
Paul's avatar
Paul committed
945
    {
Paul's avatar
Paul committed
946
        migraphx::shape vars{migraphx::shape::float_type, {channels}};
Paul's avatar
Paul committed
947
948
949
950
951
        auto scale = p.add_literal(migraphx::abs(migraphx::generate_literal(vars, 1 + channels)));
        auto bias  = p.add_literal(migraphx::abs(migraphx::generate_literal(vars, 2 + channels)));
        auto mean  = p.add_literal(migraphx::abs(migraphx::generate_literal(vars, 3 + channels)));
        auto variance =
            p.add_literal(migraphx::abs(migraphx::generate_literal(vars, 4 + channels)));
wsttiger's avatar
wsttiger committed
952
        return p.add_instruction(
Paul's avatar
Paul committed
953
            migraphx::op::batch_norm_inference{}, x, scale, bias, mean, variance);
Paul's avatar
Paul committed
954
    }
Paul's avatar
Paul committed
955
    migraphx::program create_program() const
Paul's avatar
Paul committed
956
    {
Paul's avatar
Paul committed
957
        migraphx::program p;
Paul's avatar
Paul committed
958

Paul's avatar
Paul committed
959
960
961
962
        migraphx::shape xs1{migraphx::shape::float_type, {1, 512, 7, 7}};
        migraphx::shape xs2{migraphx::shape::float_type, {1, 1024, 14, 14}};
        migraphx::shape ws1{migraphx::shape::float_type, {2048, 512, 1, 1}};
        migraphx::shape ws2{migraphx::shape::float_type, {2048, 1024, 1, 1}};
Paul's avatar
Paul committed
963
964
        auto x1    = p.add_parameter("x1", xs1);
        auto w1    = p.add_parameter("w1", ws1);
Paul's avatar
Paul committed
965
        auto conv1 = p.add_instruction(migraphx::op::convolution{{0, 0}, {1, 1}, {1, 1}}, x1, w1);
Paul's avatar
Paul committed
966
967
968
        auto bn1   = add_bn(p, conv1, 2048);
        auto x2    = p.add_parameter("x2", xs2);
        auto w2    = p.add_parameter("w2", ws2);
Paul's avatar
Paul committed
969
        auto conv2 = p.add_instruction(migraphx::op::convolution{{0, 0}, {2, 2}, {1, 1}}, x2, w2);
Paul's avatar
Paul committed
970
        auto bn2   = add_bn(p, conv2, 2048);
Paul's avatar
Paul committed
971
972
973
        auto add   = p.add_instruction(migraphx::op::add{}, bn1, bn2);
        auto relu  = p.add_instruction(migraphx::op::relu{}, add);
        p.add_instruction(migraphx::op::pooling{"average", {1, 1}, {2, 2}, {3, 3}}, relu);
Paul's avatar
Paul committed
974
975
976
977
        return p;
    }
};

Paul's avatar
Paul committed
978
979
int main()
{
Khalique's avatar
Khalique committed
980
    verify_program<test_abs>();
981
982
    verify_program<test_concat>();
    verify_program<test_concat2>();
wsttiger's avatar
wsttiger committed
983
    verify_program<test_concat_relu>();
Paul's avatar
Paul committed
984
    verify_program<test_add>();
Paul's avatar
Paul committed
985
    verify_program<test_add_half>();
Khalique's avatar
Khalique committed
986
    verify_program<test_mul>();
987
    verify_program<test_sin>();
988
989
    verify_program<test_cos>();
    verify_program<test_tan>();
990
991
    verify_program<test_sinh>();
    verify_program<test_cosh>();
992
    verify_program<test_tanh>();
993
994
995
    verify_program<test_asin>();
    verify_program<test_acos>();
    verify_program<test_atan>();
Khalique's avatar
Khalique committed
996
    verify_program<test_scale>();
Paul's avatar
Paul committed
997
998
    verify_program<test_triadd>();
    verify_program<test_triadd2>();
Paul's avatar
Paul committed
999
    verify_program<test_add_broadcast>();
Paul's avatar
Paul committed
1000
    verify_program<test_add_broadcast2>();
Paul's avatar
Latest  
Paul committed
1001
1002
    verify_program<test_add_broadcast3>();
    verify_program<test_add_broadcast4>();
Paul's avatar
Paul committed
1003
    verify_program<test_add_broadcast5>();
Paul's avatar
Paul committed
1004
    verify_program<test_triadd_broadcast>();
Paul's avatar
Paul committed
1005
    verify_program<test_softmax>();
Paul's avatar
Paul committed
1006
    verify_program<test_softmax2>();
Paul's avatar
Paul committed
1007
    verify_program<test_conv>();
Paul's avatar
Paul committed
1008
    verify_program<test_conv2>();
Paul's avatar
Paul committed
1009
    verify_program<test_conv_relu>();
Paul's avatar
Paul committed
1010
    verify_program<test_conv_relu_half>();
Paul's avatar
Paul committed
1011
    verify_program<test_add_relu>();
1012
    verify_program<test_leaky_relu>();
Khalique's avatar
Khalique committed
1013
1014
    verify_program<test_sigmoid>();
    verify_program<test_elu>();
Paul's avatar
Paul committed
1015
    verify_program<test_conv_pooling>();
1016
1017
    verify_program<test_global_avg_pooling>();
    verify_program<test_global_max_pooling>();
Paul's avatar
Paul committed
1018
    verify_program<test_gemm>();
Paul's avatar
Paul committed
1019
    verify_program<test_gemm_half>();
1020
    // verify_program<test_gemm_ld>();
1021
1022
1023
    verify_program<test_gemm_transposeb>();
    verify_program<test_gemm_transposea>();
    verify_program<test_gemm_transposeab>();
1024
1025
    verify_program<test_contiguous>();
    verify_program<test_transpose>();
1026
    verify_program<test_batchnorm_inference>();
Paul's avatar
Paul committed
1027
    verify_program<test_batchnorm_inference_2>();
Paul's avatar
Paul committed
1028
    verify_program<test_conv_bn>();
Paul's avatar
Paul committed
1029
    verify_program<test_conv_bn_relu_pooling>();
Paul's avatar
Paul committed
1030
    verify_program<test_conv_bn_relu_pooling2>();
1031
    verify_program<test_slice>();
Paul's avatar
Paul committed
1032
}