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

Paul's avatar
Paul committed
2
3
4
5
6
7
8
9
10
11
12
#include <migraphx/program.hpp>
#include <migraphx/operators.hpp>
#include <migraphx/generate.hpp>
#include <migraphx/cpu/target.hpp>
#include <migraphx/gpu/target.hpp>
#include <migraphx/gpu/miopen.hpp>
#include <migraphx/gpu/hip.hpp>
#include <migraphx/manage_ptr.hpp>
#include <migraphx/type_name.hpp>
#include <migraphx/verify_args.hpp>
#include <migraphx/instruction.hpp>
Paul's avatar
Paul committed
13
14
15

#include <miopen/miopen.h>

Paul's avatar
Paul committed
16
17
18
#include <future>
#include <thread>

Paul's avatar
Paul committed
19
#include <test.hpp>
Paul's avatar
Paul committed
20

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
template <class V>
Paul's avatar
Paul committed
130
void run_verify_program()
Paul's avatar
Paul committed
131
{
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
154
155
156
157
158
159
160
161
162
163
164
165
166
167
168
169
170
171
172
173
template<class T>
int auto_register_verify_program()
{
    test::add_test_case(migraphx::get_type_name<T>(), [] { run_verify_program<T>(); });
    return 0;
}

template<class T>
struct verify_program
{
    static int static_register;
    // This typedef ensures that the static member will be instantiated if
    // the class itself is instantiated
    using static_register_type = std::integral_constant<decltype(&static_register), &static_register>;
};

template<class T>
int verify_program<T>::static_register = auto_register_verify_program<T>();



struct test_literals : verify_program<test_literals>
Paul's avatar
Paul committed
174
{
Paul's avatar
Paul committed
175
    migraphx::program create_program() const
Paul's avatar
Paul committed
176
    {
Paul's avatar
Paul committed
177
        migraphx::program p;
Paul's avatar
Paul committed
178
        auto input = p.add_literal(
Paul's avatar
Paul committed
179
            generate_literal(migraphx::shape{migraphx::shape::float_type, {4, 3, 3, 3}}));
Paul's avatar
Paul committed
180
        auto weights = p.add_literal(
Paul's avatar
Paul committed
181
182
183
            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
184
185
186
187
        return p;
    }
};

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

Paul's avatar
Paul committed
201
struct test_add_half : verify_program<test_add_half>
Paul's avatar
Paul committed
202
{
Paul's avatar
Paul committed
203
    migraphx::program create_program() const
Paul's avatar
Paul committed
204
    {
Paul's avatar
Paul committed
205
206
        migraphx::program p;
        migraphx::shape s{migraphx::shape::half_type, {3}};
Paul's avatar
Paul committed
207
208
        auto x = p.add_parameter("x", s);
        auto y = p.add_parameter("y", s);
Paul's avatar
Paul committed
209
        p.add_instruction(migraphx::op::add{}, x, y);
Paul's avatar
Paul committed
210
211
212
213
        return p;
    }
};

Paul's avatar
Paul committed
214
struct test_mul : verify_program<test_mul>
Khalique's avatar
Khalique committed
215
{
Paul's avatar
Paul committed
216
    migraphx::program create_program() const
Khalique's avatar
Khalique committed
217
    {
Paul's avatar
Paul committed
218
219
        migraphx::program p;
        migraphx::shape s{migraphx::shape::float_type, {3}};
Khalique's avatar
Khalique committed
220
221
        auto x = p.add_parameter("x", s);
        auto y = p.add_parameter("y", s);
Paul's avatar
Paul committed
222
        p.add_instruction(migraphx::op::mul{}, x, y);
Khalique's avatar
Khalique committed
223
224
225
226
        return p;
    }
};

Paul's avatar
Paul committed
227
struct test_sin : verify_program<test_sin>
228
229
230
231
232
233
234
235
236
237
238
{
    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;
    }
};

Paul's avatar
Paul committed
239
struct test_scale : verify_program<test_scale>
Khalique's avatar
Khalique committed
240
{
Paul's avatar
Paul committed
241
    migraphx::program create_program() const
Khalique's avatar
Khalique committed
242
    {
Paul's avatar
Paul committed
243
244
        migraphx::program p;
        migraphx::shape s{migraphx::shape::float_type, {3}};
Khalique's avatar
Khalique committed
245
        auto x     = p.add_parameter("x", s);
Paul's avatar
Paul committed
246
247
248
        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
249
250
251
252
        return p;
    }
};

Paul's avatar
Paul committed
253
struct test_slice : verify_program<test_slice>
254
{
Paul's avatar
Paul committed
255
    migraphx::program create_program() const
256
    {
Paul's avatar
Paul committed
257
258
        migraphx::program p;
        migraphx::shape s{migraphx::shape::int32_type, {2, 2, 4}};
259
        auto x      = p.add_parameter("x", s);
Paul's avatar
Paul committed
260
261
262
        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);
263
264
265
266
267

        return p;
    }
};

