miopen.cpp 54.5 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
    // Program should have an output parameter
Paul's avatar
Paul committed
126
    EXPECT(bool{m.find("output") != m.end()});
Paul's avatar
Paul committed
127
128
129
130
131
132
    // Ensure the program doesn't modify the context in a dry run
    auto ctx = p.get_context();
    assert(&ctx != &p.get_context());
    EXPECT(is_shared(ctx, p.get_context()));
    p.dry_run(m);
    EXPECT(is_shared(ctx, p.get_context()));
Paul's avatar
Paul committed
133
    return migraphx::gpu::from_gpu(p.eval(m));
Paul's avatar
Paul committed
134
135
}

Paul's avatar
Paul committed
136
137
138
template <class V>
void verify_program()
{
Paul's avatar
Paul committed
139
140
141
142
    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
143
144
    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
145
    auto cpu_arg   = cpu_arg_f.get();
Paul's avatar
Paul committed
146
    bool passed    = verify_args(migraphx::get_type_name<V>(), cpu_arg, gpu_arg);
Paul's avatar
Paul committed
147
148
149
150
151
152
153
154
155
    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
156
    std::set_terminate(nullptr);
Paul's avatar
Paul committed
157
158
}

Paul's avatar
Paul committed
159
160
struct test_literals
{
Paul's avatar
Paul committed
161
    migraphx::program create_program() const
Paul's avatar
Paul committed
162
    {
Paul's avatar
Paul committed
163
        migraphx::program p;
Paul's avatar
Paul committed
164
        auto input = p.add_literal(
Paul's avatar
Paul committed
165
            generate_literal(migraphx::shape{migraphx::shape::float_type, {4, 3, 3, 3}}));
Paul's avatar
Paul committed
166
        auto weights = p.add_literal(
Paul's avatar
Paul committed
167
168
169
            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
170
171
172
173
        return p;
    }
};

Paul's avatar
Paul committed
174
175
struct test_add
{
Paul's avatar
Paul committed
176
    migraphx::program create_program() const
Paul's avatar
Paul committed
177
    {
Paul's avatar
Paul committed
178
179
        migraphx::program p;
        migraphx::shape s{migraphx::shape::float_type, {3}};
Paul's avatar
Paul committed
180
181
        auto x = p.add_parameter("x", s);
        auto y = p.add_parameter("y", s);
Paul's avatar
Paul committed
182
        p.add_instruction(migraphx::op::add{}, x, y);
Paul's avatar
Paul committed
183
184
185
186
        return p;
    }
};

Paul's avatar
Paul committed
187
188
struct test_add_half
{
Paul's avatar
Paul committed
189
    migraphx::program create_program() const
Paul's avatar
Paul committed
190
    {
Paul's avatar
Paul committed
191
192
        migraphx::program p;
        migraphx::shape s{migraphx::shape::half_type, {3}};
Paul's avatar
Paul committed
193
194
        auto x = p.add_parameter("x", s);
        auto y = p.add_parameter("y", s);
Paul's avatar
Paul committed
195
        p.add_instruction(migraphx::op::add{}, x, y);
Paul's avatar
Paul committed
196
197
198
199
        return p;
    }
};

Khalique's avatar
Khalique committed
200
201
struct test_mul
{
Paul's avatar
Paul committed
202
    migraphx::program create_program() const
Khalique's avatar
Khalique committed
203
    {
Paul's avatar
Paul committed
204
205
        migraphx::program p;
        migraphx::shape s{migraphx::shape::float_type, {3}};
Khalique's avatar
Khalique committed
206
207
        auto x = p.add_parameter("x", s);
        auto y = p.add_parameter("y", s);
Paul's avatar
Paul committed
208
        p.add_instruction(migraphx::op::mul{}, x, y);
Khalique's avatar
Khalique committed
209
210
211
212
        return p;
    }
};

Shucai Xiao's avatar
Shucai Xiao committed
213
214
215
216
217
218
219
220
221
222
223
224
225
226
227
228
229
230
231
232
233
234
235
236
237
238
struct test_exp
{
    migraphx::program create_program() const
    {
        migraphx::program p;
        migraphx::shape s{migraphx::shape::float_type, {6}};
        std::vector<float> data{0.1f, 0.2f, 1.f, 2.f, 0.6f, 10.f};
        auto x = p.add_literal(s, data);
        p.add_instruction(migraphx::op::exp{}, x);
        return p;
    }
};

struct test_log
{
    migraphx::program create_program() const
    {
        migraphx::program p;
        migraphx::shape s{migraphx::shape::float_type, {6}};
        std::vector<float> data{0.1f, 0.2f, 1.f, 2.f, 0.6f, 100.f};
        auto x = p.add_literal(s, data);
        p.add_instruction(migraphx::op::log{}, x);
        return p;
    }
};

239
240
241
242
243
244
245
246
247
248
249
250
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
251
252
253
254
255
256
257
258
259
260
261
262
263
264
265
266
267
268
269
270
271
272
273
274
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;
    }
};

275
276
277
278
279
280
281
282
283
284
285
286
287
288
289
290
291
292
293
294
295
296
297
298
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;
    }
};

299
300
301
302
303
304
305
306
307
308
309
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;
    }
};

310
311
312
313
314
315
316
317
318
319
320
321
322
323
324
325
326
327
328
329
330
331
332
333
334
335
336
337
338
339
340
341
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);
Khalique's avatar
Khalique committed
342
343
344
345
346
347
        return p;
    }
};

struct test_scale
{
Paul's avatar
Paul committed
348
    migraphx::program create_program() const
Khalique's avatar
Khalique committed
349
    {
Paul's avatar
Paul committed
350
351
        migraphx::program p;
        migraphx::shape s{migraphx::shape::float_type, {3}};
Khalique's avatar
Khalique committed
352
        auto x     = p.add_parameter("x", s);
Paul's avatar
Paul committed
353
354
355
        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
356
357
358
359
        return p;
    }
};

360
361
struct test_slice
{
Paul's avatar
Paul committed
362
    migraphx::program create_program() const
363
    {
Paul's avatar
Paul committed
364
365
        migraphx::program p;
        migraphx::shape s{migraphx::shape::int32_type, {2, 2, 4}};
366
        auto x      = p.add_parameter("x", s);
Paul's avatar
Paul committed
367
368
369
        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);
370
371
372
373
374

        return p;
    }
};

Paul's avatar
Paul committed
375
376
struct test_triadd
{
Paul's avatar
Paul committed
377
    migraphx::program create_program() const
Paul's avatar
Paul committed
378
    {
Paul's avatar
Paul committed
379
380
        migraphx::program p;
        migraphx::shape s{migraphx::shape::float_type, {3}};
Paul's avatar
Paul committed
381
382
383
        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
384
385
        auto sum = p.add_instruction(migraphx::op::add{}, x, y);
        p.add_instruction(migraphx::op::add{}, sum, z);
Paul's avatar
Paul committed
386
387
388
389
390
391
        return p;
    }
};

