miopen.cpp 30.7 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
206
207
        return p;
    }
};

struct test_scale
{
Paul's avatar
Paul committed
208
    migraphx::program create_program() const
Khalique's avatar
Khalique committed
209
    {
Paul's avatar
Paul committed
210
211
        migraphx::program p;
        migraphx::shape s{migraphx::shape::float_type, {3}};
Khalique's avatar
Khalique committed
212
        auto x     = p.add_parameter("x", s);
Paul's avatar
Paul committed
213
214
215
        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
216
217
218
219
        return p;
    }
};

220
221
struct test_slice
{
Paul's avatar
Paul committed
222
    migraphx::program create_program() const
223
    {
Paul's avatar
Paul committed
224
225
        migraphx::program p;
        migraphx::shape s{migraphx::shape::int32_type, {2, 2, 4}};
226
        auto x      = p.add_parameter("x", s);
Paul's avatar
Paul committed
227
228
229
        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);
230
231
232
233
234

        return p;
    }
};

Paul's avatar
Paul committed
235
236
struct test_triadd
{
Paul's avatar
Paul committed
237
    migraphx::program create_program() const
Paul's avatar
Paul committed
238
    {
Paul's avatar
Paul committed
239
240
        migraphx::program p;
        migraphx::shape s{migraphx::shape::float_type, {3}};
Paul's avatar
Paul committed
241
242
243
        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
244
245
        auto sum = p.add_instruction(migraphx::op::add{}, x, y);
        p.add_instruction(migraphx::op::add{}, sum, z);
Paul's avatar
Paul committed
246
247
248
249
250
251
        return p;
    }
};

struct test_triadd2
{
Paul's avatar
Paul committed
252
    migraphx::program create_program() const
Paul's avatar
Paul committed
253
    {
Paul's avatar
Paul committed
254
255
256
        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
257
258
259
        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
260
261
262
        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
263
264
265
266
        return p;
    }
};

Paul's avatar
Paul committed
267
268
struct test_add_broadcast
{
Paul's avatar
Paul committed
269
    migraphx::program create_program() const
Paul's avatar
Paul committed
270
    {
Paul's avatar
Paul committed
271
272
273
274
275
276
        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
277
278
279
280
        return p;
    }
};

Paul's avatar
Paul committed
281
282
struct test_add_broadcast2
{
Paul's avatar
Paul committed
283
    migraphx::program create_program() const
Paul's avatar
Paul committed
284
    {
Paul's avatar
Paul committed
285
286
287
288
289
290
        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
291
292
293
294
        return p;
    }
};

Paul's avatar
Latest  
Paul committed
295
296
struct test_add_broadcast3
{
Paul's avatar
Paul committed
297
    migraphx::program create_program() const
Paul's avatar
Latest  
Paul committed
298
    {
Paul's avatar
Paul committed
299
300
301
302
303
304
        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
305
306
307
308
309
310
        return p;
    }
};

struct test_add_broadcast4
{
Paul's avatar
Paul committed
311
    migraphx::program create_program() const
Paul's avatar
Latest  
Paul committed
312
    {
Paul's avatar
Paul committed
313
314
315
316
317
318
        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
319
320
321
322
        return p;
    }
};

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

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

Paul's avatar
Paul committed
353
354
struct test_softmax
{
Paul's avatar
Paul committed
355
    migraphx::program create_program() const
Paul's avatar
Paul committed
356
    {
Paul's avatar
Paul committed
357
358
359
        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
360
361
362
363
364
365
        return p;
    }
};

struct test_softmax2
{
Paul's avatar
Paul committed
366
    migraphx::program create_program() const
Paul's avatar
Paul committed
367
    {
Paul's avatar
Paul committed
368
        migraphx::program p;
Paul's avatar
Paul committed
369
370
        auto x =
            p.add_parameter("x", migraphx::shape{migraphx::shape::float_type, {1, 1000, 1, 1}});
Paul's avatar
Paul committed
371
        p.add_instruction(migraphx::op::softmax{}, x);
Paul's avatar
Paul committed
372
373
374
375
        return p;
    }
};

Paul's avatar
Paul committed
376
377
struct test_conv
{
Paul's avatar
Paul committed
378
    migraphx::program create_program() const
Paul's avatar
Paul committed
379
    {
Paul's avatar
Paul committed
380
        migraphx::program p;
Paul's avatar
Paul committed
381
382
        auto input =
            p.add_parameter("x", migraphx::shape{migraphx::shape::float_type, {4, 3, 3, 3}});
Paul's avatar
Paul committed
383
        auto weights =
Paul's avatar
Paul committed
384
385
            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
386
387
388
389
        return p;
    }
};