Paul's avatar
Paul committed
268
struct test_triadd : verify_program<test_triadd>
Paul's avatar
Paul committed
269
{
Paul's avatar
Paul committed
270
    migraphx::program create_program() const
Paul's avatar
Paul committed
271
    {
Paul's avatar
Paul committed
272
273
        migraphx::program p;
        migraphx::shape s{migraphx::shape::float_type, {3}};
Paul's avatar
Paul committed
274
275
276
        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
277
278
        auto sum = p.add_instruction(migraphx::op::add{}, x, y);
        p.add_instruction(migraphx::op::add{}, sum, z);
Paul's avatar
Paul committed
279
280
281
282
        return p;
    }
};

Paul's avatar
Paul committed
283
struct test_triadd2 : verify_program<test_triadd2>
Paul's avatar
Paul committed
284
{
Paul's avatar
Paul committed
285
    migraphx::program create_program() const
Paul's avatar
Paul committed
286
    {
Paul's avatar
Paul committed
287
288
289
        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
290
291
292
        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
293
294
295
        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
296
297
298
299
        return p;
    }
};

Paul's avatar
Paul committed
300
struct test_add_broadcast : verify_program<test_add_broadcast>
Paul's avatar
Paul committed
301
{
Paul's avatar
Paul committed
302
    migraphx::program create_program() const
Paul's avatar
Paul committed
303
    {
Paul's avatar
Paul committed
304
305
306
307
308
309
        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
310
311
312
313
        return p;
    }
};

Paul's avatar
Paul committed
314
struct test_add_broadcast2 : verify_program<test_add_broadcast2>
Paul's avatar
Paul committed
315
{
Paul's avatar
Paul committed
316
    migraphx::program create_program() const
Paul's avatar
Paul committed
317
    {
Paul's avatar
Paul committed
318
319
320
321
322
323
        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
324
325
326
327
        return p;
    }
};

Paul's avatar
Paul committed
328
struct test_add_broadcast3 : verify_program<test_add_broadcast3>
Paul's avatar
Latest  
Paul committed
329
{
Paul's avatar
Paul committed
330
    migraphx::program create_program() const
Paul's avatar
Latest  
Paul committed
331
    {
Paul's avatar
Paul committed
332
333
334
335
336
337
        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
338
339
340
341
        return p;
    }
};

Paul's avatar
Paul committed
342
struct test_add_broadcast4 : verify_program<test_add_broadcast4>
Paul's avatar
Latest  
Paul committed
343
{
Paul's avatar
Paul committed
344
    migraphx::program create_program() const
Paul's avatar
Latest  
Paul committed
345
    {
Paul's avatar
Paul committed
346
347
348
349
350
351
        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
352
353
354
355
        return p;
    }
};

Paul's avatar
Paul committed
356
struct test_add_broadcast5 : verify_program<test_add_broadcast5>
Paul's avatar
Paul committed
357
{
Paul's avatar
Paul committed
358
    migraphx::program create_program() const
Paul's avatar
Paul committed
359
    {
Paul's avatar
Paul committed
360
361
362
363
364
365
        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
366
367
368
369
        return p;
    }
};