struct test_triadd2
{
Paul's avatar
Paul committed
392
    migraphx::program create_program() const
Paul's avatar
Paul committed
393
    {
Paul's avatar
Paul committed
394
395
396
        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
397
398
399
        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
400
401
402
        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
403
404
405
406
        return p;
    }
};

Paul's avatar
Paul committed
407
408
struct test_add_broadcast
{
Paul's avatar
Paul committed
409
    migraphx::program create_program() const
Paul's avatar
Paul committed
410
    {
Paul's avatar
Paul committed
411
412
413
414
415
416
        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
417
418
419
420
        return p;
    }
};

Paul's avatar
Paul committed
421
422
struct test_add_broadcast2
{
Paul's avatar
Paul committed
423
    migraphx::program create_program() const
Paul's avatar
Paul committed
424
    {
Paul's avatar
Paul committed
425
426
427
428
429
430
        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
431
432
433
434
        return p;
    }
};

Paul's avatar
Latest  
Paul committed
435
436
struct test_add_broadcast3
{
Paul's avatar
Paul committed
437
    migraphx::program create_program() const
Paul's avatar
Latest  
Paul committed
438
    {
Paul's avatar
Paul committed
439
440
441
442
443
444
        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
445
446
447
448
449
450
        return p;
    }
};

struct test_add_broadcast4
{
Paul's avatar
Paul committed
451
    migraphx::program create_program() const
Paul's avatar
Latest  
Paul committed
452
    {
Paul's avatar
Paul committed
453
454
455
456
457
458
        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
459
460
461
462
        return p;
    }
};

Paul's avatar
Paul committed
463
464
struct test_add_broadcast5
{
Paul's avatar
Paul committed
465
    migraphx::program create_program() const
Paul's avatar
Paul committed
466
    {
Paul's avatar
Paul committed
467
468
469
470
471
472
        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
473
474
475
476
        return p;
    }
};

Paul's avatar
Paul committed
477
478
struct test_triadd_broadcast
{
Paul's avatar
Paul committed
479
    migraphx::program create_program() const
Paul's avatar
Paul committed
480
    {
Paul's avatar
Paul committed
481
482
483
484
485
486
487
488
        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
489
490
491
492
        return p;
    }
};

493
494
495
496
497
498
struct test_sub
{
    migraphx::program create_program() const
    {
        migraphx::program p;
        migraphx::shape s{migraphx::shape::float_type, {3}};
Shucai Xiao's avatar
Shucai Xiao committed
499
500
501
        auto x    = p.add_parameter("x", s);
        auto y    = p.add_parameter("y", s);
        auto z    = p.add_parameter("z", s);
502
503
504
505
506
507
508
509
510
511
512
513
514
        auto diff = p.add_instruction(migraphx::op::sub{}, x, y);
        p.add_instruction(migraphx::op::sub{}, diff, z);
        return p;
    }
};

struct test_sub2
{
    migraphx::program create_program() const
    {
        migraphx::program p;
        migraphx::shape s{migraphx::shape::float_type, {2, 3}};
        migraphx::shape b{migraphx::shape::float_type, {3}};
Shucai Xiao's avatar
Shucai Xiao committed
515
516
517
518
        auto x    = p.add_parameter("x", s);
        auto y    = p.add_parameter("y", s);
        auto z    = p.add_parameter("z", b);
        auto zb   = p.add_instruction(migraphx::op::broadcast{1, s}, z);
519
520
521
522
523
524
        auto diff = p.add_instruction(migraphx::op::sub{}, x, y);
        p.add_instruction(migraphx::op::sub{}, diff, zb);
        return p;
    }
};

Paul's avatar
Paul committed
525
526
struct test_softmax
{
Paul's avatar
Paul committed
527
    migraphx::program create_program() const
Paul's avatar
Paul committed
528
    {
Paul's avatar
Paul committed
529
530
531
        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
532
533
534
535
536
537
        return p;
    }
};

struct test_softmax2
{
Paul's avatar
Paul committed
538
    migraphx::program create_program() const
Paul's avatar
Paul committed
539
    {
Paul's avatar
Paul committed
540
        migraphx::program p;
Paul's avatar
Paul committed
541
542
        auto x =
            p.add_parameter("x", migraphx::shape{migraphx::shape::float_type, {1, 1000, 1, 1}});
Paul's avatar
Paul committed
543
        p.add_instruction(migraphx::op::softmax{}, x);
Paul's avatar
Paul committed
544
545
546
547
        return p;
    }
};

Paul's avatar
Paul committed
548
549
struct test_conv
{
Paul's avatar
Paul committed
550
    migraphx::program create_program() const
Paul's avatar
Paul committed
551
    {
Paul's avatar
Paul committed
552
        migraphx::program p;
Paul's avatar
Paul committed
553
554
        auto input =
            p.add_parameter("x", migraphx::shape{migraphx::shape::float_type, {4, 3, 3, 3}});
Paul's avatar
Paul committed
555
        auto weights =
Paul's avatar
Paul committed
556
557
            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
558
559
560
561
        return p;
    }
};

Paul's avatar
Paul committed
562
563
struct test_conv2
{
Paul's avatar
Paul committed
564
    migraphx::program create_program() const
Paul's avatar
Paul committed
565
    {
Paul's avatar
Paul committed
566
        migraphx::program p;
Paul's avatar
Paul committed
567
        auto input =
Paul's avatar
Paul committed
568
            p.add_parameter("x", migraphx::shape{migraphx::shape::float_type, {1, 512, 28, 28}});
Paul's avatar
Paul committed
569
        auto weights =
Paul's avatar
Paul committed
570
571
            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
572
573
574
575
        return p;
    }
};

Khalique's avatar
Khalique committed
576
577
578
579
580
struct test_group_conv
{
    migraphx::program create_program() const
    {
        migraphx::program p;
Paul's avatar
Paul committed
581
        auto input =
Khalique's avatar
Khalique committed
582
            p.add_parameter("x", migraphx::shape{migraphx::shape::float_type, {1, 4, 16, 16}});
Paul's avatar
Paul committed
583
        auto weights =
Khalique's avatar
Khalique committed
584
585
586
587
            p.add_parameter("w", migraphx::shape{migraphx::shape::float_type, {4, 1, 3, 3}});
        migraphx::op::convolution op;
        op.group = 4;
        p.add_instruction(op, input, weights);
Paul's avatar
Paul committed
588
589
590
591
        return p;
    }
};

Paul's avatar
Paul committed
592
struct test_conv_relu
Paul's avatar
Paul committed
593
{
Paul's avatar
Paul committed
594
    migraphx::program create_program() const
Paul's avatar
Paul committed
595
    {
Paul's avatar
Paul committed
596
        migraphx::program p;
Paul's avatar
Paul committed
597
598
        auto input =
            p.add_parameter("x", migraphx::shape{migraphx::shape::float_type, {4, 3, 3, 3}});
Paul's avatar
Paul committed
599
        auto weights =
Paul's avatar
Paul committed
600
601
602
            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
603
604
605
606
        return p;
    }
};

