ptxlint 0.0.4

Static analysis and lints for NVIDIA PTX, no GPU needed. 不需要显卡的 NVIDIA PTX 静态分析工具
Documentation 
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
94
95
96
97
98
99
100
101
102
103
104
105
106
107
108
109
110
111
112
113
114
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
168
169
170
171
172
173
174
175
176
177
178
179
180
181
182
183
184

//
// Generated by LLVM NVPTX Back-End
//

.version 7.0
.target sm_70
.address_size 64

	// .globl	mix16                   // -- Begin function mix16
                                        // @mix16
.visible .entry mix16(
	.param .u64 .ptr .align 1 mix16_param_0,
	.param .u64 .ptr .align 1 mix16_param_1,
	.param .u32 mix16_param_2
)
{
	.local .align 4 .b8 	__local_depot0[64];
	.reg .b64 	%SP;
	.reg .b64 	%SPL;
	.reg .pred 	%p<2>;
	.reg .b32 	%r<98>;
	.reg .b64 	%rd<39>;

// %bb.0:
	mov.b64 	%SPL, __local_depot0;
	ld.param.b64 	%rd4, [mix16_param_0];
	cvta.to.global.u64 	%rd1, %rd4;
	ld.param.b64 	%rd5, [mix16_param_1];
	cvta.to.global.u64 	%rd2, %rd5;
	ld.param.b32 	%r2, [mix16_param_2];
	add.u64 	%rd3, %SPL, 0;
	mov.u32 	%r3, %ctaid.x;
	mov.u32 	%r4, %ntid.x;
	mov.u32 	%r5, %tid.x;
	mad.lo.s32 	%r1, %r3, %r4, %r5;
	setp.ge.u32 	%p1, %r1, %r2;
	@%p1 bra 	$L__BB0_2;
// %bb.1:
	mul.wide.u32 	%rd6, %r1, 4;
	add.s64 	%rd7, %rd1, %rd6;
	ld.global.b32 	%r6, [%rd7];
	st.local.b32 	[%rd3], %r6;
	ld.global.b32 	%r7, [%rd7+4];
	st.local.b32 	[%rd3+4], %r7;
	ld.global.b32 	%r8, [%rd7+8];
	st.local.b32 	[%rd3+8], %r8;
	ld.global.b32 	%r9, [%rd7+12];
	st.local.b32 	[%rd3+12], %r9;
	ld.global.b32 	%r10, [%rd7+16];
	st.local.b32 	[%rd3+16], %r10;
	ld.global.b32 	%r11, [%rd7+20];
	st.local.b32 	[%rd3+20], %r11;
	ld.global.b32 	%r12, [%rd7+24];
	st.local.b32 	[%rd3+24], %r12;
	ld.global.b32 	%r13, [%rd7+28];
	st.local.b32 	[%rd3+28], %r13;
	ld.global.b32 	%r14, [%rd7+32];
	st.local.b32 	[%rd3+32], %r14;
	ld.global.b32 	%r15, [%rd7+36];
	st.local.b32 	[%rd3+36], %r15;
	ld.global.b32 	%r16, [%rd7+40];
	st.local.b32 	[%rd3+40], %r16;
	ld.global.b32 	%r17, [%rd7+44];
	st.local.b32 	[%rd3+44], %r17;
	ld.global.b32 	%r18, [%rd7+48];
	st.local.b32 	[%rd3+48], %r18;
	ld.global.b32 	%r19, [%rd7+52];
	st.local.b32 	[%rd3+52], %r19;
	ld.global.b32 	%r20, [%rd7+56];
	st.local.b32 	[%rd3+56], %r20;
	ld.global.b32 	%r21, [%rd7+60];
	st.local.b32 	[%rd3+60], %r21;
	shf.l.wrap.b32 	%r22, %r6, %r6, 3;
	add.s32 	%r23, %r22, 1;
	and.b32 	%r24, %r23, 15;
	mul.wide.u32 	%rd8, %r24, 4;
	add.s64 	%rd9, %rd3, %rd8;
	ld.local.b32 	%r25, [%rd9];
	add.s32 	%r26, %r25, %r22;
	shf.l.wrap.b32 	%r27, %r26, %r26, 3;
	add.s32 	%r28, %r27, 2;
	and.b32 	%r29, %r28, 15;
	mul.wide.u32 	%rd10, %r29, 4;
	add.s64 	%rd11, %rd3, %rd10;
	ld.local.b32 	%r30, [%rd11];
	add.s32 	%r31, %r27, %r30;
	shf.l.wrap.b32 	%r32, %r31, %r31, 3;
	add.s32 	%r33, %r32, 3;
	and.b32 	%r34, %r33, 15;
	mul.wide.u32 	%rd12, %r34, 4;
	add.s64 	%rd13, %rd3, %rd12;
	ld.local.b32 	%r35, [%rd13];
	add.s32 	%r36, %r32, %r35;
	shf.l.wrap.b32 	%r37, %r36, %r36, 3;
	add.s32 	%r38, %r37, 4;
	and.b32 	%r39, %r38, 15;
	mul.wide.u32 	%rd14, %r39, 4;
	add.s64 	%rd15, %rd3, %rd14;
	ld.local.b32 	%r40, [%rd15];
	add.s32 	%r41, %r37, %r40;
	shf.l.wrap.b32 	%r42, %r41, %r41, 3;
	add.s32 	%r43, %r42, 5;
	and.b32 	%r44, %r43, 15;
	mul.wide.u32 	%rd16, %r44, 4;
	add.s64 	%rd17, %rd3, %rd16;
	ld.local.b32 	%r45, [%rd17];
	add.s32 	%r46, %r42, %r45;
	shf.l.wrap.b32 	%r47, %r46, %r46, 3;
	add.s32 	%r48, %r47, 6;
	and.b32 	%r49, %r48, 15;
	mul.wide.u32 	%rd18, %r49, 4;
	add.s64 	%rd19, %rd3, %rd18;
	ld.local.b32 	%r50, [%rd19];
	add.s32 	%r51, %r47, %r50;
	shf.l.wrap.b32 	%r52, %r51, %r51, 3;
	add.s32 	%r53, %r52, 7;
	and.b32 	%r54, %r53, 15;
	mul.wide.u32 	%rd20, %r54, 4;
	add.s64 	%rd21, %rd3, %rd20;
	ld.local.b32 	%r55, [%rd21];
	add.s32 	%r56, %r52, %r55;
	shf.l.wrap.b32 	%r57, %r56, %r56, 3;
	and.b32 	%r58, %r57, 15;
	xor.b32 	%r59, %r58, 8;
	mul.wide.u32 	%rd22, %r59, 4;
	add.s64 	%rd23, %rd3, %rd22;
	ld.local.b32 	%r60, [%rd23];
	add.s32 	%r61, %r57, %r60;
	shf.l.wrap.b32 	%r62, %r61, %r61, 3;
	add.s32 	%r63, %r62, 9;
	and.b32 	%r64, %r63, 15;
	mul.wide.u32 	%rd24, %r64, 4;
	add.s64 	%rd25, %rd3, %rd24;
	ld.local.b32 	%r65, [%rd25];
	add.s32 	%r66, %r62, %r65;
	shf.l.wrap.b32 	%r67, %r66, %r66, 3;
	add.s32 	%r68, %r67, 10;
	and.b32 	%r69, %r68, 15;
	mul.wide.u32 	%rd26, %r69, 4;
	add.s64 	%rd27, %rd3, %rd26;
	ld.local.b32 	%r70, [%rd27];
	add.s32 	%r71, %r67, %r70;
	shf.l.wrap.b32 	%r72, %r71, %r71, 3;
	add.s32 	%r73, %r72, 11;
	and.b32 	%r74, %r73, 15;
	mul.wide.u32 	%rd28, %r74, 4;
	add.s64 	%rd29, %rd3, %rd28;
	ld.local.b32 	%r75, [%rd29];
	add.s32 	%r76, %r72, %r75;
	shf.l.wrap.b32 	%r77, %r76, %r76, 3;
	add.s32 	%r78, %r77, 12;
	and.b32 	%r79, %r78, 15;
	mul.wide.u32 	%rd30, %r79, 4;
	add.s64 	%rd31, %rd3, %rd30;
	ld.local.b32 	%r80, [%rd31];
	add.s32 	%r81, %r77, %r80;
	shf.l.wrap.b32 	%r82, %r81, %r81, 3;
	add.s32 	%r83, %r82, 13;
	and.b32 	%r84, %r83, 15;
	mul.wide.u32 	%rd32, %r84, 4;
	add.s64 	%rd33, %rd3, %rd32;
	ld.local.b32 	%r85, [%rd33];
	add.s32 	%r86, %r82, %r85;
	shf.l.wrap.b32 	%r87, %r86, %r86, 3;
	add.s32 	%r88, %r87, 14;
	and.b32 	%r89, %r88, 15;
	mul.wide.u32 	%rd34, %r89, 4;
	add.s64 	%rd35, %rd3, %rd34;
	ld.local.b32 	%r90, [%rd35];
	add.s32 	%r91, %r87, %r90;
	shf.l.wrap.b32 	%r92, %r91, %r91, 3;
	add.s32 	%r93, %r92, -1;
	and.b32 	%r94, %r93, 15;
	mul.wide.u32 	%rd36, %r94, 4;
	add.s64 	%rd37, %rd3, %rd36;
	ld.local.b32 	%r95, [%rd37];
	add.s32 	%r96, %r92, %r95;
	shf.l.wrap.b32 	%r97, %r96, %r96, 3;
	add.s64 	%rd38, %rd2, %rd6;
	st.global.b32 	[%rd38], %r97;
$L__BB0_2:
	ret;
                                        // -- End function
}