"src/include/vscode:/vscode.git/clone" did not exist on "d1481b1334bf34834cf16dbeed477945ab3f70ae"
miopen.cpp 14.6 KB
Newer Older
Paul's avatar
Paul committed
1

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

#include <miopen/miopen.h>

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

Paul's avatar
Paul committed
18
19
#include "test.hpp"

Paul's avatar
Paul committed
20
21
22
23
#ifdef __clang__
#pragma clang diagnostic push
#pragma clang diagnostic ignored "-Wglobal-constructors"
#endif
Paul's avatar
Paul committed
24

Paul's avatar
Paul committed
25
26
// An improved async, that doesn't block
template <class Function>
Paul's avatar
Paul committed
27
28
std::future<typename std::result_of<Function()>::type> detach_async(Function&& f,
                                                                    bool parallel = true)
Paul's avatar
Paul committed
29
{
Paul's avatar
Paul committed
30
31
32
33
34
35
36
37
38
39
    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);
    }
    else
    {
Paul's avatar
Paul committed
40
        return std::async(std::launch::deferred, std::forward<Function>(f));
Paul's avatar
Paul committed
41
    }
Paul's avatar
Paul committed
42
43
}

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

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

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

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

Paul's avatar
Paul committed
108
template <class V>
Paul's avatar
Paul committed
109
migraph::argument run_gpu()
Paul's avatar
Paul committed
110
{
Paul's avatar
Paul committed
111
112
    V v;
    auto p = v.create_program();
Paul's avatar
Paul committed
113
    auto_print pp{p, 1};
Paul's avatar
Paul committed
114
    compile_check(p, migraph::gpu::target{});
Paul's avatar
Paul committed
115

Paul's avatar
Paul committed
116
    migraph::program::parameter_map m;
Paul's avatar
Paul committed
117
    for(auto&& x : p.get_parameter_shapes())
Paul's avatar
Paul committed
118
    {
Paul's avatar
Paul committed
119
        m[x.first] = migraph::gpu::to_gpu(migraph::generate_argument(x.second));
Paul's avatar
Paul committed
120
121
    }

Paul's avatar
Paul committed
122
    return migraph::gpu::from_gpu(p.eval(m));
Paul's avatar
Paul committed
123
124
}

Paul's avatar
Paul committed
125
126
127
void verify_args(const std::string& name,
                 const migraph::argument& cpu_arg,
                 const migraph::argument& gpu_arg)
Paul's avatar
Paul committed
128
{
Paul's avatar
Paul committed
129
    visit_all(cpu_arg, gpu_arg)([&](auto cpu, auto gpu) {
Paul's avatar
Paul committed
130
        if(not migraph::verify_range(cpu, gpu))
Paul's avatar
Paul committed
131
        {
132
            // TODO: Check for nans
Paul's avatar
Paul committed
133
            std::cout << "FAILED: " << name << std::endl;
Paul's avatar
Paul committed
134
135
136
137
138
139
140
141
142
143
144
145
146
147
148
149
150
151
152
153
154
155
156
            // std::cout << cpu << std::endl;
            // std::cout << gpu << std::endl;
            if(migraph::range_zero(cpu))
                std::cout << "Cpu data is all zeros" << std::endl;
            if(migraph::range_zero(gpu))
                std::cout << "Gpu data is all zeros" << std::endl;

            auto idx = migraph::mismatch_idx(cpu, gpu, migraph::float_equal);
            if(idx < migraph::range_distance(cpu))
            {
                std::cout << "Mismatch at " << idx << ": " << cpu[idx] << " != " << gpu[idx]
                          << std::endl;
            }

            auto cpu_nan_idx = find_idx(cpu, migraph::not_finite);
            if(cpu_nan_idx >= 0)
                std::cout << "Non finite number found in cpu at " << cpu_nan_idx << ": "
                          << cpu[cpu_nan_idx] << std::endl;

            auto gpu_nan_idx = find_idx(gpu, migraph::not_finite);
            if(gpu_nan_idx >= 0)
                std::cout << "Non finite number found in gpu at " << gpu_nan_idx << ": "
                          << gpu[gpu_nan_idx] << std::endl;
Paul's avatar
Paul committed
157
158
        }
    });
Paul's avatar
Paul committed
159
160
161
162
163
164
165
166
167
}

template <class V>
void verify_program()
{
    auto_print::set_terminate_handler(migraph::get_type_name<V>());
    auto cpu_arg_f = detach_async([] { return run_cpu<V>(); });
    auto gpu_arg   = run_gpu<V>();
    verify_args(migraph::get_type_name<V>(), cpu_arg_f.get(), gpu_arg);
Paul's avatar
Paul committed
168
    std::set_terminate(nullptr);
Paul's avatar
Paul committed
169
170
}

