amd_inline_asm.hpp 27.8 KB
Newer Older
1
2
3
#ifndef CK_AMD_INLINE_ASM_HPP
#define CK_AMD_INLINE_ASM_HPP

Chao Liu's avatar
Chao Liu committed
4
#include "vector_type.hpp"
Jing Zhang's avatar
Jing Zhang committed
5

6
7
namespace ck {

Chao Liu's avatar
Chao Liu committed
8
// cast a pointer of LDS to its address
Chao Liu's avatar
Chao Liu committed
9
extern "C" __attribute__((address_space(3))) __device__ void* __to_local(void* p);
Chao Liu's avatar
Chao Liu committed
10

11
// global_load and global_store
12
template <typename T, index_t VectorSize>
13
14
__device__ typename vector_type<T, VectorSize>::MemoryType __global_load(
    const T* p_src_block, uint32_t src_thread_data_offset, uint32_t src_const_data_offset);
15
16

template <typename T, index_t VectorSize>
17
18
__device__ void __global_store(const typename vector_type<T, VectorSize>::MemoryType& src,
                               T* p_dst_block,
19
20
                               uint32_t dst_thread_data_offset,
                               uint32_t dst_const_data_offset);
21
22

template <>
23
__device__ float __global_load<float, 1>(const float* p_src_block,
24
25
                                         uint32_t src_thread_data_offset,
                                         uint32_t src_const_data_offset)
26
27
28
{
    float dst;

29
30
31
32
33
#if 0   // source code
    dst = p_src_block[src_const_data_offset + src_thread_data_offset];
#elif 0 // use VGPR only
    const float* src_thread_addr_offset_u64 =
        p_src_block + src_const_data_offset + src_thread_data_offset;
34
35

    asm volatile("\n \
36
     global_load_dword %0, %1 off offset:0 \n \
37
38
39
     s_waitcnt 0 \n \
     "
                 : "=v"(dst)
40
41
42
43
                 : "v"(src_thread_addr_offset_u64));
#elif 0 // use VGPR and SGPR, do compute on VALU
    uint64_t src_thread_addr_offset_u64 =
        (src_thread_data_offset + src_const_data_offset) * sizeof(float);
44

45
46
47
48
49
50
51
52
53
    asm volatile("\n \
     global_load_dword %0, %1, %2, offset:0 \n \
     s_waitcnt 0 \n \
     "
                 : "=v"(dst)
                 : "v"(src_thread_addr_offset_u64), "s"(p_src_block));
#elif 1 // use VGPR and SGPR, do compute on SALU
    uint64_t src_thread_addr_offset_u64 =
        static_cast<uint64_t>(src_thread_data_offset * sizeof(float));
54

55
    const float* p_src_block_with_offset = p_src_block + src_const_data_offset;
56
57
58

    asm volatile("\n \
     global_load_dword %0, %1, %2, offset:0 \n \
59
     s_waitcnt 0 \n \
60
61
     "
                 : "=v"(dst)
62
63
                 : "v"(src_thread_addr_offset_u64), "s"(p_src_block_with_offset));
#endif
64
65
66
67
68

    return dst;
}

template <>
69
70
__device__ vector_type<float, 2>::MemoryType __global_load<float, 2>(
    const float* p_src_block, uint32_t src_thread_data_offset, uint32_t src_const_data_offset)
71
{
72
73
74
    using vector_t = vector_type<float, 2>::MemoryType;

    vector_t dst;
75

76
77
78
79
80
#if 0   // source code
    dst = *reinterpret_cast<const vector_t*>(&p_src_block[src_const_data_offset + src_thread_data_offset]);
#elif 0 // use VGPR only
    const float* src_thread_addr_offset_u64 =
        p_src_block + src_const_data_offset + src_thread_data_offset;
81
82

    asm volatile("\n \
83
     global_load_dwordx2 %0, %1 off offset:0 \n \
84
85
86
     s_waitcnt 0 \n \
     "
                 : "=v"(dst)
87
88
89
90
                 : "v"(src_thread_addr_offset_u64));
#elif 0 // use VGPR and SGPR, do compute on VALU
    uint64_t src_thread_addr_offset_u64 =
        (src_thread_data_offset + src_const_data_offset) * sizeof(float);
91

92
93
94
95
96
97
98
99
100
    asm volatile("\n \
     global_load_dwordx2 %0, %1, %2, offset:0 \n \
     s_waitcnt 0 \n \
     "
                 : "=v"(dst)
                 : "v"(src_thread_addr_offset_u64), "s"(p_src_block));
#elif 1 // use VGPR and SGPR, do compute on SALU
    uint64_t src_thread_addr_offset_u64 =
        static_cast<uint64_t>(src_thread_data_offset * sizeof(float));
101

102
    const float* p_src_block_with_offset = p_src_block + src_const_data_offset;
103
104
105

    asm volatile("\n \
     global_load_dwordx2 %0, %1, %2, offset:0 \n \
106
     s_waitcnt 0 \n \
107
108
     "
                 : "=v"(dst)
109
110
                 : "v"(src_thread_addr_offset_u64), "s"(p_src_block_with_offset));
#endif
111
112
113
114
115

    return dst;
}

template <>
116
117
__device__ vector_type<float, 4>::MemoryType __global_load<float, 4>(
    const float* p_src_block, uint32_t src_thread_data_offset, uint32_t src_const_data_offset)
118
{
119
120
121
    using vector_t = vector_type<float, 4>::MemoryType;

    vector_t dst;
122

123
124
125
126
127
#if 0   // source code
    dst = *reinterpret_cast<const vector_t*>(&p_src_block[src_const_data_offset + src_thread_data_offset]);
#elif 0 // use VGPR only
    const float* src_thread_addr_offset_u64 =
        p_src_block + src_const_data_offset + src_thread_data_offset;
128
129

    asm volatile("\n \
130
     global_load_dwordx4 %0, %1 off offset:0 \n \
131
132
133
     s_waitcnt 0 \n \
     "
                 : "=v"(dst)
134
135
136
137
                 : "v"(src_thread_addr_offset_u64));
#elif 0 // use VGPR and SGPR, do compute on VALU
    uint64_t src_thread_addr_offset_u64 =
        (src_thread_data_offset + src_const_data_offset) * sizeof(float);
138

139
140
141
142
143
144
145
146
147
    asm volatile("\n \
     global_load_dwordx4 %0, %1, %2, offset:0 \n \
     s_waitcnt 0 \n \
     "
                 : "=v"(dst)
                 : "v"(src_thread_addr_offset_u64), "s"(p_src_block));
#elif 1 // use VGPR and SGPR, do compute on SALU
    uint64_t src_thread_addr_offset_u64 =
        static_cast<uint64_t>(src_thread_data_offset * sizeof(float));
148

149
    const float* p_src_block_with_offset = p_src_block + src_const_data_offset;
150
151
152

    asm volatile("\n \
     global_load_dwordx4 %0, %1, %2, offset:0 \n \
153
     s_waitcnt 0 \n \
154
155
     "
                 : "=v"(dst)
156
157
                 : "v"(src_thread_addr_offset_u64), "s"(p_src_block_with_offset));
#endif
158
159
160
161
162
163
164

    return dst;
}

template <>
__device__ void __global_store<float, 1>(const float& src,
                                         float* p_dst_block,
165
166
                                         uint32_t dst_thread_data_offset,
                                         uint32_t dst_const_data_offset)
167
168
{
#if 0 // compute on VALU
169
    uint64_t dst_thread_data_offset_u64 = (dst_thread_data_offset + dst_const_data_offset) * sizeof(float);
170
171
172
173
174

    asm volatile("\n \
     global_store_dword %0, %1, %2, offset:0 \n \
     "
                 :
175
                 : "v"(dst_thread_data_offset_u64), "v"(src), "s"(p_dst_block));
176
#else // compute on SALU
177
    uint64_t dst_thread_data_offset_u64 = dst_thread_data_offset * sizeof(float);
178

179
    float* p_dst_block_with_offset = p_dst_block + dst_const_data_offset;
180
181
182
183
184

    asm volatile("\n \
     global_store_dword %0, %1, %2, offset:0 \n \
     "
                 :
185
                 : "v"(dst_thread_data_offset_u64), "v"(src), "s"(p_dst_block_with_offset));
186
187
188
189
190
#endif
}

// __buffer_load and __buffer_store
template <typename T, index_t VectorSize>
191
192
__device__ typename vector_type<T, VectorSize>::MemoryType __buffer_load(
    const T* p_src_block, uint32_t src_thread_data_offset, uint32_t src_const_data_offset);
193
194
195
196

template <typename T, index_t VectorSize>
__device__ void __buffer_store(const typename vector_type<T, VectorSize>::MemoryType& src,
                               T* p_dst_block,
197
198
                               uint32_t dst_thread_data_offset,
                               uint32_t dst_const_data_offset);
199
200
201

template <>
__device__ float __buffer_load<float, 1>(const float* p_src_block,
202
203
                                         uint32_t src_thread_data_offset,
                                         uint32_t src_const_data_offset)
204
205
206
{
    float dst;

207
208
209
    uint32_t src_thread_addr_offset = src_thread_data_offset * sizeof(float);
    uint32_t src_const_addr_offset  = src_const_data_offset * sizeof(float);

210
211
212
213
214
215
216
217
218
    int32x4_t src_block_setting{0};
    // fill in byte 0 - 1
    *reinterpret_cast<float**>(&src_block_setting) = const_cast<float*>(p_src_block);
    // fill in byte 2
    reinterpret_cast<int*>(&src_block_setting)[2] = -1;
    // fill in byte 3
    reinterpret_cast<int*>(&src_block_setting)[3] = 0x00027000;

    asm volatile("\n \
219
    buffer_load_dword %0, %1, %2, %3 offen offset:0 \n \
220
    s_waitcnt 0 \n \
221
222
    "
                 : "=v"(dst)
223
                 : "v"(src_thread_addr_offset), "s"(src_block_setting), "s"(src_const_addr_offset));
224
225
226
227
228

    return dst;
}

template <>
229
230
__device__ vector_type<float, 2>::MemoryType __buffer_load<float, 2>(
    const float* p_src_block, uint32_t src_thread_data_offset, uint32_t src_const_data_offset)
231
232
233
{
    vector_type<float, 2>::MemoryType dst;

234
235
236
    uint32_t src_thread_addr_offset = src_thread_data_offset * sizeof(float);
    uint32_t src_const_addr_offset  = src_const_data_offset * sizeof(float);

237
238
239
240
241
242
243
244
245
    int32x4_t src_block_setting{0};
    // fill in byte 0 - 1
    *reinterpret_cast<float**>(&src_block_setting) = const_cast<float*>(p_src_block);
    // fill in byte 2
    reinterpret_cast<int*>(&src_block_setting)[2] = -1;
    // fill in byte 3
    reinterpret_cast<int*>(&src_block_setting)[3] = 0x00027000;

    asm volatile("\n \
246
    buffer_load_dwordx2 %0, %1, %2, %3 offen offset:0 \n \
247
    s_waitcnt 0 \n \
248
249
    "
                 : "=v"(dst)
250
                 : "v"(src_thread_addr_offset), "s"(src_block_setting), "s"(src_const_addr_offset));
251
252
253
254
255

    return dst;
}

template <>
256
257
__device__ vector_type<float, 4>::MemoryType __buffer_load<float, 4>(
    const float* p_src_block, uint32_t src_thread_data_offset, uint32_t src_const_data_offset)
258
259
260
{
    vector_type<float, 4>::MemoryType dst;

261
262
263
    uint32_t src_thread_addr_offset = src_thread_data_offset * sizeof(float);
    uint32_t src_const_addr_offset  = src_const_data_offset * sizeof(float);

264
265
266
267
268
269
270
271
272
    int32x4_t src_block_setting{0};
    // fill in byte 0 - 1
    *reinterpret_cast<float**>(&src_block_setting) = const_cast<float*>(p_src_block);
    // fill in byte 2
    reinterpret_cast<int*>(&src_block_setting)[2] = -1;
    // fill in byte 3
    reinterpret_cast<int*>(&src_block_setting)[3] = 0x00027000;

    asm volatile("\n \
273
    buffer_load_dwordx4 %0, %1, %2, %3 offen offset:0 \n \
274
    s_waitcnt 0 \n \
275
276
    "
                 : "=v"(dst)
277
                 : "v"(src_thread_addr_offset), "s"(src_block_setting), "s"(src_const_addr_offset));
278
279
280
281
282

    return dst;
}

template <>
283
284
__device__ void __buffer_store<float, 1>(const float& src,
                                         float* p_dst_block,
285
286
                                         uint32_t dst_thread_data_offset,
                                         uint32_t dst_const_data_offset)
287
{
288
289
290
    uint32_t dst_thread_addr_offset = dst_thread_data_offset * sizeof(float);
    uint32_t dst_const_addr_offset  = dst_const_data_offset * sizeof(float);

291
292
293
294
295
296
297
298
299
    int32x4_t dst_block_setting{0};
    // fill in byte 0 - 1
    *reinterpret_cast<float**>(&dst_block_setting) = p_dst_block;
    // fill in byte 2
    reinterpret_cast<int*>(&dst_block_setting)[2] = -1;
    // fill in byte 3
    reinterpret_cast<int*>(&dst_block_setting)[3] = 0x00027000;

    asm volatile("\n \
300
    buffer_store_dword %1, %2, %0, %3 offen offset:0 \n \
301
302
    "
                 :
303
304
305
306
                 : "s"(dst_block_setting),
                   "v"(src),
                   "v"(dst_thread_addr_offset),
                   "s"(dst_const_addr_offset));
307
308
}

Chao Liu's avatar
Chao Liu committed
309
310
311
312
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
342
343
344
345
346
347
348
349
350
351
352
353
354
355
356
357
358
359
360
361
362
363
364
365
366
367
368
369
370
371
372
373
374
375
376
377
378
__device__ void vmcnt(index_t 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(false);
    }
}

