"docs/git@developer.sourcefind.cn:zhaoyu6/sglang.git" did not exist on "d27a6f7092cf7caa5b99089d78c692a529ff8dd2"
miopen.cpp 11.4 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
27
28
29
30
31
32
33
34
35
// An improved async, that doesn't block
template <class Function>
std::future<typename std::result_of<Function()>::type> detach_async(Function&& f)
{
    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);
}

Paul's avatar
Paul committed
36
37
struct auto_print
{
Paul's avatar
Paul committed
38
    static std::array<std::function<void()>, 2> handlers;
Paul's avatar
Paul committed
39
    int index;
Paul's avatar
Paul committed
40
    template <class T>
Paul's avatar
Paul committed
41
    auto_print(T& x, int i) : index(i)
Paul's avatar
Paul committed
42
    {
Paul's avatar
Paul committed
43
        handlers[index] = [&x] { std::cout << x << std::endl; };
Paul's avatar
Paul committed
44
    }
Paul's avatar
Paul committed
45

Paul's avatar
Paul committed
46
    ~auto_print()
Paul's avatar
Paul committed
47
    {
Paul's avatar
Paul committed
48
        handlers[index] = [] {};
Paul's avatar
Paul committed
49
50
    }
};
Paul's avatar
Paul committed
51
std::array<std::function<void()>, 2> auto_print::handlers = {};
Paul's avatar
Paul committed
52

Paul's avatar
Paul committed
53
template <class V>
Paul's avatar
Paul committed
54
migraph::argument run_cpu()
Paul's avatar
Paul committed
55
{
Paul's avatar
Paul committed
56
57
    V v;
    auto p = v.create_program();
Paul's avatar
Paul committed
58
    auto_print pp{p, 0};
Paul's avatar
Paul committed
59
    p.compile(migraph::cpu::cpu_target{});
Paul's avatar
Paul committed
60
    migraph::program::parameter_map m;
Paul's avatar
Paul committed
61
    for(auto&& x : p.get_parameter_shapes())
Paul's avatar
Paul committed
62
63
64
    {
        m[x.first] = migraph::generate_argument(x.second);
    }
Paul's avatar
Paul committed
65
    return p.eval(m);
Paul's avatar
Paul committed
66
67
}

Paul's avatar
Paul committed
68
template <class V>
Paul's avatar
Paul committed
69
migraph::argument run_gpu()
Paul's avatar
Paul committed
70
{
Paul's avatar
Paul committed
71
72
    V v;
    auto p = v.create_program();
Paul's avatar
Paul committed
73
    auto_print pp{p, 1};
Paul's avatar
Paul committed
74
    p.compile(migraph::gpu::target{});
Paul's avatar
Paul committed
75

Paul's avatar
Paul committed
76
    migraph::program::parameter_map m;
Paul's avatar
Paul committed
77
    for(auto&& x : p.get_parameter_shapes())
Paul's avatar
Paul committed
78
    {
Paul's avatar
Paul committed
79
        m[x.first] = migraph::gpu::to_gpu(migraph::generate_argument(x.second));
Paul's avatar
Paul committed
80
81
    }

Paul's avatar
Paul committed
82
    return migraph::gpu::from_gpu(p.eval(m));
Paul's avatar
Paul committed
83
84
}

Paul's avatar
Paul committed
85
template <class V>
Paul's avatar
Paul committed
86
void verify_program()
Paul's avatar
Paul committed
87
{
Paul's avatar
Paul committed
88
    std::set_terminate(+[] {
Paul's avatar
Paul committed
89
        std::cout << "FAILED: " << migraph::get_type_name<V>() << std::endl;
Paul's avatar
Paul committed
90
91
92
93
94
95
        try
        {
            std::rethrow_exception(std::current_exception());
        }
        catch(const std::exception& e)
        {
Paul's avatar
Paul committed
96
            std::cout << "    what(): " << e.what() << std::endl;
Paul's avatar
Paul committed
97
        }
Paul's avatar
Paul committed
98
99
        std::cout << std::endl;
        for(auto&& handle : auto_print::handlers)
Paul's avatar
Paul committed
100
            handle();
Paul's avatar
Paul committed
101
    });
Paul's avatar
Paul committed
102
    auto cpu_arg_f = detach_async([] { return run_cpu<V>(); });
Paul's avatar
Paul committed
103
    auto gpu_arg   = run_gpu<V>();
Paul's avatar
Paul committed
104
    visit_all(cpu_arg_f.get(), gpu_arg)([](auto cpu, auto gpu) {
Paul's avatar
Paul committed
105
        if(not migraph::verify_range(cpu, gpu))
Paul's avatar
Paul committed
106
        {
107
            // TODO: Check for nans
Paul's avatar
Paul committed
108
109
110
            std::cout << "FAILED: " << migraph::get_type_name<V>() << std::endl;
        }
    });
Paul's avatar
Paul committed
111
    std::set_terminate(nullptr);
Paul's avatar
Paul committed
112
113
}

Paul's avatar
Paul committed
114
115
116
117
118
struct test_literals
{
    migraph::program create_program() const
    {
        migraph::program p;
Paul's avatar
Paul committed
119
120
121
122
        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
123
124
125
126
127
128
        auto conv = p.add_instruction(migraph::convolution{}, input, weights);
        p.add_instruction(migraph::activation{"relu"}, conv);
        return p;
    }
};