Paul's avatar
Paul committed
370
struct test_triadd_broadcast : verify_program<test_triadd_broadcast>
Paul's avatar
Paul committed
371
{
Paul's avatar
Paul committed
372
    migraphx::program create_program() const
Paul's avatar
Paul committed
373
    {
Paul's avatar
Paul committed
374
375
376
377
378
379
380
381
        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
382
383
384
385
        return p;
    }
};

Paul's avatar
Paul committed
386
struct test_softmax : verify_program<test_softmax>
Paul's avatar
Paul committed
387
{
Paul's avatar
Paul committed
388
    migraphx::program create_program() const
Paul's avatar
Paul committed
389
    {
Paul's avatar
Paul committed
390
391
392
        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
393
394
395
396
        return p;
    }
};

Paul's avatar
Paul committed
397
struct test_softmax2 : verify_program<test_softmax2>
Paul's avatar
Paul committed
398
{
Paul's avatar
Paul committed
399
    migraphx::program create_program() const
Paul's avatar
Paul committed
400
    {
Paul's avatar
Paul committed
401
        migraphx::program p;
Paul's avatar
Paul committed
402
403
        auto x =
            p.add_parameter("x", migraphx::shape{migraphx::shape::float_type, {1, 1000, 1, 1}});
Paul's avatar
Paul committed
404
        p.add_instruction(migraphx::op::softmax{}, x);
Paul's avatar
Paul committed
405
406
407
408
        return p;
    }
};

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

Paul's avatar
Paul committed
423
struct test_conv2 : verify_program<test_conv2>
Paul's avatar
Paul committed
424
{
Paul's avatar
Paul committed
425
    migraphx::program create_program() const
Paul's avatar
Paul committed
426
    {
Paul's avatar
Paul committed
427
        migraphx::program p;
Paul's avatar
Paul committed
428
        auto input =
Paul's avatar
Paul committed
429
            p.add_parameter("x", migraphx::shape{migraphx::shape::float_type, {1, 512, 28, 28}});
Paul's avatar
Paul committed
430
        auto weights =
Paul's avatar
Paul committed
431
432
            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
433
434
435
436
        return p;
    }
};

Paul's avatar
Paul committed
437
struct test_conv_relu : verify_program<test_conv_relu>
Paul's avatar
Paul committed
438
{
Paul's avatar
Paul committed
439
    migraphx::program create_program() const
Paul's avatar
Paul committed
440
    {
Paul's avatar
Paul committed
441
        migraphx::program p;
Paul's avatar
Paul committed
442
443
        auto input =
            p.add_parameter("x", migraphx::shape{migraphx::shape::float_type, {4, 3, 3, 3}});
Paul's avatar
Paul committed
444
        auto weights =
Paul's avatar
Paul committed
445
446
447
            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
448
449
450
451
        return p;
    }
};

Paul's avatar
Paul committed
452
struct test_conv_relu_half : verify_program<test_conv_relu_half>
Paul's avatar
Paul committed
453
{
Paul's avatar
Paul committed
454
    migraphx::program create_program() const
Paul's avatar
Paul committed
455
    {
Paul's avatar
Paul committed
456
        migraphx::program p;
Paul's avatar
Paul committed
457
458
        auto input =
            p.add_parameter("x", migraphx::shape{migraphx::shape::half_type, {4, 3, 3, 3}});
Paul's avatar
Paul committed
459
        auto weights =
Paul's avatar
Paul committed
460
461
462
            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
463
464
465
466
        return p;
    }
};

Paul's avatar
Paul committed
467
struct test_add_relu : verify_program<test_add_relu>
Paul's avatar
Paul committed
468
{
Paul's avatar
Paul committed
469
    migraphx::program create_program() const
Paul's avatar
Paul committed
470
    {
Paul's avatar
Paul committed
471
472
473
474
475
        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
476
477
478
479
        return p;
    }
};

