miopen.cpp 11.7 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
40
    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
    {
        return std::async(std::launch::deferred, std::move(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
    static std::array<std::function<void()>, 2> handlers;
Paul's avatar
Paul committed
47
    int index;
Paul's avatar
Paul committed
48
    template <class T>
Paul's avatar
Paul committed
49
    auto_print(T& x, int i) : index(i)
Paul's avatar
Paul committed
50
    {
Paul's avatar
Paul committed
51
        handlers[index] = [&x] { std::cout << x << std::endl; };
Paul's avatar
Paul committed
52
    }
Paul's avatar
Paul committed
53

Paul's avatar
Paul committed
54
    ~auto_print()
Paul's avatar
Paul committed
55
    {
Paul's avatar
Paul committed
56
        handlers[index] = [] {};
Paul's avatar
Paul committed
57
58
    }
};
Paul's avatar
Paul committed
59
std::array<std::function<void()>, 2> auto_print::handlers = {};
Paul's avatar
Paul committed
60

Paul's avatar
Paul committed
61
62
63
void compile_check(migraph::program& p, migraph::target t)
{
    auto name = t.name();
Paul's avatar
Paul committed
64
    auto s    = p.get_shape();
Paul's avatar
Paul committed
65
66
    std::stringstream ss;
    p.compile(std::move(t), migraph::tracer{ss});
Paul's avatar
Paul committed
67
    if(p.get_shape() != s)
Paul's avatar
Paul committed
68
69
70
71
72
73
    {
        std::cout << ss.str() << std::endl;
        throw std::runtime_error("Compiling program with " + name + " alters its shape");
    }
}

Paul's avatar
Paul committed
74
template <class V>
Paul's avatar
Paul committed
75
migraph::argument run_cpu()
Paul's avatar
Paul committed
76
{
Paul's avatar
Paul committed
77
78
    V v;
    auto p = v.create_program();
Paul's avatar
Paul committed
79
    auto_print pp{p, 0};
Paul's avatar
Paul committed
80
    compile_check(p, migraph::cpu::cpu_target{});
Paul's avatar
Paul committed
81
    migraph::program::parameter_map m;
Paul's avatar
Paul committed
82
    for(auto&& x : p.get_parameter_shapes())
Paul's avatar
Paul committed
83
84
85
    {
        m[x.first] = migraph::generate_argument(x.second);
    }
Paul's avatar
Paul committed
86
    return p.eval(m);
Paul's avatar
Paul committed
87
88
}

Paul's avatar
Paul committed
89
template <class V>
Paul's avatar
Paul committed
90
migraph::argument run_gpu()
Paul's avatar
Paul committed
91
{
Paul's avatar
Paul committed
92
93
    V v;
    auto p = v.create_program();
Paul's avatar
Paul committed
94
    auto_print pp{p, 1};
Paul's avatar
Paul committed
95
    compile_check(p, migraph::gpu::target{});
Paul's avatar
Paul committed
96

Paul's avatar
Paul committed
97
    migraph::program::parameter_map m;
Paul's avatar
Paul committed
98
    for(auto&& x : p.get_parameter_shapes())
Paul's avatar
Paul committed
99
    {
Paul's avatar
Paul committed
100
        m[x.first] = migraph::gpu::to_gpu(migraph::generate_argument(x.second));
Paul's avatar
Paul committed
101
102
    }

Paul's avatar
Paul committed
103
    return migraph::gpu::from_gpu(p.eval(m));
Paul's avatar
Paul committed
104
105
}

Paul's avatar
Paul committed
106
template <class V>
Paul's avatar
Paul committed
107
void verify_program()
Paul's avatar
Paul committed
108
{
Paul's avatar
Paul committed
109
    std::set_terminate(+[] {
Paul's avatar
Paul committed
110
        std::cout << "FAILED: " << migraph::get_type_name<V>() << std::endl;
Paul's avatar
Paul committed
111
112
113
114
115
116
        try
        {
            std::rethrow_exception(std::current_exception());
        }
        catch(const std::exception& e)
        {
Paul's avatar
Paul committed
117
            std::cout << "    what(): " << e.what() << std::endl;
Paul's avatar
Paul committed
118
        }
Paul's avatar
Paul committed
119
120
        std::cout << std::endl;
        for(auto&& handle : auto_print::handlers)
Paul's avatar
Paul committed
121
            handle();
Paul's avatar
Paul committed
122
    });
Paul's avatar
Paul committed
123
    auto cpu_arg_f = detach_async([] { return run_cpu<V>(); });
Paul's avatar
Paul committed
124
    auto gpu_arg   = run_gpu<V>();
Paul's avatar
Paul committed
125
    visit_all(cpu_arg_f.get(), gpu_arg)([](auto cpu, auto gpu) {
Paul's avatar
Paul committed
126
        if(not migraph::verify_range(cpu, gpu))
Paul's avatar
Paul committed
127
        {
128
            // TODO: Check for nans
Paul's avatar
Paul committed
129
            std::cout << "FAILED: " << migraph::get_type_name<V>() << std::endl;
Paul's avatar
Paul committed
130
131
132
133
134
135
136
137
138
139
140
141
142
143
144
145
146
147
148
149
150
151
152
            // 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
153
154
        }
    });
Paul's avatar
Paul committed
155
    std::set_terminate(nullptr);
Paul's avatar
Paul committed
156
157
}

Paul's avatar
Paul committed
158
159
160
161
162
struct test_literals
{
    migraph::program create_program() const
    {
        migraph::program p;
Paul's avatar
Paul committed
163
164
165
166
        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
167
168
169
170
171
172
        auto conv = p.add_instruction(migraph::convolution{}, input, weights);
        p.add_instruction(migraph::activation{"relu"}, conv);
        return p;
    }
};

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

struct test_add_broadcast
{
Paul's avatar
Paul committed
188
    migraph::program create_program() const
Paul's avatar
Paul committed
189
    {
Paul's avatar
Paul committed
190
191
192
193
194
195
        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
196
197
198
199
        return p;
    }
};

Paul's avatar
Paul committed
200
struct test_conv_relu
Paul's avatar
Paul committed
201
{
Paul's avatar
Paul committed
202
    migraph::program create_program() const
Paul's avatar
Paul committed
203
    {
Paul's avatar
Paul committed
204
        migraph::program p;
Paul's avatar
Paul committed
205
206
207
208
        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
209
        p.add_instruction(migraph::activation{"relu"}, conv);
Paul's avatar
Paul committed
210
211
212
213
        return p;
    }
};

Paul's avatar
Paul committed
214
215
struct test_conv_pooling
{
Paul's avatar
Paul committed
216
    migraph::program create_program() const
Paul's avatar
Paul committed
217
    {
Paul's avatar
Paul committed
218
        migraph::program p;
Paul's avatar
Paul committed
219
220
221
222
        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
223
224
225
        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
226
227
228
229
        return p;
    }
};

Paul's avatar
Paul committed
230
231
struct test_gemm
{
Paul's avatar
Paul committed
232
    migraph::program create_program() const
Paul's avatar
Paul committed
233
    {
Paul's avatar
Paul committed
234
235
236
237
        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
238
239
240
241
        return p;
    }
};

Paul's avatar
Paul committed
242
243
244
245
246
247
248
249
250
251
252
253
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;
    }
};

