Parse more source stuff

1 files changed, 592 insertions, 0 deletions

diff --git a/ptx/src/test/vectorAdd_kernel64.ptx b/ptx/src/test/vectorAdd_kernel64.ptx
new file mode 100644
index 0000000..100cd93
--- /dev/null
+++ b/ptx/src/test/vectorAdd_kernel64.ptx
@@ -0,0 +1,592 @@
+//
+// Generated by NVIDIA NVVM Compiler
+//
+// Compiler Build ID: CL-27506705
+// Cuda compilation tools, release 10.2, V10.2.89
+// Based on LLVM 3.4svn
+//
+
+.version 6.5
+.target sm_30, debug
+.address_size 64
+
+	// .globl	VecAdd_kernel
+
+.visible .entry VecAdd_kernel(
+	.param .u64 VecAdd_kernel_param_0,
+	.param .u64 VecAdd_kernel_param_1,
+	.param .u64 VecAdd_kernel_param_2,
+	.param .u32 VecAdd_kernel_param_3
+)
+{
+	.reg .pred 	%p<3>;
+	.reg .f32 	%f<4>;
+	.reg .b32 	%r<7>;
+	.reg .b64 	%rd<13>;
+
+
+	.loc 1 21 1
+func_begin0:
+	.loc	1 0 0
+
+	.loc 1 21 1
+
+	ld.param.u64 	%rd1, [VecAdd_kernel_param_0];
+	ld.param.u64 	%rd2, [VecAdd_kernel_param_1];
+	ld.param.u64 	%rd3, [VecAdd_kernel_param_2];
+	ld.param.u32 	%r2, [VecAdd_kernel_param_3];
+func_exec_begin0:
+	.loc	1 23 11
+tmp0:
+	mov.u32 	%r3, %ntid.x;
+	mov.u32 	%r4, %ctaid.x;
+	mul.lo.s32 	%r5, %r3, %r4;
+	mov.u32 	%r6, %tid.x;
+	add.s32 	%r1, %r5, %r6;
+tmp1:
+	.loc	1 25 5
+	setp.lt.s32	%p1, %r1, %r2;
+	not.pred 	%p2, %p1;
+	@%p2 bra 	BB0_2;
+	bra.uni 	BB0_1;
+
+BB0_1:
+	.loc	1 26 9
+tmp2:
+	cvt.s64.s32	%rd4, %r1;
+	shl.b64 	%rd5, %rd4, 2;
+	add.s64 	%rd6, %rd1, %rd5;
+	ld.f32 	%f1, [%rd6];
+	cvt.s64.s32	%rd7, %r1;
+	shl.b64 	%rd8, %rd7, 2;
+	add.s64 	%rd9, %rd2, %rd8;
+	ld.f32 	%f2, [%rd9];
+	add.f32 	%f3, %f1, %f2;
+	cvt.s64.s32	%rd10, %r1;
+	shl.b64 	%rd11, %rd10, 2;
+	add.s64 	%rd12, %rd3, %rd11;
+	st.f32 	[%rd12], %f3;
+tmp3:
+
+BB0_2:
+	.loc	1 27 1
+	ret;
+tmp4:
+func_end0:
+}
+
+	.file	1 "/home/vosen/cuda-samples/0_Simple/vectorAddMMAP/vectorAdd_kernel.cu", 1581801938, 860
+
+.section .debug_info {
+ .b32 314
+ .b8 2
+ .b8 0
+ .b32 .debug_abbrev
+ .b8 8
+ .b8 1
+
+ .b8 108
+ .b8 103
+ .b8 101
+ .b8 110
+ .b8 102
+ .b8 101
+ .b8 58
+ .b8 32
+ .b8 69
+ .b8 68
+ .b8 71
+ .b8 32
+ .b8 53
+ .b8 46
+ .b8 48
+
+ .b8 0
+ .b8 4
+ .b8 118
+ .b8 101
+ .b8 99
+ .b8 116
+ .b8 111
+ .b8 114
+ .b8 65
+ .b8 100
+ .b8 100
+ .b8 95
+ .b8 107
+ .b8 101
+ .b8 114
+ .b8 110
+ .b8 101
+ .b8 108
+ .b8 46
+ .b8 99
+ .b8 117
+
+ .b8 0
+ .b64 0
+ .b32 .debug_line
+ .b8 47
+ .b8 104
+ .b8 111
+ .b8 109
+ .b8 101
+ .b8 47
+ .b8 118
+ .b8 111
+ .b8 115
+ .b8 101
+ .b8 110
+ .b8 47
+ .b8 99
+ .b8 117
+ .b8 100
+ .b8 97
+ .b8 45
+ .b8 115
+ .b8 97
+ .b8 109
+ .b8 112
+ .b8 108
+ .b8 101
+ .b8 115
+ .b8 47
+ .b8 48
+ .b8 95
+ .b8 83
+ .b8 105
+ .b8 109
+ .b8 112
+ .b8 108
+ .b8 101
+ .b8 47
+ .b8 118
+ .b8 101
+ .b8 99
+ .b8 116
+ .b8 111
+ .b8 114
+ .b8 65
+ .b8 100
+ .b8 100
+ .b8 77
+ .b8 77
+ .b8 65
+ .b8 80
+
+ .b8 0
+ .b8 2
+
+ .b8 86
+ .b8 101
+ .b8 99
+ .b8 65
+ .b8 100
+ .b8 100
+ .b8 95
+ .b8 107
+ .b8 101
+ .b8 114
+ .b8 110
+ .b8 101
+ .b8 108
+
+ .b8 0
+ .b8 86
+ .b8 101
+ .b8 99
+ .b8 65
+ .b8 100
+ .b8 100
+ .b8 95
+ .b8 107
+ .b8 101
+ .b8 114
+ .b8 110
+ .b8 101
+ .b8 108
+
+ .b8 0
+ .b8 1
+ .b8 21
+ .b32 278
+ .b8 1
+ .b64 func_begin0
+ .b64 func_end0
+ .b8 1
+ .b8 156
+ .b8 3
+
+ .b8 65
+
+ .b8 0
+ .b8 1
+ .b8 21
+ .b32 284
+ .b8 9
+ .b8 3
+ .b64 VecAdd_kernel_param_0
+ .b8 7
+ .b8 3
+
+ .b8 66
+
+ .b8 0
+ .b8 1
+ .b8 21
+ .b32 284
+ .b8 9
+ .b8 3
+ .b64 VecAdd_kernel_param_1
+ .b8 7
+ .b8 3
+
+ .b8 67
+
+ .b8 0
+ .b8 1
+ .b8 21
+ .b32 304
+ .b8 9
+ .b8 3
+ .b64 VecAdd_kernel_param_2
+ .b8 7
+ .b8 3
+
+ .b8 78
+
+ .b8 0
+ .b8 1
+ .b8 21
+ .b32 310
+ .b8 9
+ .b8 3
+ .b64 VecAdd_kernel_param_3
+ .b8 7
+ .b8 4
+
+ .b64 tmp0
+ .b64 tmp4
+ .b8 5
+
+ .b8 105
+
+ .b8 0
+ .b8 1
+ .b8 23
+ .b32 310
+ .b8 5
+ .b8 144
+ .b8 177
+ .b8 228
+ .b8 149
+ .b8 1
+ .b8 2
+ .b8 0
+ .b8 0
+ .b8 6
+
+ .b8 118
+ .b8 111
+ .b8 105
+ .b8 100
+
+ .b8 0
+ .b8 7
+
+ .b32 290
+ .b8 12
+ .b8 8
+
+ .b32 295
+ .b8 9
+
+ .b8 102
+ .b8 108
+ .b8 111
+ .b8 97
+ .b8 116
+
+ .b8 0
+ .b8 4
+ .b8 4
+ .b8 7
+
+ .b32 295
+ .b8 12
+ .b8 9
+
+ .b8 105
+ .b8 110
+ .b8 116
+
+ .b8 0
+ .b8 5
+ .b8 4
+ .b8 0
+}
+.section .debug_abbrev {
+ .b8 1
+
+ .b8 17
+
+ .b8 1
+
+ .b8 37
+
+ .b8 8
+
+ .b8 19
+
+ .b8 11
+
+ .b8 3
+
+ .b8 8
+
+ .b8 17
+
+ .b8 1
+
+ .b8 16
+
+ .b8 6
+
+ .b8 27
+
+ .b8 8
+
+ .b8 0
+
+ .b8 0
+
+ .b8 2
+
+ .b8 46
+
+ .b8 1
+
+ .b8 135
+ .b8 64
+
+ .b8 8
+
+ .b8 3
+
+ .b8 8
+
+ .b8 58
+
+ .b8 11
+
+ .b8 59
+
+ .b8 11
+
+ .b8 73
+
+ .b8 19
+
+ .b8 63
+
+ .b8 12
+
+ .b8 17
+
+ .b8 1
+
+ .b8 18
+
+ .b8 1
+
+ .b8 64
+
+ .b8 10
+
+ .b8 0
+
+ .b8 0
+
+ .b8 3
+
+ .b8 5
+
+ .b8 0
+
+ .b8 3
+
+ .b8 8
+
+ .b8 58
+
+ .b8 11
+
+ .b8 59
+
+ .b8 11
+
+ .b8 73
+
+ .b8 19
+
+ .b8 2
+
+ .b8 10
+
+ .b8 51
+
+ .b8 11
+
+ .b8 0
+
+ .b8 0
+
+ .b8 4
+
+ .b8 11
+
+ .b8 1
+
+ .b8 17
+
+ .b8 1
+
+ .b8 18
+
+ .b8 1
+
+ .b8 0
+
+ .b8 0
+
+ .b8 5
+
+ .b8 52
+
+ .b8 0
+
+ .b8 3
+
+ .b8 8
+
+ .b8 58
+
+ .b8 11
+
+ .b8 59
+
+ .b8 11
+
+ .b8 73
+
+ .b8 19
+
+ .b8 2
+
+ .b8 10
+
+ .b8 51
+
+ .b8 11
+
+ .b8 0
+
+ .b8 0
+
+ .b8 6
+
+ .b8 59
+
+ .b8 0
+
+ .b8 3
+
+ .b8 8
+
+ .b8 0
+
+ .b8 0
+
+ .b8 7
+
+ .b8 15
+
+ .b8 0
+
+ .b8 73
+
+ .b8 19
+
+ .b8 51
+
+ .b8 11
+
+ .b8 0
+
+ .b8 0
+
+ .b8 8
+
+ .b8 38
+
+ .b8 0
+
+ .b8 73
+
+ .b8 19
+
+ .b8 0
+
+ .b8 0
+
+ .b8 9
+
+ .b8 36
+
+ .b8 0
+
+ .b8 3
+
+ .b8 8
+
+ .b8 62
+
+ .b8 11
+
+ .b8 11
+
+ .b8 11
+
+ .b8 0
+
+ .b8 0
+
+ .b8 0
+
+}
+.section .debug_ranges {
+}
+.section .debug_pubnames {
+ .b32 32
+ .b8 2
+ .b8 0
+ .b32 .debug_info
+ .b32 314
+ .b32 109
+ .b8 86
+ .b8 101
+ .b8 99
+ .b8 65
+ .b8 100
+ .b8 100
+ .b8 95
+ .b8 107
+ .b8 101
+ .b8 114
+ .b8 110
+ .b8 101
+ .b8 108
+ .b8 0
+
+ .b32 0
+}