__device__ void lgkmcnt(index_t cnt)
{
    if(cnt == 0)
    {
        asm volatile("\n \
                s_waitcnt lgkmcnt(0) \n \
                " ::);
    }
    else if(cnt == 1)
    {
        asm volatile("\n \
                s_waitcnt lgkmcnt(1) \n \
                " ::);
    }
    else if(cnt == 2)
    {
        asm volatile("\n \
                s_waitcnt lgkmcnt(2) \n \
                " ::);
    }
    else if(cnt == 3)
    {
        asm volatile("\n \
                s_waitcnt lgkmcnt(3) \n \
                " ::);
    }
    else if(cnt == 4)
    {
        asm volatile("\n \
                s_waitcnt lgkmcnt(4) \n \
                " ::);
    }
    else
    {
        assert(false);
    }
}

Chao Liu's avatar
Chao Liu committed
379
__device__ void outerProduct1x4(const float* a, const float* b, float* c)
Chao Liu's avatar
Chao Liu committed
380
{
Jing Zhang's avatar
Jing Zhang committed
381
382
383
384
385
386
    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
387
388
389
390
391
392
393
394
395
396
                 : "=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
397
398
}

Chao Liu's avatar
Chao Liu committed
399
400
401
__device__ void outerProduct1x4(const float& a,
                                const vector_type<float, 4>::MemoryType& b,
                                vector_type<float, 4>::MemoryType& c)