Paul's avatar
Paul committed
390
391
struct test_conv2
{
Paul's avatar
Paul committed
392
    migraphx::program create_program() const
Paul's avatar
Paul committed
393
    {
Paul's avatar
Paul committed
394
        migraphx::program p;
Paul's avatar
Paul committed
395
        auto input =
Paul's avatar
Paul committed
396
            p.add_parameter("x", migraphx::shape{migraphx::shape::float_type, {1, 512, 28, 28}});
Paul's avatar
Paul committed
397
        auto weights =
Paul's avatar
Paul committed
398
399
            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
400
401
402
403
        return p;
    }
};

Paul's avatar
Paul committed
404
struct test_conv_relu
Paul's avatar
Paul committed
405
{
Paul's avatar
Paul committed
406
    migraphx::program create_program() const
Paul's avatar
Paul committed
407
    {
Paul's avatar
Paul committed
408
        migraphx::program p;
Paul's avatar
Paul committed
409
410
        auto input =
            p.add_parameter("x", migraphx::shape{migraphx::shape::float_type, {4, 3, 3, 3}});
Paul's avatar
Paul committed
411
        auto weights =
Paul's avatar
Paul committed
412
413
414
            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
415
416
417
418
        return p;
    }
};

Paul's avatar
Paul committed
419
420
struct test_conv_relu_half
{
Paul's avatar
Paul committed
421
    migraphx::program create_program() const
Paul's avatar
Paul committed
422
    {
Paul's avatar
Paul committed
423
        migraphx::program p;
Paul's avatar
Paul committed
424
425
        auto input =
            p.add_parameter("x", migraphx::shape{migraphx::shape::half_type, {4, 3, 3, 3}});
Paul's avatar
Paul committed
426
        auto weights =
Paul's avatar
Paul committed
427
428
429
            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
430
431
432
433
        return p;
    }
};

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

Khalique's avatar
Khalique committed
447
448
449
450
451
452
453
454
455
456
457
458
459
460
461
462
463
464
465
466
467
468
469
470
471
472
473
474
475
476
477
478
479
struct test_sigmoid
{
    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::sigmoid{}, x);
        return p;
    }
};

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;
    }
};

struct test_abs
{
    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::abs{}, x);
        return p;
    }
};

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

Khalique's avatar
Khalique committed
491
492
493
494
495
496
497
498
499
500
501
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
502
503
struct test_conv_pooling
{
Paul's avatar
Paul committed
504
    migraphx::program create_program() const
Paul's avatar
Paul committed
505
    {
Paul's avatar
Paul committed
506
        migraphx::program p;
Paul's avatar
Paul committed
507
        auto input =
Paul's avatar
Paul committed
508
            p.add_parameter("x", migraphx::shape{migraphx::shape::float_type, {4, 3, 32, 32}});
Paul's avatar
Paul committed
509
        auto weights =
Paul's avatar
Paul committed
510
511
512
513
            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
514
515
516
517
        return p;
    }
};

518
519
struct test_global_avg_pooling
{
Paul's avatar
Paul committed
520
    migraphx::program create_program() const
521
    {
Paul's avatar
Paul committed
522
        migraphx::program p;
523
        auto input =
Paul's avatar
Paul committed
524
525
            p.add_parameter("x", migraphx::shape{migraphx::shape::float_type, {1, 3, 16, 16}});
        auto op    = migraphx::op::pooling{"average"};
526
        auto lens  = input->get_shape().lens();
Khalique's avatar
Khalique committed
527
        op.lengths = {lens[2], lens[3]};
528
529
530
531
532
533
534
        p.add_instruction(op, input);
        return p;
    }
};

struct test_global_max_pooling
{
Paul's avatar
Paul committed
535
    migraphx::program create_program() const
536
    {
Paul's avatar
Paul committed
537
        migraphx::program p;
538
        auto input =
Paul's avatar
Paul committed
539
540
            p.add_parameter("x", migraphx::shape{migraphx::shape::float_type, {1, 3, 16, 16}});
        auto op    = migraphx::op::pooling{"max"};
541
        auto lens  = input->get_shape().lens();
Khalique's avatar
Khalique committed
542
        op.lengths = {lens[2], lens[3]};
543
544
545
546
547
        p.add_instruction(op, input);
        return p;
    }
};

