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

//
// 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
)
{
	.reg .pred 	%p<2>;
	.reg .b32 	%r<53>;
	.reg .b64 	%rd<8>;

// %bb.0:
	ld.param.b64 	%rd3, [mix16_param_0];
	cvta.to.global.u64 	%rd1, %rd3;
	ld.param.b64 	%rd4, [mix16_param_1];
	cvta.to.global.u64 	%rd2, %rd4;
	ld.param.b32 	%r2, [mix16_param_2];
	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 	%rd5, %r1, 4;
	add.s64 	%rd6, %rd1, %rd5;
	ld.global.b32 	%r6, [%rd6];
	ld.global.b32 	%r7, [%rd6+4];
	ld.global.b32 	%r8, [%rd6+8];
	ld.global.b32 	%r9, [%rd6+12];
	ld.global.b32 	%r10, [%rd6+16];
	ld.global.b32 	%r11, [%rd6+20];
	ld.global.b32 	%r12, [%rd6+24];
	ld.global.b32 	%r13, [%rd6+28];
	ld.global.b32 	%r14, [%rd6+32];
	ld.global.b32 	%r15, [%rd6+36];
	ld.global.b32 	%r16, [%rd6+40];
	ld.global.b32 	%r17, [%rd6+44];
	ld.global.b32 	%r18, [%rd6+48];
	ld.global.b32 	%r19, [%rd6+52];
	ld.global.b32 	%r20, [%rd6+56];
	ld.global.b32 	%r21, [%rd6+60];
	shf.l.wrap.b32 	%r22, %r6, %r6, 3;
	add.s32 	%r23, %r22, %r7;
	shf.l.wrap.b32 	%r24, %r23, %r23, 3;
	add.s32 	%r25, %r24, %r8;
	shf.l.wrap.b32 	%r26, %r25, %r25, 3;
	add.s32 	%r27, %r26, %r9;
	shf.l.wrap.b32 	%r28, %r27, %r27, 3;
	add.s32 	%r29, %r28, %r10;
	shf.l.wrap.b32 	%r30, %r29, %r29, 3;
	add.s32 	%r31, %r30, %r11;
	shf.l.wrap.b32 	%r32, %r31, %r31, 3;
	add.s32 	%r33, %r32, %r12;
	shf.l.wrap.b32 	%r34, %r33, %r33, 3;
	add.s32 	%r35, %r34, %r13;
	shf.l.wrap.b32 	%r36, %r35, %r35, 3;
	add.s32 	%r37, %r36, %r14;
	shf.l.wrap.b32 	%r38, %r37, %r37, 3;
	add.s32 	%r39, %r38, %r15;
	shf.l.wrap.b32 	%r40, %r39, %r39, 3;
	add.s32 	%r41, %r40, %r16;
	shf.l.wrap.b32 	%r42, %r41, %r41, 3;
	add.s32 	%r43, %r42, %r17;
	shf.l.wrap.b32 	%r44, %r43, %r43, 3;
	add.s32 	%r45, %r44, %r18;
	shf.l.wrap.b32 	%r46, %r45, %r45, 3;
	add.s32 	%r47, %r46, %r19;
	shf.l.wrap.b32 	%r48, %r47, %r47, 3;
	add.s32 	%r49, %r48, %r20;
	shf.l.wrap.b32 	%r50, %r49, %r49, 3;
	add.s32 	%r51, %r50, %r21;
	shf.l.wrap.b32 	%r52, %r51, %r51, 3;
	add.s64 	%rd7, %rd2, %rd5;
	st.global.b32 	[%rd7], %r52;
$L__BB0_2:
	ret;
                                        // -- End function
}