Chao Liu's avatar
Chao Liu committed
402
{
Chao Liu's avatar
Chao Liu committed
403
    outerProduct1x4(&a, reinterpret_cast<const float*>(&b), reinterpret_cast<float*>(&c));
Jing Zhang's avatar
Jing Zhang committed
404
405
}

Chao Liu's avatar
Chao Liu committed
406
407
408
409
410
411
412
413
414
__device__ void outerProduct2x4(const vector_type<float, 2>::MemoryType& a,
                                const vector_type<float, 4>::MemoryType& b,
                                vector_type<float, 4>::MemoryType& c0,
                                vector_type<float, 4>::MemoryType& c1)
{
    outerProduct1x4(a.x, b, c0);
    outerProduct1x4(a.y, b, c1);
}

Chao Liu's avatar
Chao Liu committed
415
416
417
418
419
420
__device__ void outerProduct4x4(const vector_type<float, 4>::MemoryType& a,
                                const vector_type<float, 4>::MemoryType& b,
                                vector_type<float, 4>::MemoryType& c0,
                                vector_type<float, 4>::MemoryType& c1,
                                vector_type<float, 4>::MemoryType& c2,
                                vector_type<float, 4>::MemoryType& c3)
Chao Liu's avatar
Chao Liu committed
421
{
Jing Zhang's avatar
Jing Zhang committed
422
423
424
425
426
427
    outerProduct1x4(a.x, b, c0);
    outerProduct1x4(a.y, b, c1);
    outerProduct1x4(a.z, b, c2);
    outerProduct1x4(a.w, b, c3);
}

Chao Liu's avatar
Chao Liu committed
428
429
430
__device__ void outerProduct8x8(const vector_type<float, 4>::MemoryType* a,
                                const vector_type<float, 4>::MemoryType* b,
                                vector_type<float, 4>::MemoryType* c)
