inline_asm.hpp 10.5 KB
Newer Older
Jing Zhang's avatar
Jing Zhang committed
1
2
3
4
5
6
#pragma once

typedef float Float4 __attribute__((ext_vector_type(4)));

extern "C" __attribute__((address_space(3))) void* __to_local(void* p)[[hc]];

Jing Zhang's avatar
Jing Zhang committed
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
31
32
inline __device__ void vmcnt(int cnt) {
    if(cnt == 0) {
        asm volatile ("\n \
                s_waitcnt vmcnt(0) \n \
                "::);
    }
    else if(cnt == 1) {
        asm volatile ("\n \
                s_waitcnt vmcnt(1) \n \
                "::);
    }
    else if(cnt == 2) {
        asm volatile ("\n \
                s_waitcnt vmcnt(2) \n \
                "::);
    }
    else if(cnt == 4) {
        asm volatile ("\n \
                s_waitcnt vmcnt(2) \n \
                "::);
    }
    else {
        assert(0);
    }
}

Chao Liu's avatar
Chao Liu committed
33
34
inline __device__ void lgkmcnt(int cnt)
{
Jing Zhang's avatar
Jing Zhang committed
35
#if 1
Chao Liu's avatar
Chao Liu committed
36
37
    if(cnt == 0)
    {
Jing Zhang's avatar
Jing Zhang committed
38
39
        asm volatile("\n \
                s_waitcnt lgkmcnt(0) \n \
Chao Liu's avatar
Chao Liu committed
40
                " ::);
Jing Zhang's avatar
Jing Zhang committed
41
    }
Chao Liu's avatar
Chao Liu committed
42
43
    else if(cnt == 1)
    {
Jing Zhang's avatar
Jing Zhang committed
44
45
        asm volatile("\n \
                s_waitcnt lgkmcnt(1) \n \
Chao Liu's avatar
Chao Liu committed
46
                " ::);
Jing Zhang's avatar
Jing Zhang committed
47
    }
Chao Liu's avatar
Chao Liu committed
48
49
    else if(cnt == 2)
    {
Jing Zhang's avatar
Jing Zhang committed
50
51
        asm volatile("\n \
                s_waitcnt lgkmcnt(2) \n \
Chao Liu's avatar
Chao Liu committed
52
                " ::);
Jing Zhang's avatar
Jing Zhang committed
53
    }
Chao Liu's avatar
Chao Liu committed
54
55
    else if(cnt == 3)
    {
Jing Zhang's avatar
Jing Zhang committed
56
57
        asm volatile("\n \
                s_waitcnt lgkmcnt(3) \n \
Chao Liu's avatar
Chao Liu committed
58
                " ::);
Jing Zhang's avatar
Jing Zhang committed
59
    }
Chao Liu's avatar
Chao Liu committed
60
61
    else if(cnt == 4)
    {
Jing Zhang's avatar
Jing Zhang committed
62
63
        asm volatile("\n \
                s_waitcnt lgkmcnt(4) \n \
Chao Liu's avatar
Chao Liu committed
64
                " ::);
Jing Zhang's avatar
Jing Zhang committed
65
    }
Chao Liu's avatar
Chao Liu committed
66
67
    else
    {
Jing Zhang's avatar
Jing Zhang committed
68
        assert(0);
Jing Zhang's avatar
Jing Zhang committed
69
    }
Jing Zhang's avatar
Jing Zhang committed
70
#endif
Jing Zhang's avatar
Jing Zhang committed
71
72
}

Chao Liu's avatar
Chao Liu committed
73
74
inline __device__ void outerProduct1x4(const float* a, const float* b, float* c)
{
Jing Zhang's avatar
Jing Zhang committed
75
76
77
78
79
80
    asm volatile("\n \
            v_mac_f32 %0, %4, %5 \n \
            v_mac_f32 %1, %4, %6 \n \
            v_mac_f32 %2, %4, %7 \n \
            v_mac_f32 %3, %4, %8 \n \
            "
Chao Liu's avatar
Chao Liu committed
81
82
83
84
85
86
87
88
89
90
                 : "=v"(c[0]), "=v"(c[1]), "=v"(c[2]), "=v"(c[3])
                 : "v"(a[0]),
                   "v"(b[0]),
                   "v"(b[1]),
                   "v"(b[2]),
                   "v"(b[3]),
                   "0"(c[0]),
                   "1"(c[1]),
                   "2"(c[2]),
                   "3"(c[3]));
Jing Zhang's avatar
Jing Zhang committed
91
92
}

Chao Liu's avatar
Chao Liu committed
93
94
inline __device__ void outerProduct1x4(const float& a, const Float4& b, Float4& c)
{
Jing Zhang's avatar
Jing Zhang committed
95
96
97
98
99
100
101
102
103
104
105
106
107
#if 0
    asm volatile(
            "\n \
            v_mac_f32 %0, %4, %5 \n \
            v_mac_f32 %1, %4, %6 \n \
            v_mac_f32 %2, %4, %7 \n \
            v_mac_f32 %3, %4, %8 \n \
            "
            :
            :"v"(c.x),"v"(c.y),"v"(c.z),"v"(c.w), \
            "v"(a.x),"v"(b.x),"v"(b.y),"v"(b.z),"v"(b.w)
            );
#else
Chao Liu's avatar
Chao Liu committed
108
    outerProduct1x4(&a, (float*)&b, (float*)&c);
Jing Zhang's avatar
Jing Zhang committed
109
110
111
#endif
}

Chao Liu's avatar
Chao Liu committed
112
113
114
inline __device__ void
outerProduct4x4(const Float4& a, const Float4& b, Float4& c0, Float4& c1, Float4& c2, Float4& c3)
{
Jing Zhang's avatar
Jing Zhang committed
115
116
117
118
119
120
121
122
123
124
125
126
127
128
129
130
131
132
133
134
135
136
137
138
139
140
141
142
143
144
145
146
147
148
149
150
151
152
153
154
155
156
157
158
159
160
161
162
163
164
165
166
167
#if 0
    asm volatile(
            "\n \
            v_mac_f32 %0, %4, %5 \n \
            v_mac_f32 %1, %4, %6 \n \
            v_mac_f32 %2, %4, %7 \n \
            v_mac_f32 %3, %4, %8 \n \
            "
            :
            :"v"(c0.x),"v"(c0.y),"v"(c0.z),"v"(c0.w), \
            "v"(a.x),"v"(b.x),"v"(b.y),"v"(b.z),"v"(b.w)
            );
    asm volatile(
            "\n \
            v_mac_f32 %0, %4, %5 \n \
            v_mac_f32 %1, %4, %6 \n \
            v_mac_f32 %2, %4, %7 \n \
            v_mac_f32 %3, %4, %8 \n \
            "
            :
            :"v"(c1.x),"v"(c1.y),"v"(c1.z),"v"(c1.w), \
            "v"(a.y),"v"(b.x),"v"(b.y),"v"(b.z),"v"(b.w)
            );
    asm volatile(
            "\n \
            v_mac_f32 %0, %4, %5 \n \
            v_mac_f32 %1, %4, %6 \n \
            v_mac_f32 %2, %4, %7 \n \
            v_mac_f32 %3, %4, %8 \n \
            "
            :
            :"v"(c2.x),"v"(c2.y),"v"(c2.z),"v"(c2.w), \
            "v"(a.z),"v"(b.x),"v"(b.y),"v"(b.z),"v"(b.w)
            );
    asm volatile(
            "\n \
            v_mac_f32 %0, %4, %5 \n \
            v_mac_f32 %1, %4, %6 \n \
            v_mac_f32 %2, %4, %7 \n \
            v_mac_f32 %3, %4, %8 \n \
            "
            :
            :"v"(c3.x),"v"(c3.y),"v"(c3.z),"v"(c3.w), \
            "v"(a.w),"v"(b.x),"v"(b.y),"v"(b.z),"v"(b.w)
            );
#else
    outerProduct1x4(a.x, b, c0);
    outerProduct1x4(a.y, b, c1);
    outerProduct1x4(a.z, b, c2);
    outerProduct1x4(a.w, b, c3);
#endif
}

Chao Liu's avatar
Chao Liu committed
168
inline __device__ void outerProduct8x8(const Float4* a, const Float4* b, Float4* c)
Jing Zhang's avatar
Jing Zhang committed
169
170
171
172
173
174
175
{
    outerProduct4x4(a[0], b[0], c[0], c[2], c[4], c[6]);
    outerProduct4x4(a[0], b[1], c[1], c[3], c[5], c[7]);
    outerProduct4x4(a[1], b[0], c[8], c[10], c[12], c[14]);
    outerProduct4x4(a[1], b[1], c[9], c[11], c[13], c[15]);
}

Chao Liu's avatar
Chao Liu committed
176
inline __device__ void ds_read_b128(Float4& r, void* lds, int offset = 0)
Jing Zhang's avatar
Jing Zhang committed
177
178
179
180
181
182
{
    if(offset == 0)
    {
        asm volatile("\n \
                ds_read_b128 %0, %1 offset:0 \n \
                "
Chao Liu's avatar
Chao Liu committed
183
184
                     : "=v"(r)
                     : "v"(__to_local(lds)));
Jing Zhang's avatar
Jing Zhang committed
185
186
187
188
189
190
    }
    else if(offset == 128)
    {
        asm volatile("\n \
                ds_read_b128 %0, %1 offset:128 \n \
                "
Chao Liu's avatar
Chao Liu committed
191
192
                     : "=v"(r)
                     : "v"(__to_local(lds)));
Jing Zhang's avatar
Jing Zhang committed
193
194
195
196
197
198
    }
    else if(offset == 256)
    {
        asm volatile("\n \
                ds_read_b128 %0, %1 offset:256 \n \
                "
Chao Liu's avatar
Chao Liu committed
199
200
                     : "=v"(r)
                     : "v"(__to_local(lds)));
Jing Zhang's avatar
Jing Zhang committed
201
202
203
204
205
206
    }
    else if(offset == 384)
    {
        asm volatile("\n \
                ds_read_b128 %0, %1 offset:384 \n \
                "
Chao Liu's avatar
Chao Liu committed
207
208
                     : "=v"(r)
                     : "v"(__to_local(lds)));
Jing Zhang's avatar
Jing Zhang committed
209
210
211
212
213
214
    }
    else if(offset == 512)
    {
        asm volatile("\n \
                ds_read_b128 %0, %1 offset:512 \n \
                "
Chao Liu's avatar
Chao Liu committed
215
216
                     : "=v"(r)
                     : "v"(__to_local(lds)));
Jing Zhang's avatar
Jing Zhang committed
217
218
219
220
221
222
    }
    else if(offset == 640)
    {
        asm volatile("\n \
                ds_read_b128 %0, %1 offset:640 \n \
                "
Chao Liu's avatar
Chao Liu committed
223
224
                     : "=v"(r)
                     : "v"(__to_local(lds)));
Jing Zhang's avatar
Jing Zhang committed
225
226
227
228
229
230
    }
    else if(offset == 768)
    {
        asm volatile("\n \
                ds_read_b128 %0, %1 offset:768 \n \
                "
Chao Liu's avatar
Chao Liu committed
231
232
                     : "=v"(r)
                     : "v"(__to_local(lds)));
Jing Zhang's avatar
Jing Zhang committed
233
234
235
236
237
238
    }
    else if(offset == 896)
    {
        asm volatile("\n \
                ds_read_b128 %0, %1 offset:896 \n \
                "
Chao Liu's avatar
Chao Liu committed
239
240
                     : "=v"(r)
                     : "v"(__to_local(lds)));
Jing Zhang's avatar
Jing Zhang committed
241
242
243
244
245
246
    }
    else if(offset == 1024)
    {
        asm volatile("\n \
                ds_read_b128 %0, %1 offset:1024 \n \
                "
Chao Liu's avatar
Chao Liu committed
247
248
                     : "=v"(r)
                     : "v"(__to_local(lds)));
Jing Zhang's avatar
Jing Zhang committed
249
250
251
252
253
254
    }
    else if(offset == 1152)
    {
        asm volatile("\n \
                ds_read_b128 %0, %1 offset:1152 \n \
                "
Chao Liu's avatar
Chao Liu committed
255
256
                     : "=v"(r)
                     : "v"(__to_local(lds)));
Jing Zhang's avatar
Jing Zhang committed
257
258
259
260
261
262
    }
    else if(offset == 1280)
    {
        asm volatile("\n \
                ds_read_b128 %0, %1 offset:1280 \n \
                "
Chao Liu's avatar
Chao Liu committed
263
264
                     : "=v"(r)
                     : "v"(__to_local(lds)));
Jing Zhang's avatar
Jing Zhang committed
265
266
267
268
269
270
    }
    else if(offset == 1408)
    {
        asm volatile("\n \
                ds_read_b128 %0, %1 offset:1408 \n \
                "
Chao Liu's avatar
Chao Liu committed
271
272
                     : "=v"(r)
                     : "v"(__to_local(lds)));
Jing Zhang's avatar
Jing Zhang committed
273
274
275
276
277
278
    }
    else if(offset == 1536)
    {
        asm volatile("\n \
                ds_read_b128 %0, %1 offset:1536 \n \
                "
Chao Liu's avatar
Chao Liu committed
279
280
                     : "=v"(r)
                     : "v"(__to_local(lds)));
Jing Zhang's avatar
Jing Zhang committed
281
282
283
284
285
286
    }
    else if(offset == 1664)
    {
        asm volatile("\n \
                ds_read_b128 %0, %1 offset:1664 \n \
                "
Chao Liu's avatar
Chao Liu committed
287
288
                     : "=v"(r)
                     : "v"(__to_local(lds)));
Jing Zhang's avatar
Jing Zhang committed
289
290
291
292
293
294
    }
    else if(offset == 1792)
    {
        asm volatile("\n \
                ds_read_b128 %0, %1 offset:1792 \n \
                "
Chao Liu's avatar
Chao Liu committed
295
296
                     : "=v"(r)
                     : "v"(__to_local(lds)));
Jing Zhang's avatar
Jing Zhang committed
297
298
299
300
301
302
    }
    else if(offset == 1920)
    {
        asm volatile("\n \
                ds_read_b128 %0, %1 offset:1920 \n \
                "
Chao Liu's avatar
Chao Liu committed
303
304
                     : "=v"(r)
                     : "v"(__to_local(lds)));
Jing Zhang's avatar
Jing Zhang committed
305
306
307
308
309
310
    }
    else if(offset == 2048)
    {
        asm volatile("\n \
                ds_read_b128 %0, %1 offset:2048 \n \
                "
Chao Liu's avatar
Chao Liu committed
311
312
                     : "=v"(r)
                     : "v"(__to_local(lds)));
Jing Zhang's avatar
Jing Zhang committed
313
314
315
316
317
318
    }
    else if(offset == 2176)
    {
        asm volatile("\n \
                ds_read_b128 %0, %1 offset:2176 \n \
                "
Chao Liu's avatar
Chao Liu committed
319
320
                     : "=v"(r)
                     : "v"(__to_local(lds)));
Jing Zhang's avatar
Jing Zhang committed
321
322
323
324
325
326
    }
    else if(offset == 2304)
    {
        asm volatile("\n \
                ds_read_b128 %0, %1 offset:2304 \n \
                "
Chao Liu's avatar
Chao Liu committed
327
328
                     : "=v"(r)
                     : "v"(__to_local(lds)));
Jing Zhang's avatar
Jing Zhang committed
329
330
331
332
333
334
    }
    else if(offset == 2560)
    {
        asm volatile("\n \
                ds_read_b128 %0, %1 offset:2560 \n \
                "
Chao Liu's avatar
Chao Liu committed
335
336
                     : "=v"(r)
                     : "v"(__to_local(lds)));
Jing Zhang's avatar
Jing Zhang committed
337
338
339
340
341
342
    }
    else if(offset == 2816)
    {
        asm volatile("\n \
                ds_read_b128 %0, %1 offset:2816 \n \
                "
Chao Liu's avatar
Chao Liu committed
343
344
                     : "=v"(r)
                     : "v"(__to_local(lds)));
Jing Zhang's avatar
Jing Zhang committed
345
346
347
348
349
350
    }
    else if(offset == 3072)
    {
        asm volatile("\n \
                ds_read_b128 %0, %1 offset:3072 \n \
                "
Chao Liu's avatar
Chao Liu committed
351
352
                     : "=v"(r)
                     : "v"(__to_local(lds)));
Jing Zhang's avatar
Jing Zhang committed
353
354
355
356
357
358
    }
    else if(offset == 3328)
    {
        asm volatile("\n \
                ds_read_b128 %0, %1 offset:3328 \n \
                "
Chao Liu's avatar
Chao Liu committed
359
360
                     : "=v"(r)
                     : "v"(__to_local(lds)));
Jing Zhang's avatar
Jing Zhang committed
361
362
363
364
365
366
    }
    else if(offset == 3584)
    {
        asm volatile("\n \
                ds_read_b128 %0, %1 offset:3584 \n \
                "
Chao Liu's avatar
Chao Liu committed
367
368
                     : "=v"(r)
                     : "v"(__to_local(lds)));
Jing Zhang's avatar
Jing Zhang committed
369
370
371
372
373
374
    }
    else if(offset == 3840)
    {
        asm volatile("\n \
                ds_read_b128 %0, %1 offset:3840 \n \
                "
Chao Liu's avatar
Chao Liu committed
375
376
                     : "=v"(r)
                     : "v"(__to_local(lds)));
Jing Zhang's avatar
Jing Zhang committed
377
378
379
380
381
382
    }
    else if(offset == 4096)
    {
        asm volatile("\n \
                ds_read_b128 %0, %1 offset:4096 \n \
                "
Chao Liu's avatar
Chao Liu committed
383
384
                     : "=v"(r)
                     : "v"(__to_local(lds)));
Jing Zhang's avatar
Jing Zhang committed
385
386
387
388
389
390
    }
    else if(offset == 4352)
    {
        asm volatile("\n \
                ds_read_b128 %0, %1 offset:4352 \n \
                "
Chao Liu's avatar
Chao Liu committed
391
392
                     : "=v"(r)
                     : "v"(__to_local(lds)));
Jing Zhang's avatar
Jing Zhang committed
393
394
395
396
397
398
    }
    else
    {
        assert(0);
    }
}
Jing Zhang's avatar
Jing Zhang committed
399
400
401
402
403
404
405
406
407
408
409
410
411
412
413
414
415
416
417
418

inline __device__ void global_load(Float4 &r, Float4* ptr) {
	asm volatile("\n \
			global_load_dwordx4 %0, %1, off \n \
			"
			:"=v"(r)
			:"v"(ptr)
			);
}

inline __device__ void ds_write_b128(Float4& r, void* lds, int offset = 0)
{
    asm volatile("\n \
            ds_write_b128 %0, %1 \n \
            "
            :
            : "v"(__to_local(lds)), "v"(r)
            );
}