Paul's avatar
Paul committed
171
172
173
174
175
struct test_literals
{
    migraph::program create_program() const
    {
        migraph::program p;
Paul's avatar
Paul committed
176
177
178
179
        auto input = p.add_literal(
            generate_literal(migraph::shape{migraph::shape::float_type, {4, 3, 3, 3}}));
        auto weights = p.add_literal(
            generate_literal(migraph::shape{migraph::shape::float_type, {4, 3, 3, 3}}));
Paul's avatar
Paul committed
180
181
182
183
184
185
        auto conv = p.add_instruction(migraph::convolution{}, input, weights);
        p.add_instruction(migraph::activation{"relu"}, conv);
        return p;
    }
};

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

struct test_add_broadcast
{
Paul's avatar
Paul committed
201
    migraph::program create_program() const
Paul's avatar
Paul committed
202
    {
Paul's avatar
Paul committed
203
204
205
206
207
208
        migraph::program p;
        migraph::shape s{migraph::shape::float_type, {3}};
        auto x  = p.add_parameter("x", {migraph::shape::float_type, {2, 2, 3}});
        auto y  = p.add_parameter("y", {migraph::shape::float_type, {2, 2}});
        auto by = p.add_instruction(migraph::broadcast{0}, x, y);
        p.add_instruction(migraph::add{}, x, by);
Paul's avatar
Paul committed
209
210
211
212
        return p;
    }
};

Paul's avatar
Paul committed
213
214
215
216
217
218
219
220
221
222
223
224
225
226
struct test_add_broadcast2
{
    migraph::program create_program() const
    {
        migraph::program p;
        migraph::shape s{migraph::shape::float_type, {3}};
        auto x  = p.add_parameter("x", {migraph::shape::float_type, {2, 3, 4}});
        auto y  = p.add_parameter("y", {migraph::shape::float_type, {3}});
        auto by = p.add_instruction(migraph::broadcast{1}, x, y);
        p.add_instruction(migraph::add{}, x, by);
        return p;
    }
};

Paul's avatar
Paul committed
227
struct test_conv_relu
Paul's avatar
Paul committed
228
{
Paul's avatar
Paul committed
229
    migraph::program create_program() const
Paul's avatar
Paul committed
230
    {
Paul's avatar
Paul committed
231
        migraph::program p;
Paul's avatar
Paul committed
232
233
234
235
        auto input = p.add_parameter("x", migraph::shape{migraph::shape::float_type, {4, 3, 3, 3}});
        auto weights =
            p.add_parameter("w", migraph::shape{migraph::shape::float_type, {4, 3, 3, 3}});
        auto conv = p.add_instruction(migraph::convolution{}, input, weights);
Paul's avatar
Paul committed
236
        p.add_instruction(migraph::activation{"relu"}, conv);
Paul's avatar
Paul committed
237
238
239
240
        return p;
    }
};

Paul's avatar
Paul committed
241
242
243
244
245
struct test_add_relu
{
    migraph::program create_program() const
    {
        migraph::program p;
Paul's avatar
Paul committed
246
247
        auto x   = p.add_parameter("x", migraph::shape{migraph::shape::float_type, {4, 3, 3, 3}});
        auto y   = p.add_parameter("y", migraph::shape{migraph::shape::float_type, {4, 3, 3, 3}});
Paul's avatar
Paul committed
248
249
250
251
252
253
        auto add = p.add_instruction(migraph::add{}, x, y);
        p.add_instruction(migraph::activation{"relu"}, add);
        return p;
    }
};

Paul's avatar
Paul committed
254
255
struct test_conv_pooling
{
Paul's avatar
Paul committed
256
    migraph::program create_program() const
Paul's avatar
Paul committed
257
    {
Paul's avatar
Paul committed
258
        migraph::program p;
Paul's avatar
Paul committed
259
260
261
262
        auto input =
            p.add_parameter("x", migraph::shape{migraph::shape::float_type, {4, 3, 32, 32}});
        auto weights =
            p.add_parameter("w", migraph::shape{migraph::shape::float_type, {4, 3, 3, 3}});
Paul's avatar
Paul committed
263
264
265
        auto conv    = p.add_instruction(migraph::convolution{}, input, weights);
        auto pooling = p.add_instruction(migraph::pooling{"max"}, conv);
        p.add_instruction(migraph::activation{"relu"}, pooling);
Paul's avatar
Paul committed
266
267
268
269
        return p;
    }
};