Paul's avatar
Paul committed
480
struct test_leaky_relu : verify_program<test_leaky_relu>
481
{
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;
    }
};

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

Paul's avatar
Paul committed
507
struct test_global_avg_pooling : verify_program<test_global_avg_pooling>
508
{
Paul's avatar
Paul committed
509
    migraphx::program create_program() const
510
    {
Paul's avatar
Paul committed
511
        migraphx::program p;
512
        auto input =
Paul's avatar
Paul committed
513
514
            p.add_parameter("x", migraphx::shape{migraphx::shape::float_type, {1, 3, 16, 16}});
        auto op    = migraphx::op::pooling{"average"};
515
        auto lens  = input->get_shape().lens();
Khalique's avatar
Khalique committed
516
        op.lengths = {lens[2], lens[3]};
517
518
519
520
521
        p.add_instruction(op, input);
        return p;
    }
};

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

Paul's avatar
Paul committed
537
struct test_gemm : verify_program<test_gemm>
Paul's avatar
Paul committed
538
{
Paul's avatar
Paul committed
539
    migraphx::program create_program() const
Paul's avatar
Paul committed
540
    {
Paul's avatar
Paul committed
541
542
543
544
        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
545
546
547
548
        return p;
    }
};

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

Paul's avatar
Paul committed
561
struct test_gemm_ld // : verify_program<test_gemm_ld>
Paul's avatar
Paul committed
562
{
Paul's avatar
Paul committed
563
    migraphx::program create_program() const
Paul's avatar
Paul committed
564
    {
Paul's avatar
Paul committed
565
        migraphx::program p;
Paul's avatar
Paul committed
566
567
568
569
        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
570
        p.add_instruction(migraphx::op::dot{}, a, b);
Paul's avatar
Paul committed
571
572
573
574
        return p;
    }
};

Paul's avatar
Paul committed
575
struct test_gemm_transposeb : verify_program<test_gemm_transposeb>
576
{
Paul's avatar
Paul committed
577
    migraphx::program create_program() const
578
    {
Paul's avatar
Paul committed
579
580
581
582
583
        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);
584
585
586
587
        return p;
    }
};

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

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

Paul's avatar
Paul committed
615
struct test_contiguous : verify_program<test_contiguous>
616
{
Paul's avatar
Paul committed
617
    migraphx::program create_program() const
618
    {
Paul's avatar
Paul committed
619
620
        migraphx::program p;
        migraphx::shape s{migraphx::shape::float_type, {4, 4, 4, 3}, {48, 4, 1, 16}};
621
        auto x = p.add_parameter("x", s);
Paul's avatar
Paul committed
622
        p.add_instruction(migraphx::op::contiguous{}, x);
Paul's avatar
Paul committed
623
        EXPECT(p.get_shape().standard());
624
625
626
627
        return p;
    }
};

Paul's avatar
Paul committed
628
struct test_transpose : verify_program<test_transpose>
629
{
Paul's avatar
Paul committed
630
    migraphx::program create_program() const
631
    {
Paul's avatar
Paul committed
632
633
        migraphx::program p;
        migraphx::shape s{migraphx::shape::float_type, {4, 3, 4, 4}};
634
635
        auto x                    = p.add_parameter("x", s);
        std::vector<int64_t> perm = {0, 2, 3, 1};
Paul's avatar
Paul committed
636
637
        auto l                    = p.add_instruction(migraphx::op::transpose{perm}, x);
        p.add_instruction(migraphx::op::contiguous{}, l);
638
639
640
        return p;
    }
};
641

