Add files using upload-large-folder tool

.triton/dump/11759acf26ac56366b171628132485d6/triton_.ptx

@@ -0,0 +1,788 @@
+//
+// Generated by LLVM NVPTX Back-End
+//
+
+.version 8.2
+.target sm_89
+.address_size 64
+
+	// .globl	triton__0d1d2d3d4d5d6d7d8de9de
+.extern .shared .align 1 .b8 global_smem[];
+.global .align 1 .b8 _$_str[11] = {95, 95, 67, 85, 68, 65, 95, 70, 84, 90, 0};
+
+.visible .entry triton__0d1d2d3d4d5d6d7d8de9de(
+	.param .u64 triton__0d1d2d3d4d5d6d7d8de9de_param_0,
+	.param .u64 triton__0d1d2d3d4d5d6d7d8de9de_param_1,
+	.param .u64 triton__0d1d2d3d4d5d6d7d8de9de_param_2,
+	.param .u64 triton__0d1d2d3d4d5d6d7d8de9de_param_3,
+	.param .u64 triton__0d1d2d3d4d5d6d7d8de9de_param_4,
+	.param .u64 triton__0d1d2d3d4d5d6d7d8de9de_param_5,
+	.param .u64 triton__0d1d2d3d4d5d6d7d8de9de_param_6,
+	.param .u64 triton__0d1d2d3d4d5d6d7d8de9de_param_7,
+	.param .u32 triton__0d1d2d3d4d5d6d7d8de9de_param_8,
+	.param .u32 triton__0d1d2d3d4d5d6d7d8de9de_param_9
+)
+.maxntid 64, 1, 1
+{
+	.reg .pred 	%p<31>;
+	.reg .b16 	%rs<17>;
+	.reg .b32 	%r<103>;
+	.reg .f32 	%f<86>;
+	.reg .b64 	%rd<21>;
+	.loc	1 18 0
+$L__func_begin0:
+	.loc	1 18 0
+
+	ld.param.u64 	%rd9, [triton__0d1d2d3d4d5d6d7d8de9de_param_0];
+	ld.param.u64 	%rd10, [triton__0d1d2d3d4d5d6d7d8de9de_param_1];
+$L__tmp0:
+	.loc	1 26 26
+	mov.u32 	%r68, %tid.x;
+	and.b32  	%r69, %r68, 31;
+	ld.param.u64 	%rd11, [triton__0d1d2d3d4d5d6d7d8de9de_param_2];
+	ld.param.u64 	%rd12, [triton__0d1d2d3d4d5d6d7d8de9de_param_3];
+	ld.param.u64 	%rd13, [triton__0d1d2d3d4d5d6d7d8de9de_param_4];
+	and.b32  	%r70, %r68, 63;
+	ld.param.u64 	%rd14, [triton__0d1d2d3d4d5d6d7d8de9de_param_5];
+	shl.b32 	%r71, %r70, 2;
+	ld.param.u64 	%rd15, [triton__0d1d2d3d4d5d6d7d8de9de_param_6];
+	ld.param.u64 	%rd16, [triton__0d1d2d3d4d5d6d7d8de9de_param_7];
+	.loc	1 23 28
+	mov.u32 %r1, %ctaid.x;
+	.loc	1 30 40
+	shl.b32 	%r72, %r1, 8;
+	.loc	1 30 36
+	or.b32  	%r73, %r72, %r71;
+	.loc	1 30 30
+	mul.wide.s32 	%rd17, %r73, 4;
+	add.s64 	%rd1, %rd10, %rd17;
+	mov.b32 	%r6, 0;
+	mov.pred 	%p1, -1;
+	.loc	1 30 46
+	mov.u32 %r2, 0x0;
+	mov.u32 %r3, 0x0;
+	mov.u32 %r4, 0x0;
+	mov.u32 %r5, 0x0;
+	@%p1 ld.global.v4.b32 { %r2, %r3, %r4, %r5 }, [ %rd1 + 0 ];
+	@!%p1 mov.u32 %r2, %r6;
+	@!%p1 mov.u32 %r3, %r6;
+	@!%p1 mov.u32 %r4, %r6;
+	@!%p1 mov.u32 %r5, %r6;
+	mov.b32 	%f1, %r4;
+	mov.b32 	%f2, %r5;
+	.loc	1 31 30
+	mul.wide.s32 	%rd18, %r73, 2;
+	add.s64 	%rd2, %rd11, %rd18;
+	.loc	1 31 46
+	mov.u32 %r10, 0x0;
+	mov.u32 %r11, 0x0;
+	@%p1 ld.global.v2.b32 { %r10, %r11 }, [ %rd2 + 0 ];
+	@!%p1 mov.u32 %r10, %r6;
+	@!%p1 mov.u32 %r11, %r6;
+	cvt.u16.u32 	%rs1, %r10;
+	{ .reg .b16 tmp; mov.b32 {tmp, %rs2}, %r10; }
+	cvt.u16.u32 	%rs3, %r11;
+	{ .reg .b16 tmp; mov.b32 {tmp, %rs4}, %r11; }
+	.loc	1 31 67
+	cvt.f32.bf16 %r14, %rs1;
+	mov.b32 	%f3, %r14;
+	cvt.f32.bf16 %r15, %rs2;
+	mov.b32 	%f4, %r15;
+	cvt.f32.bf16 %r16, %rs3;
+	mov.b32 	%f5, %r16;
+	cvt.f32.bf16 %r17, %rs4;
+	mov.b32 	%f6, %r17;
+	.loc	1 32 30
+	add.s64 	%rd3, %rd12, %rd18;
+	.loc	1 32 46
+	mov.u32 %r18, 0x0;
+	mov.u32 %r19, 0x0;
+	@%p1 ld.global.v2.b32 { %r18, %r19 }, [ %rd3 + 0 ];
+	@!%p1 mov.u32 %r18, %r6;
+	@!%p1 mov.u32 %r19, %r6;
+	cvt.u16.u32 	%rs5, %r18;
+	{ .reg .b16 tmp; mov.b32 {tmp, %rs6}, %r18; }
+	cvt.u16.u32 	%rs7, %r19;
+	{ .reg .b16 tmp; mov.b32 {tmp, %rs8}, %r19; }
+	.loc	1 32 67
+	cvt.f32.bf16 %r22, %rs5;
+	mov.b32 	%f7, %r22;
+	cvt.f32.bf16 %r23, %rs6;
+	mov.b32 	%f8, %r23;
+	cvt.f32.bf16 %r24, %rs7;
+	mov.b32 	%f9, %r24;
+	cvt.f32.bf16 %r25, %rs8;
+	mov.b32 	%f10, %r25;
+	.loc	1 33 30
+	add.s64 	%rd4, %rd13, %rd18;
+	.loc	1 33 46
+	mov.u32 %r26, 0x0;
+	mov.u32 %r27, 0x0;
+	@%p1 ld.global.v2.b32 { %r26, %r27 }, [ %rd4 + 0 ];
+	@!%p1 mov.u32 %r26, %r6;
+	@!%p1 mov.u32 %r27, %r6;
+	cvt.u16.u32 	%rs9, %r26;
+	{ .reg .b16 tmp; mov.b32 {tmp, %rs10}, %r26; }
+	cvt.u16.u32 	%rs11, %r27;
+	{ .reg .b16 tmp; mov.b32 {tmp, %rs12}, %r27; }
+	.loc	1 33 67
+	cvt.f32.bf16 %r30, %rs9;
+	mov.b32 	%f11, %r30;
+	cvt.f32.bf16 %r31, %rs10;
+	mov.b32 	%f12, %r31;
+	cvt.f32.bf16 %r32, %rs11;
+	mov.b32 	%f13, %r32;
+	cvt.f32.bf16 %r33, %rs12;
+	mov.b32 	%f14, %r33;
+	.loc	1 34 31
+	mul.wide.u32 	%rd19, %r71, 4;
+	add.s64 	%rd5, %rd14, %rd19;
+	.loc	1 34 36
+	mov.u32 %r34, 0x0;
+	mov.u32 %r35, 0x0;
+	mov.u32 %r36, 0x0;
+	mov.u32 %r37, 0x0;
+	@%p1 ld.global.L1::evict_last.v4.b32 { %r34, %r35, %r36, %r37 }, [ %rd5 + 0 ];
+	@!%p1 mov.u32 %r34, %r6;
+	@!%p1 mov.u32 %r35, %r6;
+	@!%p1 mov.u32 %r36, %r6;
+	@!%p1 mov.u32 %r37, %r6;
+	.loc	1 36 18
+	add.f32 	%f15, %f5, %f1;
+	add.f32 	%f16, %f6, %f2;
+	.loc	1 38 18
+	add.f32 	%f17, %f15, %f9;
+	add.f32 	%f18, %f16, %f10;
+	.loc	1 30 46
+	mov.b32 	%f19, %r2;
+	mov.b32 	%f20, %r3;
+	.loc	1 36 18
+	add.f32 	%f21, %f4, %f20;
+	add.f32 	%f22, %f3, %f19;
+	.loc	1 38 18
+	add.f32 	%f23, %f22, %f7;
+	add.f32 	%f24, %f21, %f8;
+	.loc	1 40 18
+	add.f32 	%f25, %f24, %f12;
+	add.f32 	%f26, %f23, %f11;
+	add.f32 	%f27, %f17, %f13;
+	add.f32 	%f28, %f18, %f14;
+$L__tmp1:
+	.loc	2 233 15
+	add.f32 	%f29, %f26, %f25;
+	add.f32 	%f30, %f29, %f27;
+	add.f32 	%f31, %f30, %f28;
+$L__tmp2:
+	.loc	2 243 36
+	mov.b32 	%r74, %f31;
+	shfl.sync.bfly.b32	%r75, %r74, 16, 31, -1;
+	mov.b32 	%f32, %r75;
+$L__tmp3:
+	.loc	2 233 15
+	add.f32 	%f33, %f31, %f32;
+$L__tmp4:
+	.loc	2 243 36
+	mov.b32 	%r76, %f33;
+	shfl.sync.bfly.b32	%r77, %r76, 8, 31, -1;
+	mov.b32 	%f34, %r77;
+$L__tmp5:
+	.loc	2 233 15
+	add.f32 	%f35, %f33, %f34;
+$L__tmp6:
+	.loc	2 243 36
+	mov.b32 	%r78, %f35;
+	shfl.sync.bfly.b32	%r79, %r78, 4, 31, -1;
+	mov.b32 	%f36, %r79;
+$L__tmp7:
+	.loc	2 233 15
+	add.f32 	%f37, %f35, %f36;
+$L__tmp8:
+	.loc	2 243 36
+	mov.b32 	%r80, %f37;
+	shfl.sync.bfly.b32	%r81, %r80, 2, 31, -1;
+	mov.b32 	%f38, %r81;
+$L__tmp9:
+	.loc	2 233 15
+	add.f32 	%f39, %f37, %f38;
+$L__tmp10:
+	.loc	2 243 36
+	mov.b32 	%r82, %f39;
+	shfl.sync.bfly.b32	%r83, %r82, 1, 31, -1;
+	mov.b32 	%f40, %r83;
+$L__tmp11:
+	.loc	2 233 15
+	add.f32 	%f41, %f39, %f40;
+$L__tmp12:
+	.loc	2 243 36
+	setp.eq.s32 	%p20, %r69, 0;
+	shr.u32 	%r84, %r68, 3;
+	and.b32  	%r85, %r84, 4;
+	mov.u32 	%r86, global_smem;
+	add.s32 	%r42, %r86, %r85;
+	mov.b32 	%r43, %f41;
+	@%p20 st.shared.b32 [ %r42 + 0 ], %r43;
+	bar.sync 	0;
+	setp.lt.s32 	%p21, %r68, 2;
+	shl.b32 	%r87, %r68, 2;
+	add.s32 	%r45, %r86, %r87;
+	@%p21 ld.shared.b32 %r44, [ %r45 + 0 ];
+	mov.b32 	%f42, %r44;
+	shfl.sync.bfly.b32	%r88, %r44, 1, 31, -1;
+	mov.b32 	%f43, %r88;
+$L__tmp13:
+	.loc	2 233 15
+	add.f32 	%f44, %f42, %f43;
+$L__tmp14:
+	.loc	2 243 36
+	and.b32  	%r89, %r68, 1;
+	setp.eq.b32 	%p29, %r89, 1;
+	not.pred 	%p30, %p29;
+	and.pred  	%p22, %p21, %p30;
+	mov.b32 	%r47, %f44;
+	@%p22 st.shared.b32 [ %r45 + 0 ], %r47;
+	bar.sync 	0;
+	ld.shared.f32 	%f45, [global_smem];
+$L__tmp15:
+	.loc	3 8 15
+	add.f32 	%f46, %f45, 0f00000000;
+$L__tmp16:
+	.loc	1 48 20
+	mov.b32 	%r49, %f46;
+	mov.b32 	%r50, 1132462080;
+	div.full.f32 %r67, %r49, %r50;
+	mov.b32 	%f47, %r67;
+	.loc	1 49 20
+	sub.f32 	%f48, %f26, %f47;
+	sub.f32 	%f49, %f25, %f47;
+	sub.f32 	%f50, %f27, %f47;
+	sub.f32 	%f51, %f28, %f47;
+	.loc	1 50 20
+	mul.f32 	%f52, %f49, %f49;
+$L__tmp17:
+	.loc	2 243 36
+	bar.sync 	0;
+$L__tmp18:
+	.loc	2 233 15
+	fma.rn.f32 	%f53, %f48, %f48, %f52;
+	fma.rn.f32 	%f54, %f50, %f50, %f53;
+	fma.rn.f32 	%f55, %f51, %f51, %f54;
+$L__tmp19:
+	.loc	2 243 36
+	mov.b32 	%r90, %f55;
+	shfl.sync.bfly.b32	%r91, %r90, 16, 31, -1;
+	mov.b32 	%f56, %r91;
+$L__tmp20:
+	.loc	2 233 15
+	add.f32 	%f57, %f55, %f56;
+$L__tmp21:
+	.loc	2 243 36
+	mov.b32 	%r92, %f57;
+	shfl.sync.bfly.b32	%r93, %r92, 8, 31, -1;
+	mov.b32 	%f58, %r93;
+$L__tmp22:
+	.loc	2 233 15
+	add.f32 	%f59, %f57, %f58;
+$L__tmp23:
+	.loc	2 243 36
+	mov.b32 	%r94, %f59;
+	shfl.sync.bfly.b32	%r95, %r94, 4, 31, -1;
+	mov.b32 	%f60, %r95;
+$L__tmp24:
+	.loc	2 233 15
+	add.f32 	%f61, %f59, %f60;
+$L__tmp25:
+	.loc	2 243 36
+	mov.b32 	%r96, %f61;
+	shfl.sync.bfly.b32	%r97, %r96, 2, 31, -1;
+	mov.b32 	%f62, %r97;
+$L__tmp26:
+	.loc	2 233 15
+	add.f32 	%f63, %f61, %f62;
+$L__tmp27:
+	.loc	2 243 36
+	mov.b32 	%r98, %f63;
+	shfl.sync.bfly.b32	%r99, %r98, 1, 31, -1;
+	mov.b32 	%f64, %r99;
+$L__tmp28:
+	.loc	2 233 15
+	add.f32 	%f65, %f63, %f64;
+$L__tmp29:
+	.loc	2 243 36
+	mov.b32 	%r52, %f65;
+	@%p20 st.shared.b32 [ %r42 + 0 ], %r52;
+	bar.sync 	0;
+	@%p21 ld.shared.b32 %r53, [ %r45 + 0 ];
+	mov.b32 	%f66, %r53;
+	shfl.sync.bfly.b32	%r100, %r53, 1, 31, -1;
+	mov.b32 	%f67, %r100;
+$L__tmp30:
+	.loc	2 233 15
+	add.f32 	%f68, %f66, %f67;
+$L__tmp31:
+	.loc	2 243 36
+	mov.b32 	%r56, %f68;
+	@%p22 st.shared.b32 [ %r45 + 0 ], %r56;
+	bar.sync 	0;
+	ld.shared.f32 	%f69, [global_smem];
+$L__tmp32:
+	.loc	3 8 15
+	add.f32 	%f70, %f69, 0f00000000;
+$L__tmp33:
+	.loc	1 55 20
+	mov.b32 	%r58, %f70;
+	div.full.f32 %r57, %r58, %r50;
+	mov.b32 	%f71, %r57;
+	.loc	1 57 20
+	add.f32 	%f72, %f71, 0f3727C5AC;
+	.loc	1 58 26
+	rsqrt.approx.ftz.f32 	%f73, %f72;
+	.loc	1 34 36
+	mov.b32 	%f74, %r37;
+	mov.b32 	%f75, %r36;
+	mov.b32 	%f76, %r35;
+	mov.b32 	%f77, %r34;
+	.loc	1 60 20
+	mul.f32 	%f78, %f48, %f73;
+	mul.f32 	%f79, %f49, %f73;
+	mul.f32 	%f80, %f50, %f73;
+	mul.f32 	%f81, %f51, %f73;
+	.loc	1 61 20
+	mul.f32 	%f82, %f78, %f77;
+	mul.f32 	%f83, %f79, %f76;
+	mul.f32 	%f84, %f80, %f75;
+	mul.f32 	%f85, %f81, %f74;
+	.loc	1 63 4
+	bar.sync 	0;
+	.loc	1 64 28
+	mul.wide.s32 	%rd20, %r1, 4;
+	add.s64 	%rd6, %rd9, %rd20;
+	.loc	1 64 40
+	setp.eq.s32 	%p26, %r70, 0;
+	mov.b32 	%r60, %f73;
+	@%p26 st.global.b32 [ %rd6 + 0 ], { %r60 };
+	.loc	1 65 25
+	add.s64 	%rd7, %rd16, %rd18;
+	.loc	1 65 48
+	mov.b32 	%r61, %f82;
+	cvt.rn.bf16.f32 %rs13, %r61;
+	mov.b32 	%r62, %f83;
+	cvt.rn.bf16.f32 %rs14, %r62;
+	mov.b32 	%r63, %f84;
+	cvt.rn.bf16.f32 %rs15, %r63;
+	mov.b32 	%r64, %f85;
+	cvt.rn.bf16.f32 %rs16, %r64;
+	mov.b32 	%r101, {%rs13, %rs14};
+	mov.b32 	%r102, {%rs15, %rs16};
+	@%p1 st.global.v2.b32 [ %rd7 + 0 ], { %r101, %r102 };
+	.loc	1 66 25
+	add.s64 	%rd8, %rd15, %rd20;
+	.loc	1 66 37
+	@%p26 st.global.b32 [ %rd8 + 0 ], { %r67 };
+	.loc	1 66 4
+	ret;
+$L__tmp34:
+$L__func_end0:
+
+}
+	// .globl	__nv_rsqrtf
+.visible .func  (.param .b32 func_retval0) __nv_rsqrtf(
+	.param .b32 __nv_rsqrtf_param_0
+)
+{
+	.reg .f32 	%f<3>;
+$L__func_begin1:
+
+	ld.param.f32 	%f1, [__nv_rsqrtf_param_0];
+	rsqrt.approx.ftz.f32 	%f2, %f1;
+	st.param.f32 	[func_retval0+0], %f2;
+	ret;
+$L__func_end1:
+
+}
+	.file	1 "/tmp/torchinductor_root/dx/cdxa5yqgsimvskocpuiz4ajfrjfcwys3opyrdv53xfphj4576qx7.py"
+	.file	2 "/usr/local/lib/python3.10/dist-packages/triton/language/standard.py"
+	.file	3 "/usr/local/lib/python3.10/dist-packages/torch/_inductor/triton_helpers.py"
+	.section	.debug_abbrev
+	{
+.b8 1
+.b8 17
+.b8 1
+.b8 37
+.b8 8
+.b8 19
+.b8 5
+.b8 3
+.b8 8
+.b8 16
+.b8 6
+.b8 27
+.b8 8
+.b8 180
+.b8 66
+.b8 12
+.b8 17
+.b8 1
+.b8 18
+.b8 1
+.b8 0
+.b8 0
+.b8 2
+.b8 46
+.b8 0
+.b8 135
+.b8 64
+.b8 8
+.b8 3
+.b8 8
+.b8 58
+.b8 11
+.b8 59
+.b8 11
+.b8 63
+.b8 12
+.b8 32
+.b8 11
+.b8 0
+.b8 0
+.b8 3
+.b8 46
+.b8 1
+.b8 17
+.b8 1
+.b8 18
+.b8 1
+.b8 64
+.b8 10
+.b8 49
+.b8 19
+.b8 0
+.b8 0
+.b8 4
+.b8 29
+.b8 1
+.b8 49
+.b8 19
+.b8 17
+.b8 1
+.b8 18
+.b8 1
+.b8 88
+.b8 11
+.b8 89
+.b8 11
+.b8 87
+.b8 11
+.b8 0
+.b8 0
+.b8 5
+.b8 29
+.b8 0
+.b8 49
+.b8 19
+.b8 17
+.b8 1
+.b8 18
+.b8 1
+.b8 88
+.b8 11
+.b8 89
+.b8 11
+.b8 87
+.b8 11
+.b8 0
+.b8 0
+.b8 0
+	}
+	.section	.debug_info
+	{
+.b32 407
+.b8 2
+.b8 0
+.b32 .debug_abbrev
+.b8 8
+.b8 1
+.b8 116
+.b8 114
+.b8 105
+.b8 116
+.b8 111
+.b8 110
+.b8 0
+.b8 2
+.b8 0
+.b8 99
+.b8 100
+.b8 120
+.b8 97
+.b8 53
+.b8 121
+.b8 113
+.b8 103
+.b8 115
+.b8 105
+.b8 109
+.b8 118
+.b8 115
+.b8 107
+.b8 111
+.b8 99
+.b8 112
+.b8 117
+.b8 105
+.b8 122
+.b8 52
+.b8 97
+.b8 106
+.b8 102
+.b8 114
+.b8 106
+.b8 102
+.b8 99
+.b8 119
+.b8 121
+.b8 115
+.b8 51
+.b8 111
+.b8 112
+.b8 121
+.b8 114
+.b8 100
+.b8 118
+.b8 53
+.b8 51
+.b8 120
+.b8 102
+.b8 112
+.b8 104
+.b8 106
+.b8 52
+.b8 53
+.b8 55
+.b8 54
+.b8 113
+.b8 120
+.b8 55
+.b8 46
+.b8 112
+.b8 121
+.b8 0
+.b32 .debug_line
+.b8 47
+.b8 116
+.b8 109
+.b8 112
+.b8 47
+.b8 116
+.b8 111
+.b8 114
+.b8 99
+.b8 104
+.b8 105
+.b8 110
+.b8 100
+.b8 117
+.b8 99
+.b8 116
+.b8 111
+.b8 114
+.b8 95
+.b8 114
+.b8 111
+.b8 111
+.b8 116
+.b8 47
+.b8 100
+.b8 120
+.b8 0
+.b8 1
+.b64 $L__func_begin0
+.b64 $L__func_end0
+.b8 2
+.b8 116
+.b8 114
+.b8 105
+.b8 116
+.b8 111
+.b8 110
+.b8 95
+.b8 95
+.b8 48
+.b8 100
+.b8 49
+.b8 100
+.b8 50
+.b8 100
+.b8 51
+.b8 100
+.b8 52
+.b8 100
+.b8 53
+.b8 100
+.b8 54
+.b8 100
+.b8 55
+.b8 100
+.b8 56
+.b8 100
+.b8 101
+.b8 57
+.b8 100
+.b8 101
+.b8 0
+.b8 116
+.b8 114
+.b8 105
+.b8 116
+.b8 111
+.b8 110
+.b8 95
+.b8 95
+.b8 48
+.b8 100
+.b8 49
+.b8 100
+.b8 50
+.b8 100
+.b8 51
+.b8 100
+.b8 52
+.b8 100
+.b8 53
+.b8 100
+.b8 54
+.b8 100
+.b8 55
+.b8 100
+.b8 56
+.b8 100
+.b8 101
+.b8 57
+.b8 100
+.b8 101
+.b8 0
+.b8 1
+.b8 18
+.b8 1
+.b8 1
+.b8 3
+.b64 $L__func_begin0
+.b64 $L__func_end0
+.b8 1
+.b8 156
+.b32 125
+.b8 4
+.b32 125
+.b64 $L__tmp1
+.b64 $L__tmp14
+.b8 2
+.b8 45
+.b8 59
+.b8 5
+.b32 125
+.b64 $L__tmp1
+.b64 $L__tmp14
+.b8 2
+.b8 243
+.b8 36
+.b8 0
+.b8 5
+.b32 125
+.b64 $L__tmp2
+.b64 $L__tmp15
+.b8 2
+.b8 45
+.b8 59
+.b8 5
+.b32 125
+.b64 $L__tmp15
+.b64 $L__tmp16
+.b8 3
+.b8 45
+.b8 45
+.b8 5
+.b32 125
+.b64 $L__tmp17
+.b64 $L__tmp32
+.b8 2
+.b8 53
+.b8 59
+.b8 4
+.b32 125
+.b64 $L__tmp18
+.b64 $L__tmp31
+.b8 2
+.b8 53
+.b8 59
+.b8 5
+.b32 125
+.b64 $L__tmp18
+.b64 $L__tmp31
+.b8 2
+.b8 243
+.b8 36
+.b8 0
+.b8 5
+.b32 125
+.b64 $L__tmp32
+.b64 $L__tmp33
+.b8 3
+.b8 53
+.b8 45
+.b8 0
+.b8 0
+	}
+	.section	.debug_pubnames
+	{
+.b32 $L__pubNames_end0-$L__pubNames_start0
+$L__pubNames_start0:
+.b8 2
+.b8 0
+.b32 .debug_info
+.b32 411
+.b32 125
+.b8 116
+.b8 114
+.b8 105
+.b8 116
+.b8 111
+.b8 110
+.b8 95
+.b8 95
+.b8 48
+.b8 100
+.b8 49
+.b8 100
+.b8 50
+.b8 100
+.b8 51
+.b8 100
+.b8 52
+.b8 100
+.b8 53
+.b8 100
+.b8 54
+.b8 100
+.b8 55
+.b8 100
+.b8 56
+.b8 100
+.b8 101
+.b8 57
+.b8 100
+.b8 101
+.b8 0
+.b32 0
+$L__pubNames_end0:
+	}
+	.section	.debug_pubtypes
+	{
+.b32 $L__pubTypes_end0-$L__pubTypes_start0
+$L__pubTypes_start0:
+.b8 2
+.b8 0
+.b32 .debug_info
+.b32 411
+.b32 0
+$L__pubTypes_end0:
+	}
+	.section	.debug_loc	{	}