Paul's avatar
Paul committed
270
271
struct test_gemm
{
Paul's avatar
Paul committed
272
    migraph::program create_program() const
Paul's avatar
Paul committed
273
    {
Paul's avatar
Paul committed
274
275
276
277
        migraph::program p;
        auto a = p.add_parameter("a", migraph::shape{migraph::shape::float_type, {4, 5}});
        auto b = p.add_parameter("b", migraph::shape{migraph::shape::float_type, {5, 3}});
        p.add_instruction(migraph::gemm{}, a, b);
Paul's avatar
Paul committed
278
279
280
281
        return p;
    }
};

Paul's avatar
Paul committed
282
283
284
285
286
287
288
289
290
291
292
293
struct test_gemm_ld
{
    migraph::program create_program() const
    {
        migraph::program p;
        auto a = p.add_parameter("a", migraph::shape{migraph::shape::float_type, {4, 5}, {10, 1}});
        auto b = p.add_parameter("b", migraph::shape{migraph::shape::float_type, {5, 3}, {20, 1}});
        p.add_instruction(migraph::gemm{}, a, b);
        return p;
    }
};

294
295
296
297
298
struct test_gemm_transposeb
{
    migraph::program create_program() const
    {
        migraph::program p;
Paul's avatar
Paul committed
299
300
        auto a  = p.add_parameter("a", migraph::shape{migraph::shape::float_type, {4, 5}});
        auto b  = p.add_parameter("b", migraph::shape{migraph::shape::float_type, {3, 5}});
301
302
303
304
305
306
307
308
309
310
311
        auto bt = p.add_instruction(migraph::transpose{{1, 0}}, b);
        p.add_instruction(migraph::gemm{}, a, bt);
        return p;
    }
};

struct test_gemm_transposea
{
    migraph::program create_program() const
    {
        migraph::program p;
Paul's avatar
Paul committed
312
313
        auto a  = p.add_parameter("a", migraph::shape{migraph::shape::float_type, {5, 4}});
        auto b  = p.add_parameter("b", migraph::shape{migraph::shape::float_type, {5, 3}});
314
315
316
317
318
319
320
321
322
323
324
        auto at = p.add_instruction(migraph::transpose{{1, 0}}, a);
        p.add_instruction(migraph::gemm{}, at, b);
        return p;
    }
};

struct test_gemm_transposeab
{
    migraph::program create_program() const
    {
        migraph::program p;
Paul's avatar
Paul committed
325
326
        auto a  = p.add_parameter("a", migraph::shape{migraph::shape::float_type, {5, 4}});
        auto b  = p.add_parameter("b", migraph::shape{migraph::shape::float_type, {3, 5}});
327
328
329
330
331
332
333
        auto at = p.add_instruction(migraph::transpose{{1, 0}}, a);
        auto bt = p.add_instruction(migraph::transpose{{1, 0}}, b);
        p.add_instruction(migraph::gemm{}, at, bt);
        return p;
    }
};

334
335
336
337
338
struct test_contiguous
{
    migraph::program create_program() const
    {
        migraph::program p;
339
        migraph::shape s{migraph::shape::float_type, {4, 4, 4, 3}, {48, 4, 1, 16}};
340
341
        auto x = p.add_parameter("x", s);
        p.add_instruction(migraph::contiguous{}, x);
Paul's avatar
Paul committed
342
        EXPECT(p.get_shape().standard());
343
344
345
346
        return p;
    }
};

347
struct test_transpose
348
{
349
350
351
352
353
354
355
356
357
358
359
    migraph::program create_program() const
    {
        migraph::program p;
        migraph::shape s{migraph::shape::float_type, {4, 3, 4, 4}};
        auto x                    = p.add_parameter("x", s);
        std::vector<int64_t> perm = {0, 2, 3, 1};
        auto l                    = p.add_instruction(migraph::transpose{perm}, x);
        p.add_instruction(migraph::contiguous{}, l);
        return p;
    }
};
360

Paul's avatar
Paul committed
361
362
363
364
365
366
367
368
369
370
371
372
373
374
struct test_batchnorm_inference_2
{
    const size_t width    = 14;
    const size_t height   = 14;
    const size_t channels = 256;
    const size_t batches  = 1;