Paul's avatar
Paul committed
607
608
struct test_conv_relu_half
{
Paul's avatar
Paul committed
609
    migraphx::program create_program() const
Paul's avatar
Paul committed
610
    {
Paul's avatar
Paul committed
611
        migraphx::program p;
Paul's avatar
Paul committed
612
613
        auto input =
            p.add_parameter("x", migraphx::shape{migraphx::shape::half_type, {4, 3, 3, 3}});
Paul's avatar
Paul committed
614
        auto weights =
Paul's avatar
Paul committed
615
616
617
            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
618
619
620
621
        return p;
    }
};

Paul's avatar
Paul committed
622
623
struct test_add_relu
{
Paul's avatar
Paul committed
624
    migraphx::program create_program() const
Paul's avatar
Paul committed
625
    {
Paul's avatar
Paul committed
626
627
628
629
630
        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
631
632
633
634
        return p;
    }
};

Khalique's avatar
Khalique committed
635
636
637
638
639
struct test_sigmoid
{
    migraphx::program create_program() const
    {
        migraphx::program p;
Khalique's avatar
Khalique committed
640
        auto x = p.add_parameter("x", migraphx::shape{migraphx::shape::float_type, {4, 3, 3, 3}});
Khalique's avatar
Khalique committed
641
642
643
644
645
646
647
648
649
650
        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
651
        auto x = p.add_parameter("x", migraphx::shape{migraphx::shape::float_type, {4, 3, 3, 3}});
Khalique's avatar
Khalique committed
652
        p.add_instruction(migraphx::op::abs{}, x);
Paul's avatar
Paul committed
653
654
655
656
        return p;
    }
};

657
658
struct test_leaky_relu
{
Paul's avatar
Paul committed
659
    migraphx::program create_program() const
660
    {
Paul's avatar
Paul committed
661
662
663
        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);
664
665
666
667
        return p;
    }
};

Khalique's avatar
Khalique committed
668
669
670
671
672
673
674
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);
675
676
677
678
        return p;
    }
};

Khalique's avatar
Khalique committed
679
struct test_relu_lrn
Khalique's avatar
Khalique committed
680
{
Khalique's avatar
Khalique committed
681
    migraphx::program create_program() const
Khalique's avatar
Khalique committed
682
    {
Khalique's avatar
Khalique committed
683
684
685
686
        migraphx::program p;
        auto x = p.add_parameter("x", migraphx::shape{migraphx::shape::float_type, {1, 5, 2, 2}});
        auto y = p.add_instruction(migraphx::op::relu{}, x);
        p.add_instruction(migraphx::op::lrn{0.0001, 0.75, 1.0, 5}, y);
Khalique's avatar
Khalique committed
687
688
689
690
        return p;
    }
};

Paul's avatar
Paul committed
691
692
struct test_conv_pooling
{
Paul's avatar
Paul committed
693
    migraphx::program create_program() const
Paul's avatar
Paul committed
694
    {
Paul's avatar
Paul committed
695
        migraphx::program p;
Paul's avatar
Paul committed
696
        auto input =
Paul's avatar
Paul committed
697
            p.add_parameter("x", migraphx::shape{migraphx::shape::float_type, {4, 3, 32, 32}});
Paul's avatar
Paul committed
698
        auto weights =
Paul's avatar
Paul committed
699
700
701
702
            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
703
704
705
706
        return p;
    }
};

707
708
struct test_global_avg_pooling
{
Paul's avatar
Paul committed
709
    migraphx::program create_program() const
710
    {
Paul's avatar
Paul committed
711
        migraphx::program p;
712
        auto input =
Paul's avatar
Paul committed
713
714
            p.add_parameter("x", migraphx::shape{migraphx::shape::float_type, {1, 3, 16, 16}});
        auto op    = migraphx::op::pooling{"average"};
715
        auto lens  = input->get_shape().lens();
Khalique's avatar
Khalique committed
716
        op.lengths = {lens[2], lens[3]};
717
718
719
720
721
722
723
        p.add_instruction(op, input);
        return p;
    }
};

struct test_global_max_pooling
{
Paul's avatar
Paul committed
724
    migraphx::program create_program() const
725
    {
Paul's avatar
Paul committed
726
        migraphx::program p;
727
        auto input =
Paul's avatar
Paul committed
728
729
            p.add_parameter("x", migraphx::shape{migraphx::shape::float_type, {1, 3, 16, 16}});
        auto op    = migraphx::op::pooling{"max"};
730
        auto lens  = input->get_shape().lens();
Khalique's avatar
Khalique committed
731
        op.lengths = {lens[2], lens[3]};
732
733
734
735
736
        p.add_instruction(op, input);
        return p;
    }
};

Paul's avatar
Paul committed
737
738
struct test_gemm
{
Paul's avatar
Paul committed
739
    migraphx::program create_program() const
Paul's avatar
Paul committed
740
    {
Paul's avatar
Paul committed
741
742
743
744
        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
745
746
747
748
        return p;
    }
};

Paul's avatar
Paul committed
749
750
struct test_gemm_half
{
Paul's avatar
Paul committed
751
    migraphx::program create_program() const
Paul's avatar
Paul committed
752
    {
Paul's avatar
Paul committed
753
754
755
756
        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
757
758
759
760
        return p;
    }
};

Paul's avatar
Paul committed
761
762
struct test_gemm_ld
{
Paul's avatar
Paul committed
763
    migraphx::program create_program() const
Paul's avatar
Paul committed
764
    {
Paul's avatar
Paul committed
765
        migraphx::program p;
Paul's avatar
Paul committed
766
767
768
769
        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
770
        p.add_instruction(migraphx::op::dot{}, a, b);
Paul's avatar
Paul committed
771
772
773
774
        return p;
    }
};

775
776
struct test_gemm_transposeb
{
Paul's avatar
Paul committed
777
    migraphx::program create_program() const
778
    {
Paul's avatar
Paul committed
779
780
781
782
783
        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);
784
785
786
787
788
789
        return p;
    }
};

struct test_gemm_transposea
{
Paul's avatar
Paul committed
790
    migraphx::program create_program() const
791
    {
Paul's avatar
Paul committed
792
793
794
795
796
        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);
797
798
799
800
801
802
        return p;
    }
};

struct test_gemm_transposeab
{
Paul's avatar
Paul committed
803
    migraphx::program create_program() const
804
    {
Paul's avatar
Paul committed
805
806
807
808
809
810
        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);
811
812
813
814
        return p;
    }
};

815
816
struct test_contiguous
{
Paul's avatar
Paul committed
817
    migraphx::program create_program() const
818
    {
Paul's avatar
Paul committed
819
820
        migraphx::program p;
        migraphx::shape s{migraphx::shape::float_type, {4, 4, 4, 3}, {48, 4, 1, 16}};
821
        auto x = p.add_parameter("x", s);
Paul's avatar
Paul committed
822
        p.add_instruction(migraphx::op::contiguous{}, x);
Paul's avatar
Paul committed
823
        EXPECT(p.get_shape().standard());
824
825
826
827
        return p;
    }
};