Jing Zhang's avatar
Jing Zhang committed
431
432
433
434
435
436
437
{
    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
438
__device__ void ds_read_b128(vector_type<float, 4>::MemoryType& r, void* lds, index_t offset = 0)
Jing Zhang's avatar
Jing Zhang committed
439
440
441
442
{
    if(offset == 0)
    {
        asm volatile("\n \
Jing Zhang's avatar
Jing Zhang committed
443
                ds_read_b128 %0, %1 offset:0\n \
Jing Zhang's avatar
Jing Zhang committed
444
                "
Chao Liu's avatar
Chao Liu committed
445
446
                     : "=v"(r)
                     : "v"(__to_local(lds)));
Jing Zhang's avatar
Jing Zhang committed
447
    }
Jing Zhang's avatar
Jing Zhang committed
448
    if(offset == 64)
Jing Zhang's avatar
Jing Zhang committed
449
450
    {
        asm volatile("\n \
Jing Zhang's avatar
Jing Zhang committed
451
                ds_read_b128 %0, %1 offset:64\n \
Jing Zhang's avatar
Jing Zhang committed
452
                "
Chao Liu's avatar
Chao Liu committed
453
454
                     : "=v"(r)
                     : "v"(__to_local(lds)));
Jing Zhang's avatar
Jing Zhang committed
455
    }
Jing Zhang's avatar
Jing Zhang committed
456
    if(offset == 128)
Jing Zhang's avatar
Jing Zhang committed
457
458
    {
        asm volatile("\n \
Jing Zhang's avatar
Jing Zhang committed
459
                ds_read_b128 %0, %1 offset:128\n \
Jing Zhang's avatar
Jing Zhang committed
460
                "
Chao Liu's avatar
Chao Liu committed
461
462
                     : "=v"(r)
                     : "v"(__to_local(lds)));
Jing Zhang's avatar
Jing Zhang committed
463
    }
Jing Zhang's avatar
Jing Zhang committed
464
    if(offset == 192)
Jing Zhang's avatar
Jing Zhang committed
465
466
    {
        asm volatile("\n \
Jing Zhang's avatar
Jing Zhang committed
467
                ds_read_b128 %0, %1 offset:192\n \
Jing Zhang's avatar
Jing Zhang committed
468
                "
Chao Liu's avatar
Chao Liu committed
469
470
                     : "=v"(r)
                     : "v"(__to_local(lds)));
Jing Zhang's avatar
Jing Zhang committed
471
    }
Jing Zhang's avatar
Jing Zhang committed
472
    if(offset == 256)
Jing Zhang's avatar
Jing Zhang committed
473
474
    {
        asm volatile("\n \
Jing Zhang's avatar
Jing Zhang committed
475
                ds_read_b128 %0, %1 offset:256\n \
Jing Zhang's avatar
Jing Zhang committed
476
                "
Chao Liu's avatar
Chao Liu committed
477
478
                     : "=v"(r)
                     : "v"(__to_local(lds)));
Jing Zhang's avatar
Jing Zhang committed
479
    }
Jing Zhang's avatar
Jing Zhang committed
480
    if(offset == 320)
Jing Zhang's avatar
Jing Zhang committed
481
482
    {
        asm volatile("\n \
Jing Zhang's avatar
Jing Zhang committed
483
                ds_read_b128 %0, %1 offset:320\n \
Jing Zhang's avatar
Jing Zhang committed
484
                "
Chao Liu's avatar
Chao Liu committed
485
486
                     : "=v"(r)
                     : "v"(__to_local(lds)));
Jing Zhang's avatar
Jing Zhang committed
487
    }
Jing Zhang's avatar
Jing Zhang committed
488
    if(offset == 384)
Jing Zhang's avatar
Jing Zhang committed
489
490
    {
        asm volatile("\n \
Jing Zhang's avatar
Jing Zhang committed
491
                ds_read_b128 %0, %1 offset:384\n \
Jing Zhang's avatar
Jing Zhang committed
492
                "
Chao Liu's avatar
Chao Liu committed
493
494
                     : "=v"(r)
                     : "v"(__to_local(lds)));
Jing Zhang's avatar
Jing Zhang committed
495
    }
Jing Zhang's avatar
Jing Zhang committed
496
    if(offset == 448)
Jing Zhang's avatar
Jing Zhang committed
497
498
    {
        asm volatile("\n \
Jing Zhang's avatar
Jing Zhang committed
499
                ds_read_b128 %0, %1 offset:448\n \
Jing Zhang's avatar
Jing Zhang committed
500
                "
Chao Liu's avatar
Chao Liu committed
501
502
                     : "=v"(r)
                     : "v"(__to_local(lds)));
Jing Zhang's avatar
Jing Zhang committed
503
    }
Jing Zhang's avatar
Jing Zhang committed
504
    if(offset == 512)
Jing Zhang's avatar
Jing Zhang committed
505
506
    {
        asm volatile("\n \
Jing Zhang's avatar
Jing Zhang committed
507
                ds_read_b128 %0, %1 offset:512\n \
Jing Zhang's avatar
Jing Zhang committed
508
                "
Chao Liu's avatar
Chao Liu committed
509
510
                     : "=v"(r)
                     : "v"(__to_local(lds)));
Jing Zhang's avatar
Jing Zhang committed
511
    }
Jing Zhang's avatar
Jing Zhang committed
512
    if(offset == 576)
Jing Zhang's avatar
Jing Zhang committed
513
514
    {
        asm volatile("\n \
Jing Zhang's avatar
Jing Zhang committed
515
                ds_read_b128 %0, %1 offset:576\n \
Jing Zhang's avatar
Jing Zhang committed
516
                "
Chao Liu's avatar
Chao Liu committed
517
518
                     : "=v"(r)
                     : "v"(__to_local(lds)));
Jing Zhang's avatar
Jing Zhang committed
519
    }
Jing Zhang's avatar
Jing Zhang committed
520
    if(offset == 640)
Jing Zhang's avatar
Jing Zhang committed
521
522
    {
        asm volatile("\n \
Jing Zhang's avatar
Jing Zhang committed
523
                ds_read_b128 %0, %1 offset:640\n \
Jing Zhang's avatar
Jing Zhang committed
524
                "
Chao Liu's avatar
Chao Liu committed
525
526
                     : "=v"(r)
                     : "v"(__to_local(lds)));
Jing Zhang's avatar
Jing Zhang committed
527
    }
Jing Zhang's avatar
Jing Zhang committed
528
    if(offset == 704)
Jing Zhang's avatar
Jing Zhang committed
529
530
    {
        asm volatile("\n \
Jing Zhang's avatar
Jing Zhang committed
531
                ds_read_b128 %0, %1 offset:704\n \
Jing Zhang's avatar
Jing Zhang committed
532
                "
Chao Liu's avatar
Chao Liu committed
533
534
                     : "=v"(r)
                     : "v"(__to_local(lds)));
Jing Zhang's avatar
Jing Zhang committed
535
    }
Jing Zhang's avatar
Jing Zhang committed
536
    if(offset == 768)
Jing Zhang's avatar
Jing Zhang committed
537
538
    {
        asm volatile("\n \
Jing Zhang's avatar
Jing Zhang committed
539
                ds_read_b128 %0, %1 offset:768\n \
Jing Zhang's avatar
Jing Zhang committed
540
                "
Chao Liu's avatar
Chao Liu committed
541
542
                     : "=v"(r)
                     : "v"(__to_local(lds)));
Jing Zhang's avatar
Jing Zhang committed
543
    }
Jing Zhang's avatar
Jing Zhang committed
544
    if(offset == 832)
Jing Zhang's avatar
Jing Zhang committed
545
546
    {
        asm volatile("\n \
Jing Zhang's avatar
Jing Zhang committed
547
                ds_read_b128 %0, %1 offset:832\n \
Jing Zhang's avatar
Jing Zhang committed
548
                "
Chao Liu's avatar
Chao Liu committed
549
550
                     : "=v"(r)
                     : "v"(__to_local(lds)));
Jing Zhang's avatar
Jing Zhang committed
551
    }
Jing Zhang's avatar
Jing Zhang committed
552
    if(offset == 896)
Jing Zhang's avatar
Jing Zhang committed
553
554
    {
        asm volatile("\n \
Jing Zhang's avatar
Jing Zhang committed
555
                ds_read_b128 %0, %1 offset:896\n \
Jing Zhang's avatar
Jing Zhang committed
556
                "
Chao Liu's avatar
Chao Liu committed
557
558
                     : "=v"(r)
                     : "v"(__to_local(lds)));
Jing Zhang's avatar
Jing Zhang committed
559
    }
Jing Zhang's avatar
Jing Zhang committed
560
    if(offset == 960)
Jing Zhang's avatar
Jing Zhang committed
561
562
    {
        asm volatile("\n \
Jing Zhang's avatar
Jing Zhang committed
563
                ds_read_b128 %0, %1 offset:960\n \
Jing Zhang's avatar
Jing Zhang committed
564
                "
Chao Liu's avatar
Chao Liu committed
565
566
                     : "=v"(r)
                     : "v"(__to_local(lds)));
Jing Zhang's avatar
Jing Zhang committed
567
    }
Jing Zhang's avatar
Jing Zhang committed
568
    if(offset == 1024)
Jing Zhang's avatar
Jing Zhang committed
569
570
    {
        asm volatile("\n \
Jing Zhang's avatar
Jing Zhang committed
571
                ds_read_b128 %0, %1 offset:1024\n \
Jing Zhang's avatar
Jing Zhang committed
572
                "
Chao Liu's avatar
Chao Liu committed
573
574
                     : "=v"(r)
                     : "v"(__to_local(lds)));
Jing Zhang's avatar
Jing Zhang committed
575
    }
Jing Zhang's avatar
Jing Zhang committed
576
    if(offset == 1088)
Jing Zhang's avatar
Jing Zhang committed
577
578
    {
        asm volatile("\n \
Jing Zhang's avatar
Jing Zhang committed
579
                ds_read_b128 %0, %1 offset:1088\n \
Jing Zhang's avatar
Jing Zhang committed
580
                "
Chao Liu's avatar
Chao Liu committed
581
582
                     : "=v"(r)
                     : "v"(__to_local(lds)));
Jing Zhang's avatar
Jing Zhang committed
583
    }
Jing Zhang's avatar
Jing Zhang committed
584
    if(offset == 1152)
Jing Zhang's avatar
Jing Zhang committed
585
586
    {
        asm volatile("\n \
Jing Zhang's avatar
Jing Zhang committed
587
                ds_read_b128 %0, %1 offset:1152\n \
Jing Zhang's avatar
Jing Zhang committed
588
                "
Chao Liu's avatar
Chao Liu committed
589
590
                     : "=v"(r)
                     : "v"(__to_local(lds)));
Jing Zhang's avatar
Jing Zhang committed
591
    }
Jing Zhang's avatar
Jing Zhang committed
592
    if(offset == 1216)
Chao Liu's avatar
Chao Liu committed
593
594
    {
        asm volatile("\n \
Jing Zhang's avatar
Jing Zhang committed
595
                ds_read_b128 %0, %1 offset:1216\n \
Chao Liu's avatar
Chao Liu committed
596
                "
Chao Liu's avatar
Chao Liu committed
597
598
                     : "=v"(r)
                     : "v"(__to_local(lds)));
Chao Liu's avatar
Chao Liu committed
599
    }
Jing Zhang's avatar
Jing Zhang committed
600
    if(offset == 1280)
Jing Zhang's avatar
Jing Zhang committed
601
602
    {
        asm volatile("\n \
Jing Zhang's avatar
Jing Zhang committed
603
                ds_read_b128 %0, %1 offset:1280\n \
Jing Zhang's avatar
Jing Zhang committed
604
                "
Chao Liu's avatar
Chao Liu committed
605
606
                     : "=v"(r)
                     : "v"(__to_local(lds)));
Jing Zhang's avatar
Jing Zhang committed
607
    }
Jing Zhang's avatar
Jing Zhang committed
608
    if(offset == 1344)
Chao Liu's avatar
Chao Liu committed
609
610
    {
        asm volatile("\n \
Jing Zhang's avatar
Jing Zhang committed
611
                ds_read_b128 %0, %1 offset:1344\n \
Chao Liu's avatar
Chao Liu committed
612
                "
Chao Liu's avatar
Chao Liu committed
613
614
                     : "=v"(r)
                     : "v"(__to_local(lds)));
Chao Liu's avatar
Chao Liu committed
615
    }
Jing Zhang's avatar
Jing Zhang committed
616
    if(offset == 1408)
Jing Zhang's avatar
Jing Zhang committed
617
618
    {
        asm volatile("\n \
Jing Zhang's avatar
Jing Zhang committed
619
                ds_read_b128 %0, %1 offset:1408\n \
Jing Zhang's avatar
Jing Zhang committed
620
                "
Chao Liu's avatar
Chao Liu committed
621
622
                     : "=v"(r)
                     : "v"(__to_local(lds)));
Jing Zhang's avatar
Jing Zhang committed
623
    }
Jing Zhang's avatar
Jing Zhang committed
624
    if(offset == 1472)
Chao Liu's avatar
Chao Liu committed
625
626
    {
        asm volatile("\n \
Jing Zhang's avatar
Jing Zhang committed
627
                ds_read_b128 %0, %1 offset:1472\n \
Chao Liu's avatar
Chao Liu committed
628
                "
Chao Liu's avatar
Chao Liu committed
629
630
                     : "=v"(r)
                     : "v"(__to_local(lds)));
Chao Liu's avatar
Chao Liu committed
631
    }
Jing Zhang's avatar
Jing Zhang committed
632
    if(offset == 1536)
Jing Zhang's avatar
Jing Zhang committed
633
634
    {
        asm volatile("\n \
Jing Zhang's avatar
Jing Zhang committed
635
                ds_read_b128 %0, %1 offset:1536\n \
Jing Zhang's avatar
Jing Zhang committed
636
                "
Chao Liu's avatar
Chao Liu committed
637
638
                     : "=v"(r)
                     : "v"(__to_local(lds)));
Jing Zhang's avatar
Jing Zhang committed
639
    }
Jing Zhang's avatar
Jing Zhang committed
640
    if(offset == 1600)
Chao Liu's avatar
Chao Liu committed
641
642
    {
        asm volatile("\n \
Jing Zhang's avatar
Jing Zhang committed
643
                ds_read_b128 %0, %1 offset:1600\n \
Chao Liu's avatar
Chao Liu committed
644
                "
Chao Liu's avatar
Chao Liu committed
645
646
                     : "=v"(r)
                     : "v"(__to_local(lds)));
Chao Liu's avatar
Chao Liu committed
647
    }
Jing Zhang's avatar
Jing Zhang committed
648
    if(offset == 1664)
Jing Zhang's avatar
Jing Zhang committed
649
650
    {
        asm volatile("\n \
Jing Zhang's avatar
Jing Zhang committed
651
                ds_read_b128 %0, %1 offset:1664\n \
Jing Zhang's avatar
Jing Zhang committed
652
                "
Chao Liu's avatar
Chao Liu committed
653
654
                     : "=v"(r)
                     : "v"(__to_local(lds)));
Jing Zhang's avatar
Jing Zhang committed
655
    }
Jing Zhang's avatar
Jing Zhang committed
656
    if(offset == 1728)
Chao Liu's avatar
Chao Liu committed
657
658
    {
        asm volatile("\n \
Jing Zhang's avatar
Jing Zhang committed
659
                ds_read_b128 %0, %1 offset:1728\n \
Chao Liu's avatar
Chao Liu committed
660
                "
Chao Liu's avatar
Chao Liu committed
661
662
                     : "=v"(r)
                     : "v"(__to_local(lds)));
Chao Liu's avatar
Chao Liu committed
663
    }
Jing Zhang's avatar
Jing Zhang committed
664
    if(offset == 1792)
Jing Zhang's avatar
Jing Zhang committed
665
666
    {
        asm volatile("\n \
Jing Zhang's avatar
Jing Zhang committed
667
                ds_read_b128 %0, %1 offset:1792\n \
Jing Zhang's avatar
Jing Zhang committed
668
                "
Chao Liu's avatar
Chao Liu committed
669
670
                     : "=v"(r)
                     : "v"(__to_local(lds)));
Jing Zhang's avatar
Jing Zhang committed
671
    }
Jing Zhang's avatar
Jing Zhang committed
672
    if(offset == 1856)
Chao Liu's avatar
Chao Liu committed
673
674
    {
        asm volatile("\n \
Jing Zhang's avatar
Jing Zhang committed
675
                ds_read_b128 %0, %1 offset:1856\n \
Chao Liu's avatar
Chao Liu committed
676
                "
Chao Liu's avatar
Chao Liu committed
677
678
                     : "=v"(r)
                     : "v"(__to_local(lds)));
Chao Liu's avatar
Chao Liu committed
679
    }
Jing Zhang's avatar
Jing Zhang committed
680
    if(offset == 1920)
Jing Zhang's avatar
Jing Zhang committed
681
682
    {
        asm volatile("\n \
Jing Zhang's avatar
Jing Zhang committed
683
                ds_read_b128 %0, %1 offset:1920\n \
Jing Zhang's avatar
Jing Zhang committed
684
                "
Chao Liu's avatar
Chao Liu committed
685
686
                     : "=v"(r)
                     : "v"(__to_local(lds)));
Jing Zhang's avatar
Jing Zhang committed
687
    }
Jing Zhang's avatar
Jing Zhang committed
688
    if(offset == 1984)
Chao Liu's avatar
Chao Liu committed
689
690
    {
        asm volatile("\n \
Jing Zhang's avatar
Jing Zhang committed
691
                ds_read_b128 %0, %1 offset:1984\n \
Chao Liu's avatar
Chao Liu committed
692
                "
Chao Liu's avatar
Chao Liu committed
693
694
                     : "=v"(r)
                     : "v"(__to_local(lds)));
Chao Liu's avatar
Chao Liu committed
695
    }
Jing Zhang's avatar
Jing Zhang committed
696
    if(offset == 2048)
Jing Zhang's avatar
Jing Zhang committed
697
698
    {
        asm volatile("\n \
Jing Zhang's avatar
Jing Zhang committed
699
                ds_read_b128 %0, %1 offset:2048\n \
Jing Zhang's avatar
Jing Zhang committed
700
                "
Chao Liu's avatar
Chao Liu committed
701
702
                     : "=v"(r)
                     : "v"(__to_local(lds)));
Jing Zhang's avatar
Jing Zhang committed
703
    }
Jing Zhang's avatar
Jing Zhang committed
704
    if(offset == 2112)
Jing Zhang's avatar
Jing Zhang committed
705
706
    {
        asm volatile("\n \
Jing Zhang's avatar
Jing Zhang committed
707
                ds_read_b128 %0, %1 offset:2112\n \
Jing Zhang's avatar
Jing Zhang committed
708
                "
Chao Liu's avatar
Chao Liu committed
709
710
                     : "=v"(r)
                     : "v"(__to_local(lds)));
Jing Zhang's avatar
Jing Zhang committed
711
    }
Jing Zhang's avatar
Jing Zhang committed
712
    if(offset == 2176)
Jing Zhang's avatar
Jing Zhang committed
713
    {
Jing Zhang's avatar
Jing Zhang committed
714
715
716
        asm volatile("\n \
                ds_read_b128 %0, %1 offset:2176\n \
                "
Chao Liu's avatar
Chao Liu committed
717
718
                     : "=v"(r)
                     : "v"(__to_local(lds)));
Jing Zhang's avatar
Jing Zhang committed
719
720
721
722
723
724
    }
    if(offset == 2240)
    {
        asm volatile("\n \
                ds_read_b128 %0, %1 offset:2240\n \
                "
Chao Liu's avatar
Chao Liu committed
725
726
                     : "=v"(r)
                     : "v"(__to_local(lds)));
Jing Zhang's avatar
Jing Zhang committed
727
728
729
730
731
732
    }
    if(offset == 2304)
    {
        asm volatile("\n \
                ds_read_b128 %0, %1 offset:2304\n \
                "
Chao Liu's avatar
Chao Liu committed
733
734
                     : "=v"(r)
                     : "v"(__to_local(lds)));
Jing Zhang's avatar
Jing Zhang committed
735
736
737
738
739
740
    }
    if(offset == 2368)
    {
        asm volatile("\n \
                ds_read_b128 %0, %1 offset:2368\n \
                "
Chao Liu's avatar
Chao Liu committed
741
742
                     : "=v"(r)
                     : "v"(__to_local(lds)));
Jing Zhang's avatar
Jing Zhang committed
743
744
745
746
747
748
    }
    if(offset == 2432)
    {
        asm volatile("\n \
                ds_read_b128 %0, %1 offset:2432\n \
                "
Chao Liu's avatar
Chao Liu committed
749
750
                     : "=v"(r)
                     : "v"(__to_local(lds)));
Jing Zhang's avatar
Jing Zhang committed
751
752
753
754
755
756
    }
    if(offset == 2496)
    {
        asm volatile("\n \
                ds_read_b128 %0, %1 offset:2496\n \
                "
Chao Liu's avatar
Chao Liu committed
757
758
                     : "=v"(r)
                     : "v"(__to_local(lds)));
Jing Zhang's avatar
Jing Zhang committed
759
760
761
762
763
764
    }
    if(offset == 2560)
    {
        asm volatile("\n \
                ds_read_b128 %0, %1 offset:2560\n \
                "
Chao Liu's avatar
Chao Liu committed
765
766
                     : "=v"(r)
                     : "v"(__to_local(lds)));
Jing Zhang's avatar
Jing Zhang committed
767
768
769
770
771
772
    }
    if(offset == 2624)
    {
        asm volatile("\n \
                ds_read_b128 %0, %1 offset:2624\n \
                "
Chao Liu's avatar
Chao Liu committed
773
774
                     : "=v"(r)
                     : "v"(__to_local(lds)));
Jing Zhang's avatar
Jing Zhang committed
775
776
777
778
779
780
    }
    if(offset == 2688)
    {
        asm volatile("\n \
                ds_read_b128 %0, %1 offset:2688\n \
                "
Chao Liu's avatar
Chao Liu committed
781
782
                     : "=v"(r)
                     : "v"(__to_local(lds)));
Jing Zhang's avatar
Jing Zhang committed
783
784
785
786
787
788
    }
    if(offset == 2752)
    {
        asm volatile("\n \
                ds_read_b128 %0, %1 offset:2752\n \
                "
Chao Liu's avatar
Chao Liu committed
789
790
                     : "=v"(r)
                     : "v"(__to_local(lds)));
Jing Zhang's avatar
Jing Zhang committed
791
792
793
794
795
796
    }
    if(offset == 2816)
    {
        asm volatile("\n \
                ds_read_b128 %0, %1 offset:2816\n \
                "
Chao Liu's avatar
Chao Liu committed
797
798
                     : "=v"(r)
                     : "v"(__to_local(lds)));
Jing Zhang's avatar
Jing Zhang committed
799
800
801
802
803
804
    }
    if(offset == 2880)
    {
        asm volatile("\n \
                ds_read_b128 %0, %1 offset:2880\n \
                "
Chao Liu's avatar
Chao Liu committed
805
806
                     : "=v"(r)
                     : "v"(__to_local(lds)));
Jing Zhang's avatar
Jing Zhang committed
807
808
809
810
811
812
    }
    if(offset == 2944)
    {
        asm volatile("\n \
                ds_read_b128 %0, %1 offset:2944\n \
                "
Chao Liu's avatar
Chao Liu committed
813
814
                     : "=v"(r)
                     : "v"(__to_local(lds)));
Jing Zhang's avatar
Jing Zhang committed
815
816
817
818
819
820
    }
    if(offset == 3008)
    {
        asm volatile("\n \
                ds_read_b128 %0, %1 offset:3008\n \
                "
Chao Liu's avatar
Chao Liu committed
821
822
                     : "=v"(r)
                     : "v"(__to_local(lds)));
Jing Zhang's avatar
Jing Zhang committed
823
824
825
826
827
828
    }
    if(offset == 3072)
    {
        asm volatile("\n \
                ds_read_b128 %0, %1 offset:3072\n \
                "
Chao Liu's avatar
Chao Liu committed
829
830
                     : "=v"(r)
                     : "v"(__to_local(lds)));
Jing Zhang's avatar
Jing Zhang committed
831
832
833
834
835
836
    }
    if(offset == 3136)
    {
        asm volatile("\n \
                ds_read_b128 %0, %1 offset:3136\n \
                "
Chao Liu's avatar
Chao Liu committed
837
838
                     : "=v"(r)
                     : "v"(__to_local(lds)));
Jing Zhang's avatar
Jing Zhang committed
839
840
841
842
843
844
    }
    if(offset == 3200)
    {
        asm volatile("\n \
                ds_read_b128 %0, %1 offset:3200\n \
                "
Chao Liu's avatar
Chao Liu committed
845
846
                     : "=v"(r)
                     : "v"(__to_local(lds)));
Jing Zhang's avatar
Jing Zhang committed
847
848
849
850
851
852
    }
    if(offset == 3264)
    {
        asm volatile("\n \
                ds_read_b128 %0, %1 offset:3264\n \
                "
Chao Liu's avatar
Chao Liu committed
853
854
                     : "=v"(r)
                     : "v"(__to_local(lds)));
Jing Zhang's avatar
Jing Zhang committed
855
856
857
858
859
860
    }
    if(offset == 3328)
    {
        asm volatile("\n \
                ds_read_b128 %0, %1 offset:3328\n \
                "
Chao Liu's avatar
Chao Liu committed
861
862
                     : "=v"(r)
                     : "v"(__to_local(lds)));
Jing Zhang's avatar
Jing Zhang committed
863
864
865
866
867
868
    }
    if(offset == 3392)
    {
        asm volatile("\n \
                ds_read_b128 %0, %1 offset:3392\n \
                "
Chao Liu's avatar
Chao Liu committed
869
870
                     : "=v"(r)
                     : "v"(__to_local(lds)));
Jing Zhang's avatar
Jing Zhang committed
871
872
873
874
875
876
    }
    if(offset == 3456)
    {
        asm volatile("\n \
                ds_read_b128 %0, %1 offset:3456\n \
                "
Chao Liu's avatar
Chao Liu committed
877
878
                     : "=v"(r)
                     : "v"(__to_local(lds)));
Jing Zhang's avatar
Jing Zhang committed
879
880
881
882
883
884
    }
    if(offset == 3520)
    {
        asm volatile("\n \
                ds_read_b128 %0, %1 offset:3520\n \
                "
Chao Liu's avatar
Chao Liu committed
885
886
                     : "=v"(r)
                     : "v"(__to_local(lds)));
Jing Zhang's avatar
Jing Zhang committed
887
888
889
890
891
892
    }
    if(offset == 3584)
    {
        asm volatile("\n \
                ds_read_b128 %0, %1 offset:3584\n \
                "
Chao Liu's avatar
Chao Liu committed
893
894
                     : "=v"(r)
                     : "v"(__to_local(lds)));
Jing Zhang's avatar
Jing Zhang committed
895
896
897
898
899
900
    }
    if(offset == 3648)
    {
        asm volatile("\n \
                ds_read_b128 %0, %1 offset:3648\n \
                "
Chao Liu's avatar
Chao Liu committed
901
902
                     : "=v"(r)
                     : "v"(__to_local(lds)));
Jing Zhang's avatar
Jing Zhang committed
903
904
905
906
907
908
    }
    if(offset == 3712)
    {
        asm volatile("\n \
                ds_read_b128 %0, %1 offset:3712\n \
                "
Chao Liu's avatar
Chao Liu committed
909
910
                     : "=v"(r)
                     : "v"(__to_local(lds)));
Jing Zhang's avatar
Jing Zhang committed
911
912
913
914
915
916
    }
    if(offset == 3776)
    {
        asm volatile("\n \
                ds_read_b128 %0, %1 offset:3776\n \
                "
Chao Liu's avatar
Chao Liu committed
917
918
                     : "=v"(r)
                     : "v"(__to_local(lds)));
Jing Zhang's avatar
Jing Zhang committed
919
920
921
922
923
924
    }
    if(offset == 3840)
    {
        asm volatile("\n \
                ds_read_b128 %0, %1 offset:3840\n \
                "
Chao Liu's avatar
Chao Liu committed
925
926
                     : "=v"(r)
                     : "v"(__to_local(lds)));
Jing Zhang's avatar
Jing Zhang committed
927
928
929
930
931
932
    }
    if(offset == 3904)
    {
        asm volatile("\n \
                ds_read_b128 %0, %1 offset:3904\n \
                "
Chao Liu's avatar
Chao Liu committed
933
934
                     : "=v"(r)
                     : "v"(__to_local(lds)));
Jing Zhang's avatar
Jing Zhang committed
935
936
937
938
939
940
    }
    if(offset == 3968)
    {
        asm volatile("\n \
                ds_read_b128 %0, %1 offset:3968\n \
                "
Chao Liu's avatar
Chao Liu committed
941
942
                     : "=v"(r)
                     : "v"(__to_local(lds)));
Jing Zhang's avatar
Jing Zhang committed
943
944
945
946
947
948
    }
    if(offset == 4032)
    {
        asm volatile("\n \
                ds_read_b128 %0, %1 offset:4032\n \
                "
Chao Liu's avatar
Chao Liu committed
949
950
                     : "=v"(r)
                     : "v"(__to_local(lds)));
Jing Zhang's avatar
Jing Zhang committed
951
952
953
954
955
956
    }
    if(offset == 4096)
    {
        asm volatile("\n \
                ds_read_b128 %0, %1 offset:4096\n \
                "
Chao Liu's avatar
Chao Liu committed
957
958
                     : "=v"(r)
                     : "v"(__to_local(lds)));
Jing Zhang's avatar
Jing Zhang committed
959
    }
Jing Zhang's avatar
Jing Zhang committed
960
961
}

Chao Liu's avatar
Chao Liu committed
962
963
__device__ void
ds_write_b128(const vector_type<float, 4>::MemoryType& r, void* lds, index_t offset = 0)
Jing Zhang's avatar
Jing Zhang committed
964
{
965
966
967
    if(offset == 0)
    {
        asm volatile("\n \
Jing Zhang's avatar
Jing Zhang committed
968
969
            ds_write_b128 %0, %1 \n \
            "
970
971
972
973
974
975
976
                     :
                     : "v"(__to_local(lds)), "v"(r));
    }
    else
    {
        assert(false);
    }
Jing Zhang's avatar
Jing Zhang committed
977
}
978
979
980

} // namespace ck
#endif