    migraph::program create_program() const
    {
        migraph::program p;

        migraph::shape s{migraph::shape::float_type, {batches, channels, height, width}};
        migraph::shape vars{migraph::shape::float_type, {channels}};
        auto x        = p.add_parameter("x", s);
Paul's avatar
Paul committed
375
376
377
378
379
        auto scale    = p.add_literal(migraph::abs(migraph::generate_literal(vars, 1)));
        auto bias     = p.add_literal(migraph::abs(migraph::generate_literal(vars, 2)));
        auto mean     = p.add_literal(migraph::abs(migraph::generate_literal(vars, 3)));
        auto variance = p.add_literal(migraph::abs(migraph::generate_literal(vars, 4)));
        p.add_instruction(migraph::batch_norm_inference{}, x, scale, bias, mean, variance);
Paul's avatar
Paul committed
380
381
382
383
        return p;
    }
};

wsttiger's avatar
wsttiger committed
384
385
386
387
388
389
390
391
392
393
394
395
struct test_batchnorm_inference
{
    const size_t width    = 3;
    const size_t height   = 3;
    const size_t channels = 3;
    const size_t batches  = 4;

    migraph::program create_program() const
    {
        migraph::program p;

        migraph::shape s{migraph::shape::float_type, {batches, channels, height, width}};
396
        migraph::shape vars{migraph::shape::float_type, {channels}};
wsttiger's avatar
wsttiger committed
397
        auto x        = p.add_parameter("x", s);
Paul's avatar
Paul committed
398
399
400
401
402
        auto scale    = p.add_literal(migraph::abs(migraph::generate_literal(vars, 1)));
        auto bias     = p.add_literal(migraph::abs(migraph::generate_literal(vars, 2)));
        auto mean     = p.add_literal(migraph::abs(migraph::generate_literal(vars, 3)));
        auto variance = p.add_literal(migraph::abs(migraph::generate_literal(vars, 4)));
        p.add_instruction(migraph::batch_norm_inference{}, x, scale, bias, mean, variance);
wsttiger's avatar
wsttiger committed
403
404
405
406
        return p;
    }
};

Paul's avatar
Paul committed
407
408
409
410
411
412
413
414
struct test_conv_bn_relu_pooling
{
    migraph::program create_program() const
    {
        migraph::program p;

        migraph::shape xs{migraph::shape::float_type, {1, 3, 224, 224}};
        migraph::shape ws{migraph::shape::float_type, {64, 3, 7, 7}};
Paul's avatar
Paul committed
415
        migraph::shape vars{migraph::shape::float_type, {64}};
Paul's avatar
Paul committed
416
417
418
        auto x        = p.add_parameter("x", xs);
        auto w        = p.add_parameter("w", ws);
        auto conv     = p.add_instruction(migraph::convolution{{3, 3}, {2, 2}, {1, 1}}, x, w);
Paul's avatar
Paul committed
419
420
421
422
        auto scale    = p.add_literal(migraph::abs(migraph::generate_literal(vars, 1)));
        auto bias     = p.add_literal(migraph::abs(migraph::generate_literal(vars, 2)));
        auto mean     = p.add_literal(migraph::abs(migraph::generate_literal(vars, 3)));
        auto variance = p.add_literal(migraph::abs(migraph::generate_literal(vars, 4)));
Paul's avatar
Paul committed
423
        auto bn =
Paul's avatar
Paul committed
424
            p.add_instruction(migraph::batch_norm_inference{}, conv, scale, bias, mean, variance);
Paul's avatar
Paul committed
425
426
427
428
429
430
        auto relu = p.add_instruction(migraph::activation{"relu"}, bn);
        p.add_instruction(migraph::pooling{"average", {1, 1}, {2, 2}, {3, 3}}, relu);
        return p;
    }
};

Paul's avatar
Paul committed
431
432
int main()
{
Paul's avatar
Paul committed
433
    verify_program<test_add>();
Paul's avatar
Paul committed
434
    verify_program<test_add_broadcast>();
Paul's avatar
Paul committed
435
    verify_program<test_add_broadcast2>();
Paul's avatar
Paul committed
436
    verify_program<test_conv_relu>();
Paul's avatar
Paul committed
437
    verify_program<test_add_relu>();
Paul's avatar
Paul committed
438
439
    verify_program<test_conv_pooling>();
    verify_program<test_gemm>();
Paul's avatar
Paul committed
440
    // verify_program<test_gemm_ld>();
441
442
443
    verify_program<test_gemm_transposeb>();
    verify_program<test_gemm_transposea>();
    verify_program<test_gemm_transposeab>();
444
445
    verify_program<test_contiguous>();
    verify_program<test_transpose>();
446
    verify_program<test_batchnorm_inference>();
Paul's avatar
Paul committed
447
    verify_program<test_batchnorm_inference_2>();
Paul's avatar
Paul committed
448
    verify_program<test_conv_bn_relu_pooling>();
Paul's avatar
Paul committed
449
}