828
struct test_transpose
829
{
Paul's avatar
Paul committed
830
    migraphx::program create_program() const
831
    {
Paul's avatar
Paul committed
832
833
        migraphx::program p;
        migraphx::shape s{migraphx::shape::float_type, {4, 3, 4, 4}};
834
835
        auto x                    = p.add_parameter("x", s);
        std::vector<int64_t> perm = {0, 2, 3, 1};
Paul's avatar
Paul committed
836
837
        auto l                    = p.add_instruction(migraphx::op::transpose{perm}, x);
        p.add_instruction(migraphx::op::contiguous{}, l);
838
839
840
        return p;
    }
};
841

Paul's avatar
Paul committed
842
843
844
845
846
847
848
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
849
    migraphx::program create_program() const
Paul's avatar
Paul committed
850
    {
Paul's avatar
Paul committed
851
        migraphx::program p;
Paul's avatar
Paul committed
852

Paul's avatar
Paul committed
853
854
        migraphx::shape s{migraphx::shape::float_type, {batches, channels, height, width}};
        migraphx::shape vars{migraphx::shape::float_type, {channels}};
Paul's avatar
Paul committed
855
        auto x        = p.add_parameter("x", s);
Paul's avatar
Paul committed
856
857
858
859
860
        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
861
862
863
864
        return p;
    }
};

wsttiger's avatar
wsttiger committed
865
866
867
868
869
870
871
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
872
    migraphx::program create_program() const
wsttiger's avatar
wsttiger committed
873
    {
Paul's avatar
Paul committed
874
        migraphx::program p;
wsttiger's avatar
wsttiger committed
875

Paul's avatar
Paul committed
876
877
        migraphx::shape s{migraphx::shape::float_type, {batches, channels, height, width}};
        migraphx::shape vars{migraphx::shape::float_type, {channels}};
wsttiger's avatar
wsttiger committed
878
        auto x        = p.add_parameter("x", s);
Paul's avatar
Paul committed
879
880
881
882
883
        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
884
885
886
887
        return p;
    }
};

Paul's avatar
Paul committed
888
889
struct test_conv_bn
{
Paul's avatar
Paul committed
890
    migraphx::program create_program() const
Paul's avatar
Paul committed
891
    {
Paul's avatar
Paul committed
892
        migraphx::program p;
Paul's avatar
Paul committed
893

Paul's avatar
Paul committed
894
895
896
        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
897
898
        auto x        = p.add_parameter("x", xs);
        auto w        = p.add_parameter("w", ws);
Paul's avatar
Paul committed
899
900
901
902
903
904
        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
905
906
907
908
        return p;
    }
};

Paul's avatar
Paul committed
909
910
struct test_conv_bn_relu_pooling
{
Paul's avatar
Paul committed
911
    migraphx::program create_program() const
Paul's avatar
Paul committed
912
    {
Paul's avatar
Paul committed
913
        migraphx::program p;
Paul's avatar
Paul committed
914

Paul's avatar
Paul committed
915
916
917
        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
918
919
        auto x        = p.add_parameter("x", xs);
        auto w        = p.add_parameter("w", ws);
Paul's avatar
Paul committed
920
921
922
923
924
        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
925
        auto bn       = p.add_instruction(
Paul's avatar
Paul committed
926
927
928
            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
929
930
931
932
        return p;
    }
};

933
934
struct test_concat
{
Paul's avatar
Paul committed
935
    migraphx::program create_program() const
936
    {
Paul's avatar
Paul committed
937
        migraphx::program p;
wsttiger's avatar
wsttiger committed
938
        std::size_t axis = 1;
Paul's avatar
Paul committed
939
940
941
        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}};
942
943
944
        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
945
        p.add_instruction(migraphx::op::concat{axis}, l0, l1, l2);
946
947
948
949
950
951
        return p;
    }
};

struct test_concat2
{
Paul's avatar
Paul committed
952
    migraphx::program create_program() const
953
    {
Paul's avatar
Paul committed
954
        migraphx::program p;
wsttiger's avatar
wsttiger committed
955
        std::size_t axis = 0;
Paul's avatar
Paul committed
956
957
958
        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}};
959
960
961
        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
962
        p.add_instruction(migraphx::op::concat{axis}, l0, l1, l2);
963
964
965
966
        return p;
    }
};

wsttiger's avatar
wsttiger committed
967
968
struct test_concat_relu
{
Paul's avatar
Paul committed
969
    migraphx::program create_program() const
wsttiger's avatar
wsttiger committed
970
    {
Paul's avatar
Paul committed
971
        migraphx::program p;
wsttiger's avatar
wsttiger committed
972
        std::size_t axis = 0;
Paul's avatar
Paul committed
973
974
975
        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
976
977
978
        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
979
980
981
982
983
        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
984
985
986
987
        return p;
    }
};

988
989
990
991
992
993
struct test_pad
{
    migraphx::program create_program() const
    {
        migraphx::program p;
        migraphx::shape s0{migraphx::shape::int32_type, {1, 96, 165, 165}};
Khalique's avatar
Khalique committed
994
995
996
997
        std::vector<int64_t> pads0 = {0, 0, 0, 0, 0, 0, 1, 1};
        std::vector<int64_t> pads1 = {0, 0, 0, 0, 1, 1, 1, 1};
        std::vector<int64_t> pads2 = {1, 1, 1, 1, 0, 0, 0, 0};
        std::vector<int64_t> pads3 = {1, 0, 1, 0, 1, 0, 2, 0};
Khalique's avatar
Khalique committed
998
        auto l0                    = p.add_parameter("x", s0);
Khalique's avatar
Khalique committed
999
1000
1001
        p.add_instruction(migraphx::op::pad{pads0}, l0);
        p.add_instruction(migraphx::op::pad{pads1}, l0);
        p.add_instruction(migraphx::op::pad{pads2}, l0);
Khalique's avatar
Khalique committed
1002
        p.add_instruction(migraphx::op::pad{pads3}, l0);
Khalique's avatar
Khalique committed
1003
        return p;
Khalique's avatar
Khalique committed
1004
    }
Khalique's avatar
Khalique committed
1005
};
Khalique's avatar
Khalique committed
1006

1007
1008
1009
1010
1011
1012
1013
1014
1015
struct test_pooling_autopad
{
    migraphx::program create_program() const
    {
        migraphx::program p;
        migraphx::shape s0{migraphx::shape::float_type, {1, 3, 63, 63}};
        auto l0 = p.add_parameter("x", s0);
        migraphx::op::pooling op{"max"};
        op.padding_mode = migraphx::op::padding_mode_t::same;
Khalique's avatar
Khalique committed
1016
1017
        op.lengths      = {2, 2};
        op.stride       = {2, 2};
1018
1019
1020
1021
1022
        p.add_instruction(op, l0);
        return p;
    }
};

1023
1024
1025
1026
1027
1028
1029
1030
struct test_gather
{
    migraphx::program create_program() const
    {
        migraphx::program p;
        migraphx::shape s{migraphx::shape::float_type, {3, 3}};
        migraphx::shape s_indices{migraphx::shape::int32_type, {2, 2}};
        std::vector<int> indices{1, 2, 2, 1};
Shucai Xiao's avatar
Shucai Xiao committed
1031
1032
        auto a0  = p.add_parameter("data", s);
        auto a1  = p.add_literal(migraphx::literal{s_indices, indices});
1033
        int axis = 0;
1034
        p.add_instruction(migraphx::op::gather{axis}, a0, a1);
1035
1036
1037
1038
        return p;
    }
};