Paul's avatar
Paul committed
642
struct test_batchnorm_inference_2 : verify_program<test_batchnorm_inference_2>
Paul's avatar
Paul committed
643
644
645
646
647
648
{
    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
649
    migraphx::program create_program() const
Paul's avatar
Paul committed
650
    {
Paul's avatar
Paul committed
651
        migraphx::program p;
Paul's avatar
Paul committed
652

Paul's avatar
Paul committed
653
654
        migraphx::shape s{migraphx::shape::float_type, {batches, channels, height, width}};
        migraphx::shape vars{migraphx::shape::float_type, {channels}};
Paul's avatar
Paul committed
655
        auto x        = p.add_parameter("x", s);
Paul's avatar
Paul committed
656
657
658
659
660
        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
661
662
663
664
        return p;
    }
};

Paul's avatar
Paul committed
665
struct test_batchnorm_inference : verify_program<test_batchnorm_inference>
wsttiger's avatar
wsttiger committed
666
667
668
669
670
671
{
    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
672
    migraphx::program create_program() const
wsttiger's avatar
wsttiger committed
673
    {
Paul's avatar
Paul committed
674
        migraphx::program p;
wsttiger's avatar
wsttiger committed
675

Paul's avatar
Paul committed
676
677
        migraphx::shape s{migraphx::shape::float_type, {batches, channels, height, width}};
        migraphx::shape vars{migraphx::shape::float_type, {channels}};
wsttiger's avatar
wsttiger committed
678
        auto x        = p.add_parameter("x", s);
Paul's avatar
Paul committed
679
680
681
682
683
        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
684
685
686
687
        return p;
    }
};

Paul's avatar
Paul committed
688
struct test_conv_bn : verify_program<test_conv_bn>
Paul's avatar
Paul committed
689
{
Paul's avatar
Paul committed
690
    migraphx::program create_program() const
Paul's avatar
Paul committed
691
    {
Paul's avatar
Paul committed
692
        migraphx::program p;
Paul's avatar
Paul committed
693

Paul's avatar
Paul committed
694
695
696
        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
697
698
        auto x        = p.add_parameter("x", xs);
        auto w        = p.add_parameter("w", ws);
Paul's avatar
Paul committed
699
700
701
702
703
704
        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
705
706
707
708
        return p;
    }
};

Paul's avatar
Paul committed
709
struct test_conv_bn_relu_pooling : verify_program<test_conv_bn_relu_pooling>
Paul's avatar
Paul committed
710
{
Paul's avatar
Paul committed
711
    migraphx::program create_program() const
Paul's avatar
Paul committed
712
    {
Paul's avatar
Paul committed
713
        migraphx::program p;
Paul's avatar
Paul committed
714

Paul's avatar
Paul committed
715
716
717
        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
718
719
        auto x        = p.add_parameter("x", xs);
        auto w        = p.add_parameter("w", ws);
Paul's avatar
Paul committed
720
721
722
723
724
        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
725
        auto bn       = p.add_instruction(
Paul's avatar
Paul committed
726
727
728
            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
729
730
731
732
        return p;
    }
};

Paul's avatar
Paul committed
733
struct test_concat : verify_program<test_concat>
734
{
Paul's avatar
Paul committed
735
    migraphx::program create_program() const
736
    {
Paul's avatar
Paul committed
737
        migraphx::program p;
wsttiger's avatar
wsttiger committed
738
        std::size_t axis = 1;
Paul's avatar
Paul committed
739
740
741
        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}};
742
743
744
        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
745
        p.add_instruction(migraphx::op::concat{axis}, l0, l1, l2);
746
747
748
749
        return p;
    }
};

Paul's avatar
Paul committed
750
struct test_concat2 : verify_program<test_concat2>
751
{
Paul's avatar
Paul committed
752
    migraphx::program create_program() const
753
    {
Paul's avatar
Paul committed
754
        migraphx::program p;
wsttiger's avatar
wsttiger committed
755
        std::size_t axis = 0;
Paul's avatar
Paul committed
756
757
758
        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}};
759
760
761
        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
762
        p.add_instruction(migraphx::op::concat{axis}, l0, l1, l2);
763
764
765
766
        return p;
    }
};