Paul's avatar
Paul committed
548
549
struct test_gemm
{
Paul's avatar
Paul committed
550
    migraphx::program create_program() const
Paul's avatar
Paul committed
551
    {
Paul's avatar
Paul committed
552
553
554
555
        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
556
557
558
559
        return p;
    }
};

Paul's avatar
Paul committed
560
561
struct test_gemm_half
{
Paul's avatar
Paul committed
562
    migraphx::program create_program() const
Paul's avatar
Paul committed
563
    {
Paul's avatar
Paul committed
564
565
566
567
        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
568
569
570
571
        return p;
    }
};

Paul's avatar
Paul committed
572
573
struct test_gemm_ld
{
Paul's avatar
Paul committed
574
    migraphx::program create_program() const
Paul's avatar
Paul committed
575
    {
Paul's avatar
Paul committed
576
        migraphx::program p;
Paul's avatar
Paul committed
577
578
579
580
        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
581
        p.add_instruction(migraphx::op::dot{}, a, b);
Paul's avatar
Paul committed
582
583
584
585
        return p;
    }
};

586
587
struct test_gemm_transposeb
{
Paul's avatar
Paul committed
588
    migraphx::program create_program() const
589
    {
Paul's avatar
Paul committed
590
591
592
593
594
        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);
595
596
597
598
599
600
        return p;
    }
};

struct test_gemm_transposea
{
Paul's avatar
Paul committed
601
    migraphx::program create_program() const
602
    {
Paul's avatar
Paul committed
603
604
605
606
607
        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);
608
609
610
611
612
613
        return p;
    }
};

struct test_gemm_transposeab
{
Paul's avatar
Paul committed
614
    migraphx::program create_program() const
615
    {
Paul's avatar
Paul committed
616
617
618
619
620
621
        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);
622
623
624
625
        return p;
    }
};

626
627
struct test_contiguous
{
Paul's avatar
Paul committed
628
    migraphx::program create_program() const
629
    {
Paul's avatar
Paul committed
630
631
        migraphx::program p;
        migraphx::shape s{migraphx::shape::float_type, {4, 4, 4, 3}, {48, 4, 1, 16}};
632
        auto x = p.add_parameter("x", s);
Paul's avatar
Paul committed
633
        p.add_instruction(migraphx::op::contiguous{}, x);
Paul's avatar
Paul committed
634
        EXPECT(p.get_shape().standard());
635
636
637
638
        return p;
    }
};

639
struct test_transpose
640
{
Paul's avatar
Paul committed
641
    migraphx::program create_program() const
642
    {
Paul's avatar
Paul committed
643
644
        migraphx::program p;
        migraphx::shape s{migraphx::shape::float_type, {4, 3, 4, 4}};
645
646
        auto x                    = p.add_parameter("x", s);
        std::vector<int64_t> perm = {0, 2, 3, 1};
Paul's avatar
Paul committed
647
648
        auto l                    = p.add_instruction(migraphx::op::transpose{perm}, x);
        p.add_instruction(migraphx::op::contiguous{}, l);
649
650
651
        return p;
    }
};
652

Paul's avatar
Paul committed
653
654
655
656
657
658
659
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
660
    migraphx::program create_program() const
Paul's avatar
Paul committed
661
    {
Paul's avatar
Paul committed
662
        migraphx::program p;
Paul's avatar
Paul committed
663

Paul's avatar
Paul committed
664
665
        migraphx::shape s{migraphx::shape::float_type, {batches, channels, height, width}};
        migraphx::shape vars{migraphx::shape::float_type, {channels}};
Paul's avatar
Paul committed
666
        auto x        = p.add_parameter("x", s);
Paul's avatar
Paul committed
667
668
669
670
671
        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
672
673
674
675
        return p;
    }
};

wsttiger's avatar
wsttiger committed
676
677
678
679
680
681
682
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
683
    migraphx::program create_program() const
wsttiger's avatar
wsttiger committed
684
    {
Paul's avatar
Paul committed
685
        migraphx::program p;
wsttiger's avatar
wsttiger committed
686

Paul's avatar
Paul committed
687
688
        migraphx::shape s{migraphx::shape::float_type, {batches, channels, height, width}};
        migraphx::shape vars{migraphx::shape::float_type, {channels}};
wsttiger's avatar
wsttiger committed
689
        auto x        = p.add_parameter("x", s);
Paul's avatar
Paul committed
690
691
692
693
694
        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
695
696
697
698
        return p;
    }
};

