-
Notifications
You must be signed in to change notification settings - Fork 2
/
transpose.ptx
94 lines (83 loc) · 2.16 KB
/
transpose.ptx
1
2
3
4
5
6
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
33
34
35
36
37
38
39
40
41
42
43
44
45
46
47
48
49
50
51
52
53
54
55
56
57
58
59
60
61
62
63
64
65
66
67
68
69
70
71
72
73
74
75
76
77
78
79
80
81
82
83
84
85
86
87
88
89
90
91
92
93
//
// Generated by NVIDIA NVVM Compiler
//
// Compiler Build ID: CL-33281558
// Cuda compilation tools, release 12.3, V12.3.52
// Based on NVVM 7.0.1
//
.version 8.3
.target sm_89
.address_size 64
// .globl _Z9transposePfS_m
// _ZZ9transposePfS_mE12sharedMemory has been demoted
.visible .entry _Z9transposePfS_m(
.param .u64 _Z9transposePfS_m_param_0,
.param .u64 _Z9transposePfS_m_param_1,
.param .u64 _Z9transposePfS_m_param_2
)
{
.reg .pred %p<7>;
.reg .f32 %f<3>;
.reg .b32 %r<26>;
.reg .b64 %rd<14>;
// demoted variable
.shared .align 4 .b8 _ZZ9transposePfS_mE12sharedMemory[4096];
ld.param.u64 %rd1, [_Z9transposePfS_m_param_0];
ld.param.u64 %rd2, [_Z9transposePfS_m_param_1];
ld.param.u64 %rd3, [_Z9transposePfS_m_param_2];
mov.u32 %r8, %ntid.x;
mov.u32 %r9, %ctaid.x;
mov.u32 %r1, %tid.x;
mad.lo.s32 %r2, %r9, %r8, %r1;
mov.u32 %r10, %ntid.y;
mov.u32 %r11, %ctaid.y;
mov.u32 %r3, %tid.y;
mad.lo.s32 %r4, %r11, %r10, %r3;
mad.lo.s32 %r5, %r11, %r8, %r1;
mad.lo.s32 %r6, %r9, %r10, %r3;
cvt.s64.s32 %rd4, %r2;
setp.lt.u64 %p1, %rd4, %rd3;
cvt.s64.s32 %rd5, %r4;
setp.lt.u64 %p2, %rd5, %rd3;
and.pred %p3, %p1, %p2;
shl.b32 %r12, %r1, 7;
mov.u32 %r13, _ZZ9transposePfS_mE12sharedMemory;
add.s32 %r14, %r13, %r12;
shl.b32 %r15, %r3, 2;
add.s32 %r7, %r14, %r15;
@%p3 bra $L__BB0_2;
bra.uni $L__BB0_1;
$L__BB0_2:
cvta.to.global.u64 %rd6, %rd1;
cvt.u32.u64 %r17, %rd3;
mad.lo.s32 %r18, %r4, %r17, %r2;
mul.wide.s32 %rd7, %r18, 4;
add.s64 %rd8, %rd6, %rd7;
ld.global.f32 %f1, [%rd8];
st.shared.f32 [%r7], %f1;
bra.uni $L__BB0_3;
$L__BB0_1:
mov.u32 %r16, 0;
st.shared.u32 [%r7], %r16;
$L__BB0_3:
bar.sync 0;
cvt.s64.s32 %rd9, %r5;
setp.ge.u64 %p4, %rd9, %rd3;
cvt.s64.s32 %rd10, %r6;
setp.ge.u64 %p5, %rd10, %rd3;
or.pred %p6, %p4, %p5;
@%p6 bra $L__BB0_5;
cvt.u32.u64 %r19, %rd3;
mad.lo.s32 %r20, %r6, %r19, %r5;
shl.b32 %r21, %r3, 7;
add.s32 %r23, %r13, %r21;
shl.b32 %r24, %r1, 2;
add.s32 %r25, %r23, %r24;
ld.shared.f32 %f2, [%r25];
cvta.to.global.u64 %rd11, %rd2;
mul.wide.s32 %rd12, %r20, 4;
add.s64 %rd13, %rd11, %rd12;
st.global.f32 [%rd13], %f2;
$L__BB0_5:
ret;
}