Paul's avatar
Paul committed
129
130
struct test_add
{
Paul's avatar
Paul committed
131
    migraph::program create_program() const
Paul's avatar
Paul committed
132
    {
Paul's avatar
Paul committed
133
134
        migraph::program p;
        migraph::shape s{migraph::shape::float_type, {3}};
Paul's avatar
Paul committed
135
136
        auto x = p.add_parameter("x", s);
        auto y = p.add_parameter("y", s);
Paul's avatar
Paul committed
137
        p.add_instruction(migraph::add{}, x, y);
Paul's avatar
Paul committed
138
139
140
141
142
143
        return p;
    }
};

struct test_add_broadcast
{
Paul's avatar
Paul committed
144
    migraph::program create_program() const
Paul's avatar
Paul committed
145
    {
Paul's avatar
Paul committed
146
147
148
149
150
151
        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
152
153
154
155
        return p;
    }
};

Paul's avatar
Paul committed
156
struct test_conv_relu
Paul's avatar
Paul committed
157
{
Paul's avatar
Paul committed
158
    migraph::program create_program() const
Paul's avatar
Paul committed
159
    {
Paul's avatar
Paul committed
160
        migraph::program p;
Paul's avatar
Paul committed
161
162
163
164
        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
165
        p.add_instruction(migraph::activation{"relu"}, conv);
Paul's avatar
Paul committed
166
167
168
169
        return p;
    }
};

Paul's avatar
Paul committed
170
171
struct test_conv_pooling
{
Paul's avatar
Paul committed
172
    migraph::program create_program() const
Paul's avatar
Paul committed
173
    {
Paul's avatar
Paul committed
174
        migraph::program p;
Paul's avatar
Paul committed
175
176
177
178
        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
179
180
181
        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
182
183
184
185
        return p;
    }
};

Paul's avatar
Paul committed
186
187
struct test_gemm
{
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
        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
194
195
196
197
        return p;
    }
};

Paul's avatar
Paul committed
198
199
200
201
202
203
204
205
206
207
208
209
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;
    }
};

210
211
212
213
214
struct test_gemm_transposeb
{
    migraph::program create_program() const
    {
        migraph::program p;
Paul's avatar
Paul committed
215
216
        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}});
217
218
219
220
221
222
223
224
225
226
227
        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
228
229
        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}});
230
231
232
233
234
235
236
237
238
239
240
        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
241
242
        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}});
243
244
245
246
247
248
249
        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;
    }
};

250
251
252
253
254
struct test_contiguous
{
    migraph::program create_program() const
    {
        migraph::program p;
255
        migraph::shape s{migraph::shape::float_type, {4, 4, 4, 3}, {48, 4, 1, 16}};
256
257
258
259
260
261
        auto x = p.add_parameter("x", s);
        p.add_instruction(migraph::contiguous{}, x);
        return p;
    }
};

262
struct test_transpose
263
{
264
265
266
267
268
269
270
271
272
273
274
    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;
    }
};
275

Paul's avatar
Paul committed
276
277
278
279
280
281
282
283
284
285
286
287
288
289
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);
290
291
292
293
294
        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
295
296
297
298
        return p;
    }
};

wsttiger's avatar
wsttiger committed
299
300
301
302
303
304
305
306
307
308
309
310
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}};
311
        migraph::shape vars{migraph::shape::float_type, {channels}};
wsttiger's avatar
wsttiger committed
312
        auto x        = p.add_parameter("x", s);
313
314
315
316
317
318
319
320
321
322
323
324
325
326
327
328
329
330
331
332
333
334
335
336
337
338
339
340
341
        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);
        return p;
    }
};

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}};
        migraph::shape vars{migraph::shape::float_type, {64}};
        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);
        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)));
        auto bn =
            p.add_instruction(migraph::batch_norm_inference{}, conv, scale, bias, mean, variance);
        auto relu = p.add_instruction(migraph::activation{"relu"}, bn);
        p.add_instruction(migraph::pooling{"average", {1, 1}, {2, 2}, {3, 3}}, relu);
wsttiger's avatar
wsttiger committed
342
343
344
345
        return p;
    }
};

Paul's avatar
Paul committed
346
347
int main()
{
Paul's avatar
Paul committed
348
    verify_program<test_add>();
Paul's avatar
Paul committed
349
    verify_program<test_add_broadcast>();
Paul's avatar
Paul committed
350
351
352
    verify_program<test_conv_relu>();
    verify_program<test_conv_pooling>();
    verify_program<test_gemm>();
Paul's avatar
Paul committed
353
    // verify_program<test_gemm_ld>();
354
355
356
    verify_program<test_gemm_transposeb>();
    verify_program<test_gemm_transposea>();
    verify_program<test_gemm_transposeab>();
357
358
    verify_program<test_contiguous>();
    verify_program<test_transpose>();
359
    verify_program<test_batchnorm_inference>();
Paul's avatar
Paul committed
360
    verify_program<test_batchnorm_inference_2>();
Paul's avatar
Paul committed
361
}