Paul's avatar
Paul committed
699
700
struct test_conv_bn
{
Paul's avatar
Paul committed
701
    migraphx::program create_program() const
Paul's avatar
Paul committed
702
    {
Paul's avatar
Paul committed
703
        migraphx::program p;
Paul's avatar
Paul committed
704

Paul's avatar
Paul committed
705
706
707
        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
708
709
        auto x        = p.add_parameter("x", xs);
        auto w        = p.add_parameter("w", ws);
Paul's avatar
Paul committed
710
711
712
713
714
715
        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
716
717
718
719
        return p;
    }
};

Paul's avatar
Paul committed
720
721
struct test_conv_bn_relu_pooling
{
Paul's avatar
Paul committed
722
    migraphx::program create_program() const
Paul's avatar
Paul committed
723
    {
Paul's avatar
Paul committed
724
        migraphx::program p;
Paul's avatar
Paul committed
725

Paul's avatar
Paul committed
726
727
728
        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
729
730
        auto x        = p.add_parameter("x", xs);
        auto w        = p.add_parameter("w", ws);
Paul's avatar
Paul committed
731
732
733
734
735
        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
736
        auto bn       = p.add_instruction(
Paul's avatar
Paul committed
737
738
739
            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
740
741
742
743
        return p;
    }
};

744
745
struct test_concat
{
Paul's avatar
Paul committed
746
    migraphx::program create_program() const
747
    {
Paul's avatar
Paul committed
748
        migraphx::program p;
wsttiger's avatar
wsttiger committed
749
        std::size_t axis = 1;
Paul's avatar
Paul committed
750
751
752
        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}};
753
754
755
        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
756
        p.add_instruction(migraphx::op::concat{axis}, l0, l1, l2);
757
758
759
760
761
762
        return p;
    }
};

struct test_concat2
{
Paul's avatar
Paul committed
763
    migraphx::program create_program() const
764
    {
Paul's avatar
Paul committed
765
        migraphx::program p;
wsttiger's avatar
wsttiger committed
766
        std::size_t axis = 0;
Paul's avatar
Paul committed
767
768
769
        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}};
770
771
772
        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
773
        p.add_instruction(migraphx::op::concat{axis}, l0, l1, l2);
774
775
776
777
        return p;
    }
};

wsttiger's avatar
wsttiger committed
778
779
struct test_concat_relu
{
Paul's avatar
Paul committed
780
    migraphx::program create_program() const
wsttiger's avatar
wsttiger committed
781
    {
Paul's avatar
Paul committed
782
        migraphx::program p;
wsttiger's avatar
wsttiger committed
783
        std::size_t axis = 0;
Paul's avatar
Paul committed
784
785
786
        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
787
788
789
        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
790
791
792
793
794
        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
795
796
797
798
799
800
        return p;
    }
};

void manual_identity()
{
Paul's avatar
Paul committed
801
    migraphx::program p;
wsttiger's avatar
wsttiger committed
802
    std::vector<float> data0 = {0, 1, 2, 3};
Paul's avatar
Paul committed
803
804
805
806
807
    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
808
809
    for(auto&& x : p.get_parameter_shapes())
    {
Paul's avatar
Paul committed
810
        m[x.first] = migraphx::gpu::to_gpu(migraphx::generate_argument(x.second));
wsttiger's avatar
wsttiger committed
811
    }
Paul's avatar
Paul committed
812
    auto result = migraphx::gpu::from_gpu(p.eval(m));
wsttiger's avatar
wsttiger committed
813
814
815
816
817
    std::cout << result << std::endl;
}

void manual_test_concat_relu()
{
Paul's avatar
Paul committed
818
    migraphx::program p;
wsttiger's avatar
wsttiger committed
819
    std::size_t axis         = 0;
wsttiger's avatar
wsttiger committed
820
821
822
    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
823
824
825
826
827
828
829
830
831
832
833
834
835
836
    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
837
838
    for(auto&& x : p.get_parameter_shapes())
    {
Paul's avatar
Paul committed
839
        m[x.first] = migraphx::gpu::to_gpu(migraphx::generate_argument(x.second));
wsttiger's avatar
wsttiger committed
840
    }
Paul's avatar
Paul committed
841
    auto result = migraphx::gpu::from_gpu(p.eval(m));
wsttiger's avatar
wsttiger committed
842
843
844
    std::cout << result << std::endl;
}