1039
1040
1041
1042
1043
1044
1045
1046
1047
1048
1049
struct test_gather_neg_axis
{
    migraphx::program create_program() const
    {
        migraphx::program p;
        migraphx::shape s{migraphx::shape::float_type, {3, 3}};
        migraphx::shape s_indices{migraphx::shape::int32_type, {2, 2}};
        std::vector<int> indices{1, 2, 2, 1};
        auto a0  = p.add_parameter("data", s);
        auto a1  = p.add_literal(migraphx::literal{s_indices, indices});
        int axis = -1;
1050
        p.add_instruction(migraphx::op::gather{axis}, a0, a1);
wsttiger's avatar
wsttiger committed
1051
1052
1053
1054
1055
1056
        return p;
    }
};

void manual_identity()
{
Paul's avatar
Paul committed
1057
    migraphx::program p;
wsttiger's avatar
wsttiger committed
1058
    std::vector<float> data0 = {0, 1, 2, 3};
Paul's avatar
Paul committed
1059
1060
1061
1062
1063
    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
1064
1065
    for(auto&& x : p.get_parameter_shapes())
    {
Paul's avatar
Paul committed
1066
        m[x.first] = migraphx::gpu::to_gpu(migraphx::generate_argument(x.second));
wsttiger's avatar
wsttiger committed
1067
    }
Paul's avatar
Paul committed
1068
    auto result = migraphx::gpu::from_gpu(p.eval(m));
wsttiger's avatar
wsttiger committed
1069
1070
1071
1072
1073
    std::cout << result << std::endl;
}

void manual_test_concat_relu()
{
Paul's avatar
Paul committed
1074
    migraphx::program p;
wsttiger's avatar
wsttiger committed
1075
    std::size_t axis         = 0;
wsttiger's avatar
wsttiger committed
1076
1077
1078
    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
1079
1080
1081
1082
1083
1084
1085
1086
1087
1088
1089
1090
1091
1092
    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
1093
1094
    for(auto&& x : p.get_parameter_shapes())
    {
Paul's avatar
Paul committed
1095
        m[x.first] = migraphx::gpu::to_gpu(migraphx::generate_argument(x.second));
wsttiger's avatar
wsttiger committed
1096
    }
Paul's avatar
Paul committed
1097
    auto result = migraphx::gpu::from_gpu(p.eval(m));
wsttiger's avatar
wsttiger committed
1098
1099
1100
    std::cout << result << std::endl;
}

Paul's avatar
Paul committed
1101
1102
struct test_conv_bn_relu_pooling2
{
Paul's avatar
Paul committed
1103
1104
    static migraphx::instruction_ref
    add_bn(migraphx::program& p, migraphx::instruction_ref x, std::size_t channels)
Paul's avatar
Paul committed
1105
    {
Paul's avatar
Paul committed
1106
        migraphx::shape vars{migraphx::shape::float_type, {channels}};
Paul's avatar
Paul committed
1107
1108
1109
1110
1111
        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
1112
        return p.add_instruction(
Paul's avatar
Paul committed
1113
            migraphx::op::batch_norm_inference{}, x, scale, bias, mean, variance);
Paul's avatar
Paul committed
1114
    }
Paul's avatar
Paul committed
1115
    migraphx::program create_program() const
Paul's avatar
Paul committed
1116
    {
Paul's avatar
Paul committed
1117
        migraphx::program p;
Paul's avatar
Paul committed
1118

Paul's avatar
Paul committed
1119
1120
1121
1122
        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
1123
1124
        auto x1    = p.add_parameter("x1", xs1);
        auto w1    = p.add_parameter("w1", ws1);
Paul's avatar
Paul committed
1125
        auto conv1 = p.add_instruction(migraphx::op::convolution{{0, 0}, {1, 1}, {1, 1}}, x1, w1);
Paul's avatar
Paul committed
1126
1127
1128
        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
1129
        auto conv2 = p.add_instruction(migraphx::op::convolution{{0, 0}, {2, 2}, {1, 1}}, x2, w2);
Paul's avatar
Paul committed
1130
        auto bn2   = add_bn(p, conv2, 2048);
Paul's avatar
Paul committed
1131
1132
1133
        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
1134
1135
1136
1137
        return p;
    }
};

1138
1139
1140
1141
1142
struct test_rnn_forward
{
    migraphx::program create_program() const
    {
        std::size_t batch_size  = 2;
1143
        std::size_t seq_len     = 1;
1144
1145
1146
        std::size_t hidden_size = 4;
        std::size_t input_size  = 3;
        std::size_t num_dirct   = 1;
Shucai Xiao's avatar
Shucai Xiao committed
1147
        float clip              = 0.0f;
1148
1149
1150
1151
1152
1153

        migraphx::program p;
        migraphx::shape in_shape{migraphx::shape::float_type, {seq_len, batch_size, input_size}};
        migraphx::shape w_shape{migraphx::shape::float_type, {num_dirct, hidden_size, input_size}};
        migraphx::shape r_shape{migraphx::shape::float_type, {num_dirct, hidden_size, hidden_size}};
        migraphx::shape b_shape{migraphx::shape::float_type, {num_dirct, 2 * hidden_size}};
1154
        migraphx::shape ih_shape{migraphx::shape::float_type, {num_dirct, batch_size, hidden_size}};
1155

Shucai Xiao's avatar
Shucai Xiao committed
1156
        auto seq  = p.add_parameter("seq", in_shape);
1157
1158
1159
        auto w    = p.add_parameter("w", w_shape);
        auto r    = p.add_parameter("r", r_shape);
        auto bias = p.add_parameter("bias", b_shape);
Shucai Xiao's avatar
Shucai Xiao committed
1160
        auto ih   = p.add_parameter("ih", ih_shape);
1161
        auto und  = p.add_instruction(migraphx::op::undefined{});
Shucai Xiao's avatar
Shucai Xiao committed
1162

1163
1164
1165
1166
1167
1168
1169
1170
1171
        auto output =
            p.add_instruction(migraphx::op::rnn{hidden_size,
                                                {migraphx::op::tanh{}, migraphx::op::tanh{}},
                                                migraphx::op::rnn::forward,
                                                clip},
                              seq,
                              w,
                              r,
                              bias,
1172
                              und,
1173
1174
1175
1176
1177
1178
1179
1180
1181
1182
1183
1184
1185
1186
1187
1188
1189
1190
1191
1192
1193
1194
1195
1196
1197
1198
                              ih);
        p.add_instruction(migraphx::op::rnn_last_output{}, output);

        return p;
    }
};