254
255
256
257
258
struct test_gemm_transposeb
{
    migraph::program create_program() const
    {
        migraph::program p;
Paul's avatar
Paul committed
259
260
        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}});
261
262
263
264
265
266
267
268
269
270
271
        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
272
273
        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}});
274
275
276
277
278
279
280
281
282
283
284
        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
285
286
        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}});
287
288
289
290
291
292
293
        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;
    }
};

294
295
296
297
298
struct test_contiguous
{
    migraph::program create_program() const
    {
        migraph::program p;
299
        migraph::shape s{migraph::shape::float_type, {4, 4, 4, 3}, {48, 4, 1, 16}};
300
301
        auto x = p.add_parameter("x", s);
        p.add_instruction(migraph::contiguous{}, x);
Paul's avatar
Paul committed
302
        EXPECT(p.get_shape().standard());
303
304
305
306
        return p;
    }
};

307
struct test_transpose
308
{
309
310
311
312
313
314
315
316
317
318
319
    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;
    }
};
320

Paul's avatar
Paul committed
321
322
323
324
325
326
327
328
329
330
331
332
333
334
335
336
337
338
339
340
341
342
343
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);
        auto mean     = p.add_parameter("mean", vars);
        auto variance = p.add_parameter("variance", vars);
        auto scale    = p.add_parameter("scale", vars);
        auto bias     = p.add_parameter("bias", vars);
        p.add_instruction(migraph::batch_norm_inference{}, x, mean, variance, scale, bias);
        return p;
    }
};

wsttiger's avatar
wsttiger committed
344
345
346
347
348
349
350
351
352
353
354
355
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}};
356
        migraph::shape vars{migraph::shape::float_type, {channels}};
wsttiger's avatar
wsttiger committed
357
358
359
360
361
362
363
364
365
366
        auto x        = p.add_parameter("x", s);
        auto mean     = p.add_parameter("mean", vars);
        auto variance = p.add_parameter("variance", vars);
        auto scale    = p.add_parameter("scale", vars);
        auto bias     = p.add_parameter("bias", vars);
        p.add_instruction(migraph::batch_norm_inference{}, x, mean, variance, scale, bias);
        return p;
    }
};

Paul's avatar
Paul committed
367
368
int main()
{
Paul's avatar
Paul committed
369
    verify_program<test_add>();
Paul's avatar
Paul committed
370
    verify_program<test_add_broadcast>();
Paul's avatar
Paul committed
371
372
373
    verify_program<test_conv_relu>();
    verify_program<test_conv_pooling>();
    verify_program<test_gemm>();
Paul's avatar
Paul committed
374
    // verify_program<test_gemm_ld>();
375
376
377
    verify_program<test_gemm_transposeb>();
    verify_program<test_gemm_transposea>();
    verify_program<test_gemm_transposeab>();
378
379
    verify_program<test_contiguous>();
    verify_program<test_transpose>();
380
    verify_program<test_batchnorm_inference>();
Paul's avatar
Paul committed
381
    verify_program<test_batchnorm_inference_2>();
Paul's avatar
Paul committed
382
}