Paul's avatar
Paul committed
845
846
struct test_conv_bn_relu_pooling2
{
Paul's avatar
Paul committed
847
848
    static migraphx::instruction_ref
    add_bn(migraphx::program& p, migraphx::instruction_ref x, std::size_t channels)
Paul's avatar
Paul committed
849
    {
Paul's avatar
Paul committed
850
        migraphx::shape vars{migraphx::shape::float_type, {channels}};
Paul's avatar
Paul committed
851
852
853
854
855
        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
856
        return p.add_instruction(
Paul's avatar
Paul committed
857
            migraphx::op::batch_norm_inference{}, x, scale, bias, mean, variance);
Paul's avatar
Paul committed
858
    }
Paul's avatar
Paul committed
859
    migraphx::program create_program() const
Paul's avatar
Paul committed
860
    {
Paul's avatar
Paul committed
861
        migraphx::program p;
Paul's avatar
Paul committed
862

Paul's avatar
Paul committed
863
864
865
866
        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
867
868
        auto x1    = p.add_parameter("x1", xs1);
        auto w1    = p.add_parameter("w1", ws1);
Paul's avatar
Paul committed
869
        auto conv1 = p.add_instruction(migraphx::op::convolution{{0, 0}, {1, 1}, {1, 1}}, x1, w1);
Paul's avatar
Paul committed
870
871
872
        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
873
        auto conv2 = p.add_instruction(migraphx::op::convolution{{0, 0}, {2, 2}, {1, 1}}, x2, w2);
Paul's avatar
Paul committed
874
        auto bn2   = add_bn(p, conv2, 2048);
Paul's avatar
Paul committed
875
876
877
        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
878
879
880
881
        return p;
    }
};

Paul's avatar
Paul committed
882
883
int main()
{
Khalique's avatar
Khalique committed
884
    verify_program<test_abs>();
885
886
    verify_program<test_concat>();
    verify_program<test_concat2>();
wsttiger's avatar
wsttiger committed
887
    verify_program<test_concat_relu>();
Paul's avatar
Paul committed
888
    verify_program<test_add>();
Paul's avatar
Paul committed
889
    verify_program<test_add_half>();
Khalique's avatar
Khalique committed
890
891
    verify_program<test_mul>();
    verify_program<test_scale>();
Paul's avatar
Paul committed
892
893
    verify_program<test_triadd>();
    verify_program<test_triadd2>();
Paul's avatar
Paul committed
894
    verify_program<test_add_broadcast>();
Paul's avatar
Paul committed
895
    verify_program<test_add_broadcast2>();
Paul's avatar
Latest  
Paul committed
896
897
    verify_program<test_add_broadcast3>();
    verify_program<test_add_broadcast4>();
Paul's avatar
Paul committed
898
    verify_program<test_add_broadcast5>();
Paul's avatar
Paul committed
899
    verify_program<test_triadd_broadcast>();
Paul's avatar
Paul committed
900
    verify_program<test_softmax>();
Paul's avatar
Paul committed
901
    verify_program<test_softmax2>();
Paul's avatar
Paul committed
902
    verify_program<test_conv>();
Paul's avatar
Paul committed
903
    verify_program<test_conv2>();
Paul's avatar
Paul committed
904
    verify_program<test_conv_relu>();
Paul's avatar
Paul committed
905
    verify_program<test_conv_relu_half>();
Paul's avatar
Paul committed
906
    verify_program<test_add_relu>();
907
    verify_program<test_leaky_relu>();
Khalique's avatar
Khalique committed
908
909
910
    verify_program<test_sigmoid>();
    verify_program<test_tanh>();
    verify_program<test_elu>();
Paul's avatar
Paul committed
911
    verify_program<test_conv_pooling>();
912
913
    verify_program<test_global_avg_pooling>();
    verify_program<test_global_max_pooling>();
Paul's avatar
Paul committed
914
    verify_program<test_gemm>();
Paul's avatar
Paul committed
915
    verify_program<test_gemm_half>();
916
    // verify_program<test_gemm_ld>();
917
918
919
    verify_program<test_gemm_transposeb>();
    verify_program<test_gemm_transposea>();
    verify_program<test_gemm_transposeab>();
920
921
    verify_program<test_contiguous>();
    verify_program<test_transpose>();
922
    verify_program<test_batchnorm_inference>();
Paul's avatar
Paul committed
923
    verify_program<test_batchnorm_inference_2>();
Paul's avatar
Paul committed
924
    verify_program<test_conv_bn>();
Paul's avatar
Paul committed
925
    verify_program<test_conv_bn_relu_pooling>();
Paul's avatar
Paul committed
926
    verify_program<test_conv_bn_relu_pooling2>();
927
    verify_program<test_slice>();
Paul's avatar
Paul committed
928
}