struct test_rnn_forward10
{
    migraphx::program create_program() const
    {
        std::size_t batch_size  = 2;
        std::size_t seq_len     = 10;
        std::size_t hidden_size = 4;
        std::size_t input_size  = 3;
        std::size_t num_dirct   = 1;
        float clip              = 0.0f;

        migraphx::program p;
        migraphx::shape in_shape{migraphx::shape::float_type, {seq_len, batch_size, input_size}};
        migraphx::shape w_shape{migraphx::shape::float_type, {num_dirct, hidden_size, input_size}};
        migraphx::shape r_shape{migraphx::shape::float_type, {num_dirct, hidden_size, hidden_size}};
        migraphx::shape b_shape{migraphx::shape::float_type, {num_dirct, 2 * hidden_size}};
        migraphx::shape ih_shape{migraphx::shape::float_type, {num_dirct, batch_size, hidden_size}};

        auto seq  = p.add_parameter("seq", in_shape);
Shucai Xiao's avatar
Shucai Xiao committed
1199
1200
        auto w    = p.add_parameter("w", w_shape);
        auto r    = p.add_parameter("r", r_shape);
1201
        auto bias = p.add_parameter("bias", b_shape);
1202
        auto ih   = p.add_parameter("ih", ih_shape);
1203
        auto und  = p.add_instruction(migraphx::op::undefined{});
1204

Shucai Xiao's avatar
Shucai Xiao committed
1205
1206
1207
1208
1209
1210
1211
1212
1213
        auto output =
            p.add_instruction(migraphx::op::rnn{hidden_size,
                                                {migraphx::op::tanh{}, migraphx::op::tanh{}},
                                                migraphx::op::rnn::forward,
                                                clip},
                              seq,
                              w,
                              r,
                              bias,
1214
                              und,
Shucai Xiao's avatar
Shucai Xiao committed
1215
                              ih);
Shucai Xiao's avatar
Shucai Xiao committed
1216
        p.add_instruction(migraphx::op::rnn_last_output{}, output);
1217
1218
1219
1220
1221
1222
1223
1224
1225
1226

        return p;
    }
};

struct test_rnn_reverse
{
    migraphx::program create_program() const
    {
        std::size_t batch_size  = 2;
1227
        std::size_t seq_len     = 1;
1228
1229
1230
        std::size_t hidden_size = 4;
        std::size_t input_size  = 3;
        std::size_t num_dirct   = 1;
Shucai Xiao's avatar
Shucai Xiao committed
1231
        float clip              = 0.0f;
1232
1233
1234
1235
1236
1237

        migraphx::program p;
        migraphx::shape in_shape{migraphx::shape::float_type, {seq_len, batch_size, input_size}};
        migraphx::shape w_shape{migraphx::shape::float_type, {num_dirct, hidden_size, input_size}};
        migraphx::shape r_shape{migraphx::shape::float_type, {num_dirct, hidden_size, hidden_size}};
        migraphx::shape b_shape{migraphx::shape::float_type, {num_dirct, 2 * hidden_size}};
1238
        migraphx::shape ih_shape{migraphx::shape::float_type, {num_dirct, batch_size, hidden_size}};
1239

Shucai Xiao's avatar
Shucai Xiao committed
1240
        auto seq  = p.add_parameter("seq", in_shape);
1241
1242
1243
        auto w    = p.add_parameter("w", w_shape);
        auto r    = p.add_parameter("r", r_shape);
        auto bias = p.add_parameter("bias", b_shape);
Shucai Xiao's avatar
Shucai Xiao committed
1244
        auto ih   = p.add_parameter("ih", ih_shape);
1245
        auto und  = p.add_instruction(migraphx::op::undefined{});
1246
1247
1248
1249
1250
1251
1252
1253
1254

        p.add_instruction(migraphx::op::rnn{hidden_size,
                                            {migraphx::op::tanh{}, migraphx::op::tanh{}},
                                            migraphx::op::rnn::reverse,
                                            clip},
                          seq,
                          w,
                          r,
                          bias,
1255
                          und,
1256
1257
1258
1259
1260
1261
1262
1263
1264
1265
1266
1267
1268
1269
1270
1271
1272
1273
1274
1275
1276
1277
1278
1279
1280
                          ih);

        return p;
    }
};

struct test_rnn_reverse2
{
    migraphx::program create_program() const
    {
        std::size_t batch_size  = 2;
        std::size_t seq_len     = 2;
        std::size_t hidden_size = 4;
        std::size_t input_size  = 3;
        std::size_t num_dirct   = 1;
        float clip              = 0.0f;

        migraphx::program p;
        migraphx::shape in_shape{migraphx::shape::float_type, {seq_len, batch_size, input_size}};
        migraphx::shape w_shape{migraphx::shape::float_type, {num_dirct, hidden_size, input_size}};
        migraphx::shape r_shape{migraphx::shape::float_type, {num_dirct, hidden_size, hidden_size}};
        migraphx::shape b_shape{migraphx::shape::float_type, {num_dirct, 2 * hidden_size}};
        migraphx::shape ih_shape{migraphx::shape::float_type, {num_dirct, batch_size, hidden_size}};

        auto seq  = p.add_parameter("seq", in_shape);
Shucai Xiao's avatar
Shucai Xiao committed
1281
1282
        auto w    = p.add_parameter("w", w_shape);
        auto r    = p.add_parameter("r", r_shape);
1283
        auto bias = p.add_parameter("bias", b_shape);
1284
        auto ih   = p.add_parameter("ih", ih_shape);
1285
        auto und  = p.add_instruction(migraphx::op::undefined{});
1286

Shucai Xiao's avatar
Shucai Xiao committed
1287
1288
1289
1290
1291
1292
1293
1294
        p.add_instruction(migraphx::op::rnn{hidden_size,
                                            {migraphx::op::tanh{}, migraphx::op::tanh{}},
                                            migraphx::op::rnn::reverse,
                                            clip},
                          seq,
                          w,
                          r,
                          bias,
1295
                          und,
Shucai Xiao's avatar
Shucai Xiao committed
1296
                          ih);
Shucai Xiao's avatar
Shucai Xiao committed
1297

1298
1299
1300
1301
        return p;
    }
};

1302
1303
1304
1305
1306
1307
1308
1309
1310
1311
1312
1313
1314
1315
1316
1317
struct test_rnn_3args
{
    migraphx::program create_program() const
    {
        std::size_t batch_size  = 2;
        std::size_t seq_len     = 1;
        std::size_t hidden_size = 4;
        std::size_t input_size  = 3;
        std::size_t num_dirct   = 1;
        float clip              = 0.0f;

        migraphx::program p;
        migraphx::shape in_shape{migraphx::shape::float_type, {seq_len, batch_size, input_size}};
        migraphx::shape w_shape{migraphx::shape::float_type, {num_dirct, hidden_size, input_size}};
        migraphx::shape r_shape{migraphx::shape::float_type, {num_dirct, hidden_size, hidden_size}};

Shucai Xiao's avatar
Shucai Xiao committed
1318
1319
1320
        auto seq = p.add_parameter("seq", in_shape);
        auto w   = p.add_parameter("w", w_shape);
        auto r   = p.add_parameter("r", r_shape);
1321
1322
1323
1324
1325
1326
1327
1328
1329
1330
1331
1332
1333
1334
1335
1336
1337
1338
1339
1340
1341
1342
1343
1344
1345
1346
1347
1348
1349
1350
1351
1352
1353
1354
1355
1356
1357
1358
1359
1360
1361
1362
1363
1364
1365
1366
1367
1368
1369
1370
1371
1372
1373
1374
1375
1376
1377
1378
1379
1380
1381
1382
1383
1384
1385
1386
1387
1388
1389

        p.add_instruction(migraphx::op::rnn{hidden_size,
                                            {migraphx::op::tanh{}, migraphx::op::tanh{}},
                                            migraphx::op::rnn::reverse,
                                            clip},
                          seq,
                          w,
                          r);

        return p;
    }
};

