inline_asm.hpp 10.8 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
#define NO_VM_WAIT 0
#define NO_LGKM_WAIT 0
#define NO_DS_READ 0
#define NO_DS_WRITE 0
#define NO_GLB_READ 0

13
14
inline __device__ void vmcnt(int cnt)
{
Jing Zhang's avatar
Jing Zhang committed
15
#if !NO_VM_WAIT
16
17
18
    if(cnt == 0)
    {
        asm volatile("\n \
Jing Zhang's avatar
Jing Zhang committed
19
                s_waitcnt vmcnt(0) \n \
20
                " ::);
Jing Zhang's avatar
Jing Zhang committed
21
    }
22
23
24
    else if(cnt == 1)
    {
        asm volatile("\n \
Jing Zhang's avatar
Jing Zhang committed
25
                s_waitcnt vmcnt(1) \n \
26
                " ::);
Jing Zhang's avatar
Jing Zhang committed
27
    }
28
29
30
    else if(cnt == 2)
    {
        asm volatile("\n \
Jing Zhang's avatar
Jing Zhang committed
31
                s_waitcnt vmcnt(2) \n \
32
                " ::);
Jing Zhang's avatar
Jing Zhang committed
33
    }
34
35
36
    else if(cnt == 4)
    {
        asm volatile("\n \
Jing Zhang's avatar
Jing Zhang committed
37
                s_waitcnt vmcnt(2) \n \
38
                " ::);
Jing Zhang's avatar
Jing Zhang committed
39
    }
40
41
    else
    {
Jing Zhang's avatar
Jing Zhang committed
42
43
        assert(0);
    }
Jing Zhang's avatar
Jing Zhang committed
44
#endif
Jing Zhang's avatar
Jing Zhang committed
45
46
}

Chao Liu's avatar
Chao Liu committed
47
48
inline __device__ void lgkmcnt(int cnt)
{
Jing Zhang's avatar
Jing Zhang committed
49
#if !NO_LGKM_WAIT
Chao Liu's avatar
Chao Liu committed
50
51
    if(cnt == 0)
    {
Jing Zhang's avatar
Jing Zhang committed
52
53
        asm volatile("\n \
                s_waitcnt lgkmcnt(0) \n \
Chao Liu's avatar
Chao Liu committed
54
                " ::);
Jing Zhang's avatar
Jing Zhang committed
55
    }
Chao Liu's avatar
Chao Liu committed
56
57
    else if(cnt == 1)
    {
Jing Zhang's avatar
Jing Zhang committed
58
59
        asm volatile("\n \
                s_waitcnt lgkmcnt(1) \n \
Chao Liu's avatar
Chao Liu committed
60
                " ::);
Jing Zhang's avatar
Jing Zhang committed
61
    }
Chao Liu's avatar
Chao Liu committed
62
63
    else if(cnt == 2)
    {
Jing Zhang's avatar
Jing Zhang committed
64
65
        asm volatile("\n \
                s_waitcnt lgkmcnt(2) \n \
Chao Liu's avatar
Chao Liu committed
66
                " ::);
Jing Zhang's avatar
Jing Zhang committed
67
    }
Chao Liu's avatar
Chao Liu committed
68
69
    else if(cnt == 3)
    {
Jing Zhang's avatar
Jing Zhang committed
70
71
        asm volatile("\n \
                s_waitcnt lgkmcnt(3) \n \
Chao Liu's avatar
Chao Liu committed
72
                " ::);
Jing Zhang's avatar
Jing Zhang committed
73
    }
Chao Liu's avatar
Chao Liu committed
74
75
    else if(cnt == 4)
    {
Jing Zhang's avatar
Jing Zhang committed
76
77
        asm volatile("\n \
                s_waitcnt lgkmcnt(4) \n \
Chao Liu's avatar
Chao Liu committed
78
                " ::);
Jing Zhang's avatar
Jing Zhang committed
79
    }
Chao Liu's avatar
Chao Liu committed
80
81
    else
    {
Jing Zhang's avatar
Jing Zhang committed
82
        assert(0);
Jing Zhang's avatar
Jing Zhang committed
83
    }
Jing Zhang's avatar
Jing Zhang committed
84
#endif
Jing Zhang's avatar
Jing Zhang committed
85
86
}

Chao Liu's avatar
Chao Liu committed
87
88
inline __device__ void outerProduct1x4(const float* a, const float* b, float* c)
{
Jing Zhang's avatar
Jing Zhang committed
89
90
91
92
93
94
    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
95
96
97
98
99
100
101
102
103
104
                 : "=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
105
106
}

Chao Liu's avatar
Chao Liu committed
107
108
inline __device__ void outerProduct1x4(const float& a, const Float4& b, Float4& c)
{
Jing Zhang's avatar
Jing Zhang committed
109
110
111
112
113
114
115
116
117
118
119
120
121
#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
122
    outerProduct1x4(&a, (float*)&b, (float*)&c);
Jing Zhang's avatar
Jing Zhang committed
123
124
125
#endif
}

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

416
417
inline __device__ void global_load(Float4& r, Float4* ptr)
{
Jing Zhang's avatar
Jing Zhang committed
418
#if !NO_GLB_READ
419
    asm volatile("\n \
Jing Zhang's avatar
Jing Zhang committed
420
421
			global_load_dwordx4 %0, %1, off \n \
			"
422
423
                 : "=v"(r)
                 : "v"(ptr));
Jing Zhang's avatar
Jing Zhang committed
424
#endif
Jing Zhang's avatar
Jing Zhang committed
425
426
427
428
}

inline __device__ void ds_write_b128(Float4& r, void* lds, int offset = 0)
{
Jing Zhang's avatar
Jing Zhang committed
429
#if !NO_DS_WRITE
Jing Zhang's avatar
Jing Zhang committed
430
431
432
    asm volatile("\n \
            ds_write_b128 %0, %1 \n \
            "
433
434
                 :
                 : "v"(__to_local(lds)), "v"(r));
Jing Zhang's avatar
Jing Zhang committed
435
#endif
Jing Zhang's avatar
Jing Zhang committed
436
}