Paul's avatar
Paul committed
767
struct test_concat_relu : verify_program<test_concat_relu>
wsttiger's avatar
wsttiger committed
768
{
Paul's avatar
Paul committed
769
    migraphx::program create_program() const
wsttiger's avatar
wsttiger committed
770
    {
Paul's avatar
Paul committed
771
        migraphx::program p;
wsttiger's avatar
wsttiger committed
772
        std::size_t axis = 0;
Paul's avatar
Paul committed
773
774
775
        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
776
777
778
        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
779
780
781
782
783
        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
784
785
786
787
788
789
        return p;
    }
};

void manual_identity()
{
Paul's avatar
Paul committed
790
    migraphx::program p;
wsttiger's avatar
wsttiger committed
791
    std::vector<float> data0 = {0, 1, 2, 3};
Paul's avatar
Paul committed
792
793
794
795
796
    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
797
798
    for(auto&& x : p.get_parameter_shapes())
    {
Paul's avatar
Paul committed
799
        m[x.first] = migraphx::gpu::to_gpu(migraphx::generate_argument(x.second));
wsttiger's avatar
wsttiger committed
800
    }
Paul's avatar
Paul committed
801
    auto result = migraphx::gpu::from_gpu(p.eval(m));
wsttiger's avatar
wsttiger committed
802
803
804
805
806
    std::cout << result << std::endl;
}

void manual_test_concat_relu()
{
Paul's avatar
Paul committed
807
    migraphx::program p;
wsttiger's avatar
wsttiger committed
808
    std::size_t axis         = 0;
wsttiger's avatar
wsttiger committed
809
810
811
    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
812
813
814
815
816
817
818
819
820
821
822
823
824
825
    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
826
827
    for(auto&& x : p.get_parameter_shapes())
    {
Paul's avatar
Paul committed
828
        m[x.first] = migraphx::gpu::to_gpu(migraphx::generate_argument(x.second));
wsttiger's avatar
wsttiger committed
829
    }
Paul's avatar
Paul committed
830
    auto result = migraphx::gpu::from_gpu(p.eval(m));
wsttiger's avatar
wsttiger committed
831
832
833
    std::cout << result << std::endl;
}

Paul's avatar
Paul committed
834
struct test_conv_bn_relu_pooling2 : verify_program<test_conv_bn_relu_pooling2>
Paul's avatar
Paul committed
835
{
Paul's avatar
Paul committed
836
837
    static migraphx::instruction_ref
    add_bn(migraphx::program& p, migraphx::instruction_ref x, std::size_t channels)
Paul's avatar
Paul committed
838
    {
Paul's avatar
Paul committed
839
        migraphx::shape vars{migraphx::shape::float_type, {channels}};
Paul's avatar
Paul committed
840
841
842
843
844
        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
845
        return p.add_instruction(
Paul's avatar
Paul committed
846
            migraphx::op::batch_norm_inference{}, x, scale, bias, mean, variance);
Paul's avatar
Paul committed
847
    }
Paul's avatar
Paul committed
848
    migraphx::program create_program() const
Paul's avatar
Paul committed
849
    {
Paul's avatar
Paul committed
850
        migraphx::program p;
Paul's avatar
Paul committed
851

Paul's avatar
Paul committed
852
853
854
855
        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
856
857
        auto x1    = p.add_parameter("x1", xs1);
        auto w1    = p.add_parameter("w1", ws1);
Paul's avatar
Paul committed
858
        auto conv1 = p.add_instruction(migraphx::op::convolution{{0, 0}, {1, 1}, {1, 1}}, x1, w1);
Paul's avatar
Paul committed
859
860
861
        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
862
        auto conv2 = p.add_instruction(migraphx::op::convolution{{0, 0}, {2, 2}, {1, 1}}, x2, w2);
Paul's avatar
Paul committed
863
        auto bn2   = add_bn(p, conv2, 2048);
Paul's avatar
Paul committed
864
865
866
        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
867
868
869
870
        return p;
    }
};

Paul's avatar
Paul committed
871
int main(int argc, const char* argv[]) { test::run(argc, argv); }