struct test_rnn_4args
{
    migraphx::program create_program() const
    {
        std::size_t batch_size  = 2;
        std::size_t seq_len     = 5;
        std::size_t hidden_size = 4;
        std::size_t input_size  = 3;
        std::size_t num_dirct   = 1;
        float clip              = 0.0f;

        migraphx::program p;
        migraphx::shape in_shape{migraphx::shape::float_type, {seq_len, batch_size, input_size}};
        migraphx::shape w_shape{migraphx::shape::float_type, {num_dirct, hidden_size, input_size}};
        migraphx::shape r_shape{migraphx::shape::float_type, {num_dirct, hidden_size, hidden_size}};
        migraphx::shape b_shape{migraphx::shape::float_type, {num_dirct, 2 * hidden_size}};

        auto seq  = p.add_parameter("seq", in_shape);
        auto w    = p.add_parameter("w", w_shape);
        auto r    = p.add_parameter("r", r_shape);
        auto bias = p.add_parameter("bias", b_shape);

        p.add_instruction(migraphx::op::rnn{hidden_size,
                                            {migraphx::op::tanh{}, migraphx::op::tanh{}},
                                            migraphx::op::rnn::reverse,
                                            clip},
                          seq,
                          w,
                          r,
                          bias);

        return p;
    }
};

struct test_rnn_5args
{
    migraphx::program create_program() const
    {
        std::size_t batch_size  = 2;
        std::size_t seq_len     = 10;
        std::size_t hidden_size = 4;
        std::size_t input_size  = 3;
        std::size_t num_dirct   = 1;
        float clip              = 0.0f;

        migraphx::program p;
        migraphx::shape in_shape{migraphx::shape::float_type, {seq_len, batch_size, input_size}};
        migraphx::shape w_shape{migraphx::shape::float_type, {num_dirct, hidden_size, input_size}};
        migraphx::shape r_shape{migraphx::shape::float_type, {num_dirct, hidden_size, hidden_size}};
        migraphx::shape b_shape{migraphx::shape::float_type, {num_dirct, 2 * hidden_size}};

        auto seq  = p.add_parameter("seq", in_shape);
        auto w    = p.add_parameter("w", w_shape);
        auto r    = p.add_parameter("r", r_shape);
        auto bias = p.add_parameter("bias", b_shape);
1390
        auto und  = p.add_instruction(migraphx::op::undefined{});
1391
1392
1393
1394
1395
1396
1397
1398
1399

        auto output =
            p.add_instruction(migraphx::op::rnn{hidden_size,
                                                {migraphx::op::tanh{}, migraphx::op::tanh{}},
                                                migraphx::op::rnn::forward,
                                                clip},
                              seq,
                              w,
                              r,
1400
1401
                              bias,
                              und);
1402
1403
1404
1405
1406
1407
        p.add_instruction(migraphx::op::rnn_last_output{}, output);

        return p;
    }
};

1408
1409
1410
1411
1412
struct test_rnn_bidirectional
{
    migraphx::program create_program() const
    {
        std::size_t batch_size  = 2;
1413
        std::size_t seq_len     = 1;
1414
1415
1416
        std::size_t hidden_size = 4;
        std::size_t input_size  = 3;
        std::size_t num_dirct   = 2;
Shucai Xiao's avatar
Shucai Xiao committed
1417
        float clip              = 0.0f;
1418
1419
1420
1421
1422
1423

        migraphx::program p;
        migraphx::shape in_shape{migraphx::shape::float_type, {seq_len, batch_size, input_size}};
        migraphx::shape w_shape{migraphx::shape::float_type, {num_dirct, hidden_size, input_size}};
        migraphx::shape r_shape{migraphx::shape::float_type, {num_dirct, hidden_size, hidden_size}};
        migraphx::shape b_shape{migraphx::shape::float_type, {num_dirct, 2 * hidden_size}};
1424
        migraphx::shape ih_shape{migraphx::shape::float_type, {num_dirct, batch_size, hidden_size}};
1425

Shucai Xiao's avatar
Shucai Xiao committed
1426
        auto seq  = p.add_parameter("seq", in_shape);
1427
1428
1429
        auto w    = p.add_parameter("w", w_shape);
        auto r    = p.add_parameter("r", r_shape);
        auto bias = p.add_parameter("bias", b_shape);
Shucai Xiao's avatar
Shucai Xiao committed
1430
        auto ih   = p.add_parameter("ih", ih_shape);
1431
        auto und  = p.add_instruction(migraphx::op::undefined{});
1432
1433
1434
1435
1436
1437
1438
1439
1440
1441

        auto output =
            p.add_instruction(migraphx::op::rnn{hidden_size,
                                                {migraphx::op::tanh{}, migraphx::op::tanh{}},
                                                migraphx::op::rnn::bidirectional,
                                                clip},
                              seq,
                              w,
                              r,
                              bias,
1442
                              und,
1443
1444
1445
1446
1447
1448
1449
1450
1451
1452
1453
1454
1455
1456
1457
1458
1459
1460
1461
1462
1463
1464
1465
1466
1467
1468
                              ih);
        p.add_instruction(migraphx::op::rnn_last_output{}, output);

        return p;
    }
};

struct test_rnn_bidirectional10
{
    migraphx::program create_program() const
    {
        std::size_t batch_size  = 2;
        std::size_t seq_len     = 10;
        std::size_t hidden_size = 4;
        std::size_t input_size  = 3;
        std::size_t num_dirct   = 2;
        float clip              = 0.0f;

        migraphx::program p;
        migraphx::shape in_shape{migraphx::shape::float_type, {seq_len, batch_size, input_size}};
        migraphx::shape w_shape{migraphx::shape::float_type, {num_dirct, hidden_size, input_size}};
        migraphx::shape r_shape{migraphx::shape::float_type, {num_dirct, hidden_size, hidden_size}};
        migraphx::shape b_shape{migraphx::shape::float_type, {num_dirct, 2 * hidden_size}};
        migraphx::shape ih_shape{migraphx::shape::float_type, {num_dirct, batch_size, hidden_size}};

        auto seq  = p.add_parameter("seq", in_shape);
Shucai Xiao's avatar
Shucai Xiao committed
1469
1470
        auto w    = p.add_parameter("w", w_shape);
        auto r    = p.add_parameter("r", r_shape);
1471
        auto bias = p.add_parameter("bias", b_shape);
1472
        auto ih   = p.add_parameter("ih", ih_shape);
1473
        auto und  = p.add_instruction(migraphx::op::undefined{});
Shucai Xiao's avatar
Shucai Xiao committed
1474
1475
1476
1477
1478
1479
1480
1481
1482
        auto output =
            p.add_instruction(migraphx::op::rnn{hidden_size,
                                                {migraphx::op::tanh{}, migraphx::op::tanh{}},
                                                migraphx::op::rnn::bidirectional,
                                                clip},
                              seq,
                              w,
                              r,
                              bias,
1483
                              und,
Shucai Xiao's avatar
Shucai Xiao committed
1484
                              ih);
Shucai Xiao's avatar
Shucai Xiao committed
1485
        p.add_instruction(migraphx::op::rnn_last_output{}, output);
Shucai Xiao's avatar
Shucai Xiao committed
1486

1487
1488
1489
1490
        return p;
    }
};

1491
1492
1493
1494
1495
1496
1497
1498
1499
1500
1501
1502
1503
1504
1505
struct test_rnn_bi_3args
{
    migraphx::program create_program() const
    {
        std::size_t batch_size  = 2;
        std::size_t seq_len     = 10;
        std::size_t hidden_size = 4;
        std::size_t input_size  = 3;
        std::size_t num_dirct   = 2;
        float clip              = 0.0f;

        migraphx::program p;
        migraphx::shape in_shape{migraphx::shape::float_type, {seq_len, batch_size, input_size}};
        migraphx::shape w_shape{migraphx::shape::float_type, {num_dirct, hidden_size, input_size}};
        migraphx::shape r_shape{migraphx::shape::float_type, {num_dirct, hidden_size, hidden_size}};
1506
1507
        migraphx::shape b_shape{migraphx::shape::float_type, {num_dirct, 2 * hidden_size}};
        migraphx::shape ih_shape{migraphx::shape::float_type, {num_dirct, batch_size, hidden_size}};
1508

Shucai Xiao's avatar
Shucai Xiao committed
1509
1510
1511
        auto seq = p.add_parameter("seq", in_shape);
        auto w   = p.add_parameter("w", w_shape);
        auto r   = p.add_parameter("r", r_shape);
1512
1513
1514
1515
1516
1517
1518
1519
1520
1521
1522
1523
1524
1525
        auto output =
            p.add_instruction(migraphx::op::rnn{hidden_size,
                                                {migraphx::op::tanh{}, migraphx::op::tanh{}},
                                                migraphx::op::rnn::bidirectional,
                                                clip},
                              seq,
                              w,
                              r);
        p.add_instruction(migraphx::op::rnn_last_output{}, output);

        return p;
    }
};

Paul's avatar
Paul committed
1526
1527
int main()
{
Khalique's avatar
Khalique committed
1528
    verify_program<test_relu_lrn>();
1529
    verify_program<test_pooling_autopad>();
Khalique's avatar
Khalique committed
1530
    verify_program<test_abs>();
1531
1532
    verify_program<test_concat>();
    verify_program<test_concat2>();
wsttiger's avatar
wsttiger committed
1533
    verify_program<test_concat_relu>();
1534
    verify_program<test_pad>();
Paul's avatar
Paul committed
1535
    verify_program<test_add>();
Paul's avatar
Paul committed
1536
    verify_program<test_add_half>();
Khalique's avatar
Khalique committed
1537
    verify_program<test_mul>();
Shucai Xiao's avatar
Shucai Xiao committed
1538
1539
    verify_program<test_exp>();
    verify_program<test_log>();
1540
    verify_program<test_sin>();
1541
1542
    verify_program<test_cos>();
    verify_program<test_tan>();
1543
1544
    verify_program<test_sinh>();
    verify_program<test_cosh>();
1545
    verify_program<test_tanh>();
1546
1547
1548
    verify_program<test_asin>();
    verify_program<test_acos>();
    verify_program<test_atan>();
Khalique's avatar
Khalique committed
1549
    verify_program<test_scale>();
Paul's avatar
Paul committed
1550
1551
    verify_program<test_triadd>();
    verify_program<test_triadd2>();
Paul's avatar
Paul committed
1552
    verify_program<test_add_broadcast>();
Paul's avatar
Paul committed
1553
    verify_program<test_add_broadcast2>();
Paul's avatar
Latest  
Paul committed
1554
1555
    verify_program<test_add_broadcast3>();
    verify_program<test_add_broadcast4>();
Paul's avatar
Paul committed
1556
    verify_program<test_add_broadcast5>();
Paul's avatar
Paul committed
1557
    verify_program<test_triadd_broadcast>();
1558
1559
    verify_program<test_sub>();
    verify_program<test_sub2>();
Paul's avatar
Paul committed
1560
    verify_program<test_softmax>();
Paul's avatar
Paul committed
1561
    verify_program<test_softmax2>();
Paul's avatar
Paul committed
1562
    verify_program<test_conv>();
Paul's avatar
Paul committed
1563
    verify_program<test_conv2>();
Khalique's avatar
Khalique committed
1564
    verify_program<test_group_conv>();
Paul's avatar
Paul committed
1565
    verify_program<test_conv_relu>();
Paul's avatar
Paul committed
1566
    verify_program<test_conv_relu_half>();
Paul's avatar
Paul committed
1567
    verify_program<test_add_relu>();
1568
    verify_program<test_leaky_relu>();
Khalique's avatar
Khalique committed
1569
1570
    verify_program<test_sigmoid>();
    verify_program<test_elu>();
Paul's avatar
Paul committed
1571
    verify_program<test_conv_pooling>();
1572
1573
    verify_program<test_global_avg_pooling>();
    verify_program<test_global_max_pooling>();
Paul's avatar
Paul committed
1574
    verify_program<test_gemm>();
Paul's avatar
Paul committed
1575
    verify_program<test_gemm_half>();
1576
    // verify_program<test_gemm_ld>();
1577
1578
1579
    verify_program<test_gemm_transposeb>();
    verify_program<test_gemm_transposea>();
    verify_program<test_gemm_transposeab>();
1580
1581
    verify_program<test_contiguous>();
    verify_program<test_transpose>();
1582
    verify_program<test_batchnorm_inference>();
Paul's avatar
Paul committed
1583
    verify_program<test_batchnorm_inference_2>();
Paul's avatar
Paul committed
1584
    verify_program<test_conv_bn>();
Paul's avatar
Paul committed
1585
    verify_program<test_conv_bn_relu_pooling>();
Paul's avatar
Paul committed
1586
    verify_program<test_conv_bn_relu_pooling2>();
1587
    verify_program<test_slice>();
1588
1589
    verify_program<test_gather>();
    verify_program<test_gather_neg_axis>();
1590
    verify_program<test_rnn_forward>();
1591
    verify_program<test_rnn_forward10>();
Shucai Xiao's avatar
Shucai Xiao committed
1592
    verify_program<test_rnn_reverse>();
1593
1594
1595
1596
    verify_program<test_rnn_reverse2>();
    verify_program<test_rnn_3args>();
    verify_program<test_rnn_4args>();
    verify_program<test_rnn_5args>();
Shucai Xiao's avatar
Shucai Xiao committed
1597
    verify_program<test_rnn_bidirectional>();
1598
    verify_program<test_rnn_bidirectional10>();
1599
    verify_program<test_rnn_bi_3args>();
Paul's avatar
Paul committed
1600
}