v18/my_solution/msieve/lanczos_kernel.ptx

//
// Generated by NVIDIA NVVM Compiler
//
// Compiler Build ID: CL-32267302
// Cuda compilation tools, release 12.0, V12.0.140
// Based on NVVM 7.0.1
//

.version 8.0
.target sm_70
.address_size 64

	// .globl	lanczos_kernel_mask
// _ZZ25lanczos_kernel_inner_prodE1c has been demoted
// _ZZ25lanczos_kernel_outer_prodE7scratch has been demoted

.visible .entry lanczos_kernel_mask(
	.param .u64 lanczos_kernel_mask_param_0,
	.param .align 8 .b8 lanczos_kernel_mask_param_1[32],
	.param .u32 lanczos_kernel_mask_param_2
)
{
	.reg .pred 	%p<3>;
	.reg .b32 	%r<11>;
	.reg .b64 	%rd<21>;


	ld.param.u64 	%rd6, [lanczos_kernel_mask_param_0];
	ld.param.u32 	%r6, [lanczos_kernel_mask_param_2];
	ld.param.u64 	%rd10, [lanczos_kernel_mask_param_1+24];
	ld.param.u64 	%rd9, [lanczos_kernel_mask_param_1+16];
	ld.param.u64 	%rd8, [lanczos_kernel_mask_param_1+8];
	ld.param.u64 	%rd7, [lanczos_kernel_mask_param_1];
	mov.u32 	%r1, %ntid.x;
	mov.u32 	%r7, %ctaid.x;
	mov.u32 	%r8, %tid.x;
	mad.lo.s32 	%r10, %r7, %r1, %r8;
	setp.ge.u32 	%p1, %r10, %r6;
	@%p1 bra 	$L__BB0_3;

	mov.u32 	%r9, %nctaid.x;
	mul.lo.s32 	%r3, %r9, %r1;
	cvta.to.global.u64 	%rd5, %rd6;

$L__BB0_2:
	mul.wide.u32 	%rd11, %r10, 32;
	add.s64 	%rd12, %rd5, %rd11;
	ld.global.u64 	%rd13, [%rd12];
	and.b64  	%rd14, %rd13, %rd7;
	ld.global.u64 	%rd15, [%rd12+8];
	and.b64  	%rd16, %rd15, %rd8;
	ld.global.u64 	%rd17, [%rd12+16];
	and.b64  	%rd18, %rd17, %rd9;
	ld.global.u64 	%rd19, [%rd12+24];
	and.b64  	%rd20, %rd19, %rd10;
	st.global.u64 	[%rd12], %rd14;
	st.global.u64 	[%rd12+8], %rd16;
	st.global.u64 	[%rd12+16], %rd18;
	st.global.u64 	[%rd12+24], %rd20;
	add.s32 	%r10, %r10, %r3;
	setp.lt.u32 	%p2, %r10, %r6;
	@%p2 bra 	$L__BB0_2;

$L__BB0_3:
	ret;

}
	// .globl	lanczos_kernel_xor
.visible .entry lanczos_kernel_xor(
	.param .u64 lanczos_kernel_xor_param_0,
	.param .u64 lanczos_kernel_xor_param_1,
	.param .u32 lanczos_kernel_xor_param_2
)
{
	.reg .pred 	%p<3>;
	.reg .b32 	%r<11>;
	.reg .b64 	%rd<20>;


	ld.param.u64 	%rd3, [lanczos_kernel_xor_param_0];
	ld.param.u64 	%rd4, [lanczos_kernel_xor_param_1];
	ld.param.u32 	%r6, [lanczos_kernel_xor_param_2];
	mov.u32 	%r1, %ntid.x;
	mov.u32 	%r7, %ctaid.x;
	mov.u32 	%r8, %tid.x;
	mad.lo.s32 	%r10, %r7, %r1, %r8;
	setp.ge.u32 	%p1, %r10, %r6;
	@%p1 bra 	$L__BB1_3;

	mov.u32 	%r9, %nctaid.x;
	mul.lo.s32 	%r3, %r9, %r1;
	cvta.to.global.u64 	%rd1, %rd3;
	cvta.to.global.u64 	%rd2, %rd4;

$L__BB1_2:
	mul.wide.u32 	%rd5, %r10, 32;
	add.s64 	%rd6, %rd1, %rd5;
	add.s64 	%rd7, %rd2, %rd5;
	ld.global.u64 	%rd8, [%rd7];
	ld.global.u64 	%rd9, [%rd6];
	xor.b64  	%rd10, %rd8, %rd9;
	ld.global.u64 	%rd11, [%rd7+8];
	ld.global.u64 	%rd12, [%rd6+8];
	xor.b64  	%rd13, %rd11, %rd12;
	ld.global.u64 	%rd14, [%rd7+16];
	ld.global.u64 	%rd15, [%rd6+16];
	xor.b64  	%rd16, %rd14, %rd15;
	ld.global.u64 	%rd17, [%rd7+24];
	ld.global.u64 	%rd18, [%rd6+24];
	xor.b64  	%rd19, %rd17, %rd18;
	st.global.u64 	[%rd6], %rd10;
	st.global.u64 	[%rd6+8], %rd13;
	st.global.u64 	[%rd6+16], %rd16;
	st.global.u64 	[%rd6+24], %rd19;
	add.s32 	%r10, %r10, %r3;
	setp.lt.u32 	%p2, %r10, %r6;
	@%p2 bra 	$L__BB1_2;

$L__BB1_3:
	ret;

}
	// .globl	lanczos_kernel_inner_prod
.visible .entry lanczos_kernel_inner_prod(
	.param .u64 lanczos_kernel_inner_prod_param_0,
	.param .u64 lanczos_kernel_inner_prod_param_1,
	.param .u64 lanczos_kernel_inner_prod_param_2,
	.param .u32 lanczos_kernel_inner_prod_param_3
)
{
	.local .align 8 .b8 	__local_depot2[32];
	.reg .b64 	%SP;
	.reg .b64 	%SPL;
	.reg .pred 	%p<14>;
	.reg .b32 	%r<100>;
	.reg .b64 	%rd<197>;
	// demoted variable
	.shared .align 8 .b8 _ZZ25lanczos_kernel_inner_prodE1c[12288];

	mov.u64 	%SPL, __local_depot2;
	ld.param.u32 	%r23, [lanczos_kernel_inner_prod_param_3];
	add.u64 	%rd1, %SPL, 0;
	mov.u32 	%r1, %ntid.x;
	mov.u32 	%r2, %tid.x;
	setp.gt.u32 	%p1, %r2, 127;
	@%p1 bra 	$L__BB2_3;

	mov.u32 	%r95, %tid.x;
	ld.param.u64 	%rd159, [lanczos_kernel_inner_prod_param_2];
	cvta.to.global.u64 	%rd2, %rd159;

$L__BB2_2:
	shl.b32 	%r25, %r95, 1;
	mul.wide.u32 	%rd79, %r25, 32;
	add.s64 	%rd80, %rd2, %rd79;
	ld.global.u64 	%rd81, [%rd80];
	ld.global.u64 	%rd82, [%rd80+8];
	ld.global.u64 	%rd83, [%rd80+16];
	ld.global.u64 	%rd84, [%rd80+24];
	mov.u32 	%r26, _ZZ25lanczos_kernel_inner_prodE1c;
	mad.lo.s32 	%r27, %r95, 96, %r26;
	st.shared.u64 	[%r27], %rd81;
	st.shared.u64 	[%r27+8], %rd82;
	st.shared.u64 	[%r27+16], %rd83;
	st.shared.u64 	[%r27+24], %rd84;
	or.b32  	%r28, %r25, 1;
	mul.wide.u32 	%rd85, %r28, 32;
	add.s64 	%rd86, %rd2, %rd85;
	ld.global.u64 	%rd87, [%rd86];
	xor.b64  	%rd88, %rd87, %rd81;
	ld.global.u64 	%rd89, [%rd86+8];
	xor.b64  	%rd90, %rd89, %rd82;
	ld.global.u64 	%rd91, [%rd86+16];
	xor.b64  	%rd92, %rd91, %rd83;
	ld.global.u64 	%rd93, [%rd86+24];
	xor.b64  	%rd94, %rd93, %rd84;
	st.shared.u64 	[%r27+64], %rd88;
	st.shared.u64 	[%r27+72], %rd90;
	st.shared.u64 	[%r27+80], %rd92;
	st.shared.u64 	[%r27+88], %rd94;
	st.shared.u64 	[%r27+32], %rd87;
	st.shared.u64 	[%r27+40], %rd89;
	st.shared.u64 	[%r27+48], %rd91;
	st.shared.u64 	[%r27+56], %rd93;
	add.s32 	%r95, %r95, %r1;
	setp.lt.u32 	%p2, %r95, 128;
	@%p2 bra 	$L__BB2_2;

$L__BB2_3:
	mov.u32 	%r93, %tid.x;
	mov.u32 	%r90, %ctaid.x;
	mad.lo.s32 	%r89, %r90, %r1, %r93;
	bar.sync 	0;
	setp.ge.u32 	%p3, %r89, %r23;
	@%p3 bra 	$L__BB2_24;

	mov.u32 	%r92, %ctaid.x;
	mad.lo.s32 	%r96, %r92, %r1, %r2;
	ld.param.u64 	%rd158, [lanczos_kernel_inner_prod_param_1];
	ld.param.u64 	%rd157, [lanczos_kernel_inner_prod_param_0];
	cvta.to.global.u64 	%rd3, %rd157;
	cvta.to.global.u64 	%rd4, %rd158;
	mov.u32 	%r29, %nctaid.x;
	mul.lo.s32 	%r6, %r29, %r1;

$L__BB2_5:
	mul.wide.u32 	%rd99, %r96, 32;
	add.s64 	%rd100, %rd4, %rd99;
	ld.global.u64 	%rd101, [%rd100];
	mov.u64 	%rd165, 0;
	ld.global.u64 	%rd102, [%rd100+8];
	ld.global.u64 	%rd103, [%rd100+16];
	ld.global.u64 	%rd104, [%rd100+24];
	st.local.u64 	[%rd1], %rd101;
	st.local.u64 	[%rd1+8], %rd102;
	st.local.u64 	[%rd1+16], %rd103;
	st.local.u64 	[%rd1+24], %rd104;
	mov.u32 	%r33, _ZZ25lanczos_kernel_inner_prodE1c;
	add.s32 	%r98, %r33, 664;
	mov.u32 	%r99, 352;
	mov.u32 	%r97, 7;
	mov.u64 	%rd166, %rd165;
	mov.u64 	%rd167, %rd165;
	mov.u64 	%rd168, %rd165;

$L__BB2_6:
	add.s32 	%r34, %r97, -7;
	shr.u32 	%r35, %r34, 5;
	mul.wide.u32 	%rd105, %r35, 8;
	add.s64 	%rd106, %rd1, %rd105;
	shl.b32 	%r36, %r34, 1;
	and.b32  	%r37, %r36, 48;
	ld.local.u64 	%rd10, [%rd106];
	shr.u64 	%rd107, %rd10, %r37;
	cvt.u32.u64 	%r38, %rd107;
	and.b32  	%r11, %r38, 3;
	setp.eq.s32 	%p4, %r11, 0;
	@%p4 bra 	$L__BB2_8;

	shl.b32 	%r39, %r11, 5;
	add.s32 	%r40, %r98, %r39;
	ld.shared.u64 	%rd108, [%r40+-696];
	xor.b64  	%rd168, %rd108, %rd168;
	ld.shared.u64 	%rd109, [%r40+-688];
	xor.b64  	%rd167, %rd109, %rd167;
	ld.shared.u64 	%rd110, [%r40+-680];
	xor.b64  	%rd166, %rd110, %rd166;
	ld.shared.u64 	%rd111, [%r40+-672];
	xor.b64  	%rd165, %rd111, %rd165;

$L__BB2_8:
	shl.b32 	%r41, %r97, 1;
	add.s32 	%r42, %r41, -12;
	and.b32  	%r43, %r42, 50;
	shr.u64 	%rd112, %rd10, %r43;
	cvt.u32.u64 	%r44, %rd112;
	and.b32  	%r12, %r44, 3;
	setp.eq.s32 	%p5, %r12, 0;
	@%p5 bra 	$L__BB2_10;

	shl.b32 	%r45, %r12, 5;
	add.s32 	%r46, %r98, %r45;
	ld.shared.u64 	%rd113, [%r46+-600];
	xor.b64  	%rd168, %rd113, %rd168;
	ld.shared.u64 	%rd114, [%r46+-592];
	xor.b64  	%rd167, %rd114, %rd167;
	ld.shared.u64 	%rd115, [%r46+-584];
	xor.b64  	%rd166, %rd115, %rd166;
	ld.shared.u64 	%rd116, [%r46+-576];
	xor.b64  	%rd165, %rd116, %rd165;

$L__BB2_10:
	shl.b32 	%r83, %r97, 1;
	add.s32 	%r48, %r83, -10;
	and.b32  	%r49, %r48, 52;
	shr.u64 	%rd117, %rd10, %r49;
	cvt.u32.u64 	%r50, %rd117;
	and.b32  	%r13, %r50, 3;
	setp.eq.s32 	%p6, %r13, 0;
	@%p6 bra 	$L__BB2_12;

	shl.b32 	%r51, %r13, 5;
	add.s32 	%r52, %r98, %r51;
	ld.shared.u64 	%rd118, [%r52+-504];
	xor.b64  	%rd168, %rd118, %rd168;
	ld.shared.u64 	%rd119, [%r52+-496];
	xor.b64  	%rd167, %rd119, %rd167;
	ld.shared.u64 	%rd120, [%r52+-488];
	xor.b64  	%rd166, %rd120, %rd166;
	ld.shared.u64 	%rd121, [%r52+-480];
	xor.b64  	%rd165, %rd121, %rd165;

$L__BB2_12:
	shl.b32 	%r84, %r97, 1;
	add.s32 	%r54, %r84, -8;
	and.b32  	%r55, %r54, 54;
	shr.u64 	%rd122, %rd10, %r55;
	cvt.u32.u64 	%r56, %rd122;
	and.b32  	%r14, %r56, 3;
	setp.eq.s32 	%p7, %r14, 0;
	@%p7 bra 	$L__BB2_14;

	shl.b32 	%r57, %r14, 5;
	add.s32 	%r58, %r98, %r57;
	ld.shared.u64 	%rd123, [%r58+-408];
	xor.b64  	%rd168, %rd123, %rd168;
	ld.shared.u64 	%rd124, [%r58+-400];
	xor.b64  	%rd167, %rd124, %rd167;
	ld.shared.u64 	%rd125, [%r58+-392];
	xor.b64  	%rd166, %rd125, %rd166;
	ld.shared.u64 	%rd126, [%r58+-384];
	xor.b64  	%rd165, %rd126, %rd165;

$L__BB2_14:
	shl.b32 	%r85, %r97, 1;
	add.s32 	%r60, %r85, -6;
	and.b32  	%r61, %r60, 56;
	shr.u64 	%rd127, %rd10, %r61;
	cvt.u32.u64 	%r62, %rd127;
	and.b32  	%r15, %r62, 3;
	setp.eq.s32 	%p8, %r15, 0;
	@%p8 bra 	$L__BB2_16;

	shl.b32 	%r63, %r15, 5;
	add.s32 	%r64, %r98, %r63;
	ld.shared.u64 	%rd128, [%r64+-312];
	xor.b64  	%rd168, %rd128, %rd168;
	ld.shared.u64 	%rd129, [%r64+-304];
	xor.b64  	%rd167, %rd129, %rd167;
	ld.shared.u64 	%rd130, [%r64+-296];
	xor.b64  	%rd166, %rd130, %rd166;
	ld.shared.u64 	%rd131, [%r64+-288];
	xor.b64  	%rd165, %rd131, %rd165;

$L__BB2_16:
	shl.b32 	%r86, %r97, 1;
	add.s32 	%r66, %r86, -4;
	and.b32  	%r67, %r66, 58;
	shr.u64 	%rd132, %rd10, %r67;
	cvt.u32.u64 	%r68, %rd132;
	and.b32  	%r16, %r68, 3;
	setp.eq.s32 	%p9, %r16, 0;
	@%p9 bra 	$L__BB2_18;

	shl.b32 	%r69, %r16, 5;
	add.s32 	%r70, %r98, %r69;
	ld.shared.u64 	%rd133, [%r70+-216];
	xor.b64  	%rd168, %rd133, %rd168;
	ld.shared.u64 	%rd134, [%r70+-208];
	xor.b64  	%rd167, %rd134, %rd167;
	ld.shared.u64 	%rd135, [%r70+-200];
	xor.b64  	%rd166, %rd135, %rd166;
	ld.shared.u64 	%rd136, [%r70+-192];
	xor.b64  	%rd165, %rd136, %rd165;

$L__BB2_18:
	shl.b32 	%r87, %r97, 1;
	add.s32 	%r72, %r87, -2;
	and.b32  	%r73, %r72, 60;
	shr.u64 	%rd137, %rd10, %r73;
	cvt.u32.u64 	%r74, %rd137;
	and.b32  	%r17, %r74, 3;
	setp.eq.s32 	%p10, %r17, 0;
	@%p10 bra 	$L__BB2_20;

	shl.b32 	%r75, %r17, 5;
	add.s32 	%r76, %r98, %r75;
	ld.shared.u64 	%rd138, [%r76+-120];
	xor.b64  	%rd168, %rd138, %rd168;
	ld.shared.u64 	%rd139, [%r76+-112];
	xor.b64  	%rd167, %rd139, %rd167;
	ld.shared.u64 	%rd140, [%r76+-104];
	xor.b64  	%rd166, %rd140, %rd166;
	ld.shared.u64 	%rd141, [%r76+-96];
	xor.b64  	%rd165, %rd141, %rd165;

$L__BB2_20:
	shl.b32 	%r88, %r97, 1;
	and.b32  	%r78, %r88, 62;
	shr.u64 	%rd142, %rd10, %r78;
	cvt.u32.u64 	%r79, %rd142;
	and.b32  	%r18, %r79, 3;
	setp.eq.s32 	%p11, %r18, 0;
	@%p11 bra 	$L__BB2_22;

	shl.b32 	%r80, %r18, 5;
	add.s32 	%r81, %r98, %r80;
	ld.shared.u64 	%rd143, [%r81+-24];
	xor.b64  	%rd168, %rd143, %rd168;
	ld.shared.u64 	%rd144, [%r81+-16];
	xor.b64  	%rd167, %rd144, %rd167;
	ld.shared.u64 	%rd145, [%r81+-8];
	xor.b64  	%rd166, %rd145, %rd166;
	ld.shared.u64 	%rd146, [%r81];
	xor.b64  	%rd165, %rd146, %rd165;

$L__BB2_22:
	add.s32 	%r98, %r98, 768;
	add.s32 	%r97, %r97, 8;
	add.s32 	%r99, %r99, 768;
	setp.ne.s32 	%p12, %r99, 12640;
	@%p12 bra 	$L__BB2_6;

	cvt.u64.u32 	%rd160, %r96;
	cvt.u32.u64 	%r82, %rd160;
	shl.b64 	%rd147, %rd160, 5;
	add.s64 	%rd148, %rd3, %rd147;
	ld.global.u64 	%rd149, [%rd148];
	xor.b64  	%rd150, %rd168, %rd149;
	ld.global.u64 	%rd151, [%rd148+8];
	xor.b64  	%rd152, %rd167, %rd151;
	ld.global.u64 	%rd153, [%rd148+16];
	xor.b64  	%rd154, %rd166, %rd153;
	ld.global.u64 	%rd155, [%rd148+24];
	xor.b64  	%rd156, %rd165, %rd155;
	st.global.u64 	[%rd148], %rd150;
	st.global.u64 	[%rd148+8], %rd152;
	st.global.u64 	[%rd148+16], %rd154;
	st.global.u64 	[%rd148+24], %rd156;
	add.s32 	%r96, %r82, %r6;
	setp.lt.u32 	%p13, %r96, %r23;
	@%p13 bra 	$L__BB2_5;

$L__BB2_24:
	ret;

}
	// .globl	lanczos_kernel_outer_prod
.visible .entry lanczos_kernel_outer_prod(
	.param .u64 lanczos_kernel_outer_prod_param_0,
	.param .u64 lanczos_kernel_outer_prod_param_1,
	.param .u64 lanczos_kernel_outer_prod_param_2,
	.param .u32 lanczos_kernel_outer_prod_param_3
)
{
	.reg .pred 	%p<43>;
	.reg .b32 	%r<454>;
	.reg .b64 	%rd<160>;
	// demoted variable
	.shared .align 8 .b8 _ZZ25lanczos_kernel_outer_prodE7scratch[6144];

	ld.param.u64 	%rd10, [lanczos_kernel_outer_prod_param_0];
	ld.param.u64 	%rd11, [lanczos_kernel_outer_prod_param_1];
	ld.param.u64 	%rd12, [lanczos_kernel_outer_prod_param_2];
	ld.param.u32 	%r46, [lanczos_kernel_outer_prod_param_3];
	cvta.to.global.u64 	%rd1, %rd11;
	cvta.to.global.u64 	%rd2, %rd10;
	cvta.to.global.u64 	%rd3, %rd12;
	mov.u32 	%r48, %ntid.x;
	mov.u32 	%r49, %nctaid.x;
	mul.lo.s32 	%r1, %r49, %r48;
	mov.u32 	%r50, %ctaid.x;
	mov.u32 	%r2, %tid.x;
	mad.lo.s32 	%r3, %r50, %r48, %r2;
	and.b32  	%r51, %r2, -32;
	shl.b32 	%r52, %r2, 3;
	mov.u32 	%r53, _ZZ25lanczos_kernel_outer_prodE7scratch;
	add.s32 	%r4, %r53, %r52;
	and.b32  	%r5, %r2, 31;
	shl.b32 	%r6, %r2, 1;
	and.b32  	%r54, %r6, 62;
	cvt.u64.u32 	%rd4, %r54;
	mov.u32 	%r55, 64;
	sub.s32 	%r56, %r55, %r54;
	cvt.u64.u32 	%rd5, %r56;
	or.b32  	%r7, %r5, %r51;
	add.s32 	%r57, %r2, 1;
	and.b32  	%r58, %r57, 31;
	or.b32  	%r8, %r58, %r51;
	add.s32 	%r59, %r2, 2;
	and.b32  	%r60, %r59, 31;
	or.b32  	%r9, %r60, %r51;
	add.s32 	%r61, %r2, 3;
	and.b32  	%r62, %r61, 31;
	or.b32  	%r10, %r62, %r51;
	add.s32 	%r63, %r2, 4;
	and.b32  	%r64, %r63, 31;
	or.b32  	%r11, %r64, %r51;
	add.s32 	%r65, %r2, 5;
	and.b32  	%r66, %r65, 31;
	or.b32  	%r12, %r66, %r51;
	add.s32 	%r67, %r2, 6;
	and.b32  	%r68, %r67, 31;
	or.b32  	%r13, %r68, %r51;
	add.s32 	%r69, %r2, 7;
	and.b32  	%r70, %r69, 31;
	or.b32  	%r14, %r70, %r51;
	add.s32 	%r71, %r2, 8;
	and.b32  	%r72, %r71, 31;
	or.b32  	%r15, %r72, %r51;
	add.s32 	%r73, %r2, 9;
	and.b32  	%r74, %r73, 31;
	or.b32  	%r16, %r74, %r51;
	add.s32 	%r75, %r2, 10;
	and.b32  	%r76, %r75, 31;
	or.b32  	%r17, %r76, %r51;
	add.s32 	%r77, %r2, 11;
	and.b32  	%r78, %r77, 31;
	or.b32  	%r18, %r78, %r51;
	add.s32 	%r79, %r2, 12;
	and.b32  	%r80, %r79, 31;
	or.b32  	%r19, %r80, %r51;
	add.s32 	%r81, %r2, 13;
	and.b32  	%r82, %r81, 31;
	or.b32  	%r20, %r82, %r51;
	add.s32 	%r83, %r2, 14;
	and.b32  	%r84, %r83, 31;
	or.b32  	%r21, %r84, %r51;
	add.s32 	%r85, %r2, 15;
	and.b32  	%r86, %r85, 31;
	or.b32  	%r22, %r86, %r51;
	add.s32 	%r87, %r2, 16;
	and.b32  	%r88, %r87, 31;
	or.b32  	%r23, %r88, %r51;
	add.s32 	%r89, %r2, 17;
	and.b32  	%r90, %r89, 31;
	or.b32  	%r24, %r90, %r51;
	add.s32 	%r91, %r2, 18;
	and.b32  	%r92, %r91, 31;
	or.b32  	%r25, %r92, %r51;
	add.s32 	%r93, %r2, 19;
	and.b32  	%r94, %r93, 31;
	or.b32  	%r26, %r94, %r51;
	add.s32 	%r95, %r2, 20;
	and.b32  	%r96, %r95, 31;
	or.b32  	%r27, %r96, %r51;
	add.s32 	%r97, %r2, 21;
	and.b32  	%r98, %r97, 31;
	or.b32  	%r28, %r98, %r51;
	add.s32 	%r99, %r2, 22;
	and.b32  	%r100, %r99, 31;
	or.b32  	%r29, %r100, %r51;
	add.s32 	%r101, %r2, 23;
	and.b32  	%r102, %r101, 31;
	or.b32  	%r30, %r102, %r51;
	add.s32 	%r103, %r2, 24;
	and.b32  	%r104, %r103, 31;
	or.b32  	%r31, %r104, %r51;
	add.s32 	%r105, %r2, 25;
	and.b32  	%r106, %r105, 31;
	or.b32  	%r32, %r106, %r51;
	add.s32 	%r107, %r2, 26;
	and.b32  	%r108, %r107, 31;
	or.b32  	%r33, %r108, %r51;
	add.s32 	%r109, %r2, 27;
	and.b32  	%r110, %r109, 31;
	or.b32  	%r34, %r110, %r51;
	add.s32 	%r111, %r2, 28;
	and.b32  	%r112, %r111, 31;
	or.b32  	%r35, %r112, %r51;
	add.s32 	%r113, %r2, 29;
	and.b32  	%r114, %r113, 31;
	or.b32  	%r36, %r114, %r51;
	add.s32 	%r115, %r2, 30;
	and.b32  	%r116, %r115, 31;
	or.b32  	%r37, %r116, %r51;
	add.s32 	%r117, %r2, -1;
	and.b32  	%r118, %r117, 31;
	or.b32  	%r38, %r118, %r51;
	mov.u32 	%r451, 0;

$L__BB3_1:
	mov.u32 	%r452, 0;
	mov.u32 	%r448, %tid.x;
	shl.b32 	%r447, %r448, 1;
	cvt.u64.u32 	%rd6, %r451;
	shl.b32 	%r120, %r451, 6;
	add.s32 	%r121, %r120, %r447;
	cvt.u64.u32 	%rd7, %r121;
	add.s32 	%r122, %r121, -63;
	cvt.u64.u32 	%rd8, %r122;

$L__BB3_2:
	mov.u64 	%rd13, 0;
	st.shared.u64 	[%r4], %rd13;
	st.shared.u64 	[%r4+2048], %rd13;
	st.shared.u64 	[%r4+4096], %rd13;
	setp.ge.u32 	%p1, %r3, %r46;
	@%p1 bra 	$L__BB3_5;

	mov.u32 	%r453, %r3;

$L__BB3_4:
	mov.u32 	%r450, _ZZ25lanczos_kernel_outer_prodE7scratch;
	cvt.u64.u32 	%rd159, %r452;
	mul.wide.u32 	%rd14, %r453, 32;
	add.s64 	%rd15, %rd2, %rd14;
	shl.b64 	%rd16, %rd6, 3;
	add.s64 	%rd17, %rd15, %rd16;
	add.s64 	%rd18, %rd1, %rd14;
	shl.b64 	%rd19, %rd159, 3;
	add.s64 	%rd20, %rd18, %rd19;
	ld.global.u64 	%rd21, [%rd20];
	ld.global.u64 	%rd22, [%rd17];
	cvt.u32.u64 	%r251, %rd4;
	shr.u64 	%rd23, %rd22, %r251;
	cvt.u32.u64 	%r252, %rd5;
	shl.b64 	%rd24, %rd22, %r252;
	or.b64  	%rd25, %rd23, %rd24;
	setp.eq.s32 	%p2, %r5, 0;
	mov.u32 	%r189, 0;
	selp.b64 	%rd26, %rd22, %rd25, %p2;
	cvt.u32.u64 	%r184, %rd26;
	mov.u32 	%r250, 2;
	// begin inline asm
	bfe.u32 %r123, %r184, %r189, %r250; 
	
	// end inline asm
	setp.eq.s32 	%p3, %r123, 0;
	selp.b64 	%rd27, 0, %rd21, %p3;
	shl.b32 	%r253, %r123, 8;
	mov.u32 	%r205, 8;
	add.s32 	%r254, %r253, -256;
	selp.b32 	%r255, 0, %r254, %p3;
	add.s32 	%r256, %r7, %r255;
	shl.b32 	%r257, %r256, 3;
	add.s32 	%r259, %r450, %r257;
	ld.shared.u64 	%rd28, [%r259];
	xor.b64  	%rd29, %rd27, %rd28;
	st.shared.u64 	[%r259], %rd29;
	// begin inline asm
	bfe.u32 %r127, %r184, %r250, %r250; 
	
	// end inline asm
	setp.eq.s32 	%p4, %r127, 0;
	selp.b64 	%rd30, 0, %rd21, %p4;
	shl.b32 	%r260, %r127, 8;
	add.s32 	%r261, %r260, -256;
	selp.b32 	%r262, 0, %r261, %p4;
	add.s32 	%r263, %r8, %r262;
	shl.b32 	%r264, %r263, 3;
	add.s32 	%r265, %r450, %r264;
	ld.shared.u64 	%rd31, [%r265];
	xor.b64  	%rd32, %rd30, %rd31;
	st.shared.u64 	[%r265], %rd32;
	mov.u32 	%r197, 4;
	// begin inline asm
	bfe.u32 %r131, %r184, %r197, %r250; 
	
	// end inline asm
	setp.eq.s32 	%p5, %r131, 0;
	selp.b64 	%rd33, 0, %rd21, %p5;
	shl.b32 	%r266, %r131, 8;
	add.s32 	%r267, %r266, -256;
	selp.b32 	%r268, 0, %r267, %p5;
	add.s32 	%r269, %r9, %r268;
	shl.b32 	%r270, %r269, 3;
	add.s32 	%r271, %r450, %r270;
	ld.shared.u64 	%rd34, [%r271];
	xor.b64  	%rd35, %rd33, %rd34;
	st.shared.u64 	[%r271], %rd35;
	mov.u32 	%r201, 6;
	// begin inline asm
	bfe.u32 %r135, %r184, %r201, %r250; 
	
	// end inline asm
	setp.eq.s32 	%p6, %r135, 0;
	selp.b64 	%rd36, 0, %rd21, %p6;
	shl.b32 	%r272, %r135, 8;
	add.s32 	%r273, %r272, -256;
	selp.b32 	%r274, 0, %r273, %p6;
	add.s32 	%r275, %r10, %r274;
	shl.b32 	%r276, %r275, 3;
	add.s32 	%r277, %r450, %r276;
	ld.shared.u64 	%rd37, [%r277];
	xor.b64  	%rd38, %rd36, %rd37;
	st.shared.u64 	[%r277], %rd38;
	// begin inline asm
	bfe.u32 %r139, %r184, %r205, %r250; 
	
	// end inline asm
	setp.eq.s32 	%p7, %r139, 0;
	selp.b64 	%rd39, 0, %rd21, %p7;
	shl.b32 	%r278, %r139, 8;
	add.s32 	%r279, %r278, -256;
	selp.b32 	%r280, 0, %r279, %p7;
	add.s32 	%r281, %r11, %r280;
	shl.b32 	%r282, %r281, 3;
	add.s32 	%r283, %r450, %r282;
	ld.shared.u64 	%rd40, [%r283];
	xor.b64  	%rd41, %rd39, %rd40;
	st.shared.u64 	[%r283], %rd41;
	mov.u32 	%r209, 10;
	// begin inline asm
	bfe.u32 %r143, %r184, %r209, %r250; 
	
	// end inline asm
	setp.eq.s32 	%p8, %r143, 0;
	selp.b64 	%rd42, 0, %rd21, %p8;
	shl.b32 	%r284, %r143, 8;
	add.s32 	%r285, %r284, -256;
	selp.b32 	%r286, 0, %r285, %p8;
	add.s32 	%r287, %r12, %r286;
	shl.b32 	%r288, %r287, 3;
	add.s32 	%r289, %r450, %r288;
	ld.shared.u64 	%rd43, [%r289];
	xor.b64  	%rd44, %rd42, %rd43;
	st.shared.u64 	[%r289], %rd44;
	mov.u32 	%r213, 12;
	// begin inline asm
	bfe.u32 %r147, %r184, %r213, %r250; 
	
	// end inline asm
	setp.eq.s32 	%p9, %r147, 0;
	selp.b64 	%rd45, 0, %rd21, %p9;
	shl.b32 	%r290, %r147, 8;
	add.s32 	%r291, %r290, -256;
	selp.b32 	%r292, 0, %r291, %p9;
	add.s32 	%r293, %r13, %r292;
	shl.b32 	%r294, %r293, 3;
	add.s32 	%r295, %r450, %r294;
	ld.shared.u64 	%rd46, [%r295];
	xor.b64  	%rd47, %rd45, %rd46;
	st.shared.u64 	[%r295], %rd47;
	mov.u32 	%r217, 14;
	// begin inline asm
	bfe.u32 %r151, %r184, %r217, %r250; 
	
	// end inline asm
	setp.eq.s32 	%p10, %r151, 0;
	selp.b64 	%rd48, 0, %rd21, %p10;
	shl.b32 	%r296, %r151, 8;
	add.s32 	%r297, %r296, -256;
	selp.b32 	%r298, 0, %r297, %p10;
	add.s32 	%r299, %r14, %r298;
	shl.b32 	%r300, %r299, 3;
	add.s32 	%r301, %r450, %r300;
	ld.shared.u64 	%rd49, [%r301];
	xor.b64  	%rd50, %rd48, %rd49;
	st.shared.u64 	[%r301], %rd50;
	mov.u32 	%r221, 16;
	// begin inline asm
	bfe.u32 %r155, %r184, %r221, %r250; 
	
	// end inline asm
	setp.eq.s32 	%p11, %r155, 0;
	selp.b64 	%rd51, 0, %rd21, %p11;
	shl.b32 	%r302, %r155, 8;
	add.s32 	%r303, %r302, -256;
	selp.b32 	%r304, 0, %r303, %p11;
	add.s32 	%r305, %r15, %r304;
	shl.b32 	%r306, %r305, 3;
	add.s32 	%r307, %r450, %r306;
	ld.shared.u64 	%rd52, [%r307];
	xor.b64  	%rd53, %rd51, %rd52;
	st.shared.u64 	[%r307], %rd53;
	mov.u32 	%r225, 18;
	// begin inline asm
	bfe.u32 %r159, %r184, %r225, %r250; 
	
	// end inline asm
	setp.eq.s32 	%p12, %r159, 0;
	selp.b64 	%rd54, 0, %rd21, %p12;
	shl.b32 	%r308, %r159, 8;
	add.s32 	%r309, %r308, -256;
	selp.b32 	%r310, 0, %r309, %p12;
	add.s32 	%r311, %r16, %r310;
	shl.b32 	%r312, %r311, 3;
	add.s32 	%r313, %r450, %r312;
	ld.shared.u64 	%rd55, [%r313];
	xor.b64  	%rd56, %rd54, %rd55;
	st.shared.u64 	[%r313], %rd56;
	mov.u32 	%r229, 20;
	// begin inline asm
	bfe.u32 %r163, %r184, %r229, %r250; 
	
	// end inline asm
	setp.eq.s32 	%p13, %r163, 0;
	selp.b64 	%rd57, 0, %rd21, %p13;
	shl.b32 	%r314, %r163, 8;
	add.s32 	%r315, %r314, -256;
	selp.b32 	%r316, 0, %r315, %p13;
	add.s32 	%r317, %r17, %r316;
	shl.b32 	%r318, %r317, 3;
	add.s32 	%r319, %r450, %r318;
	ld.shared.u64 	%rd58, [%r319];
	xor.b64  	%rd59, %rd57, %rd58;
	st.shared.u64 	[%r319], %rd59;
	mov.u32 	%r233, 22;
	// begin inline asm
	bfe.u32 %r167, %r184, %r233, %r250; 
	
	// end inline asm
	setp.eq.s32 	%p14, %r167, 0;
	selp.b64 	%rd60, 0, %rd21, %p14;
	shl.b32 	%r320, %r167, 8;
	add.s32 	%r321, %r320, -256;
	selp.b32 	%r322, 0, %r321, %p14;
	add.s32 	%r323, %r18, %r322;
	shl.b32 	%r324, %r323, 3;
	add.s32 	%r325, %r450, %r324;
	ld.shared.u64 	%rd61, [%r325];
	xor.b64  	%rd62, %rd60, %rd61;
	st.shared.u64 	[%r325], %rd62;
	mov.u32 	%r237, 24;
	// begin inline asm
	bfe.u32 %r171, %r184, %r237, %r250; 
	
	// end inline asm
	setp.eq.s32 	%p15, %r171, 0;
	selp.b64 	%rd63, 0, %rd21, %p15;
	shl.b32 	%r326, %r171, 8;
	add.s32 	%r327, %r326, -256;
	selp.b32 	%r328, 0, %r327, %p15;
	add.s32 	%r329, %r19, %r328;
	shl.b32 	%r330, %r329, 3;
	add.s32 	%r331, %r450, %r330;
	ld.shared.u64 	%rd64, [%r331];
	xor.b64  	%rd65, %rd63, %rd64;
	st.shared.u64 	[%r331], %rd65;
	mov.u32 	%r241, 26;
	// begin inline asm
	bfe.u32 %r175, %r184, %r241, %r250; 
	
	// end inline asm
	setp.eq.s32 	%p16, %r175, 0;
	selp.b64 	%rd66, 0, %rd21, %p16;
	shl.b32 	%r332, %r175, 8;
	add.s32 	%r333, %r332, -256;
	selp.b32 	%r334, 0, %r333, %p16;
	add.s32 	%r335, %r20, %r334;
	shl.b32 	%r336, %r335, 3;
	add.s32 	%r337, %r450, %r336;
	ld.shared.u64 	%rd67, [%r337];
	xor.b64  	%rd68, %rd66, %rd67;
	st.shared.u64 	[%r337], %rd68;
	mov.u32 	%r245, 28;
	// begin inline asm
	bfe.u32 %r179, %r184, %r245, %r250; 
	
	// end inline asm
	setp.eq.s32 	%p17, %r179, 0;
	selp.b64 	%rd69, 0, %rd21, %p17;
	shl.b32 	%r338, %r179, 8;
	add.s32 	%r339, %r338, -256;
	selp.b32 	%r340, 0, %r339, %p17;
	add.s32 	%r341, %r21, %r340;
	shl.b32 	%r342, %r341, 3;
	add.s32 	%r343, %r450, %r342;
	ld.shared.u64 	%rd70, [%r343];
	xor.b64  	%rd71, %rd69, %rd70;
	st.shared.u64 	[%r343], %rd71;
	mov.u32 	%r249, 30;
	// begin inline asm
	bfe.u32 %r183, %r184, %r249, %r250; 
	
	// end inline asm
	setp.eq.s32 	%p18, %r183, 0;
	selp.b64 	%rd72, 0, %rd21, %p18;
	shl.b32 	%r344, %r183, 8;
	add.s32 	%r345, %r344, -256;
	selp.b32 	%r346, 0, %r345, %p18;
	add.s32 	%r347, %r22, %r346;
	shl.b32 	%r348, %r347, 3;
	add.s32 	%r349, %r450, %r348;
	ld.shared.u64 	%rd73, [%r349];
	xor.b64  	%rd74, %rd72, %rd73;
	st.shared.u64 	[%r349], %rd74;
	shr.u64 	%rd75, %rd26, 32;
	cvt.u32.u64 	%r248, %rd75;
	// begin inline asm
	bfe.u32 %r187, %r248, %r189, %r250; 
	
	// end inline asm
	setp.eq.s32 	%p19, %r187, 0;
	selp.b64 	%rd76, 0, %rd21, %p19;
	shl.b32 	%r350, %r187, 8;
	add.s32 	%r351, %r350, -256;
	selp.b32 	%r352, 0, %r351, %p19;
	add.s32 	%r353, %r23, %r352;
	shl.b32 	%r354, %r353, 3;
	add.s32 	%r355, %r450, %r354;
	ld.shared.u64 	%rd77, [%r355];
	xor.b64  	%rd78, %rd76, %rd77;
	st.shared.u64 	[%r355], %rd78;
	// begin inline asm
	bfe.u32 %r191, %r248, %r250, %r250; 
	
	// end inline asm
	setp.eq.s32 	%p20, %r191, 0;
	selp.b64 	%rd79, 0, %rd21, %p20;
	shl.b32 	%r356, %r191, 8;
	add.s32 	%r357, %r356, -256;
	selp.b32 	%r358, 0, %r357, %p20;
	add.s32 	%r359, %r24, %r358;
	shl.b32 	%r360, %r359, 3;
	add.s32 	%r361, %r450, %r360;
	ld.shared.u64 	%rd80, [%r361];
	xor.b64  	%rd81, %rd79, %rd80;
	st.shared.u64 	[%r361], %rd81;
	// begin inline asm
	bfe.u32 %r195, %r248, %r197, %r250; 
	
	// end inline asm
	setp.eq.s32 	%p21, %r195, 0;
	selp.b64 	%rd82, 0, %rd21, %p21;
	shl.b32 	%r362, %r195, 8;
	add.s32 	%r363, %r362, -256;
	selp.b32 	%r364, 0, %r363, %p21;
	add.s32 	%r365, %r25, %r364;
	shl.b32 	%r366, %r365, 3;
	add.s32 	%r367, %r450, %r366;
	ld.shared.u64 	%rd83, [%r367];
	xor.b64  	%rd84, %rd82, %rd83;
	st.shared.u64 	[%r367], %rd84;
	// begin inline asm
	bfe.u32 %r199, %r248, %r201, %r250; 
	
	// end inline asm
	setp.eq.s32 	%p22, %r199, 0;
	selp.b64 	%rd85, 0, %rd21, %p22;
	shl.b32 	%r368, %r199, 8;
	add.s32 	%r369, %r368, -256;
	selp.b32 	%r370, 0, %r369, %p22;
	add.s32 	%r371, %r26, %r370;
	shl.b32 	%r372, %r371, 3;
	add.s32 	%r373, %r450, %r372;
	ld.shared.u64 	%rd86, [%r373];
	xor.b64  	%rd87, %rd85, %rd86;
	st.shared.u64 	[%r373], %rd87;
	// begin inline asm
	bfe.u32 %r203, %r248, %r205, %r250; 
	
	// end inline asm
	setp.eq.s32 	%p23, %r203, 0;
	selp.b64 	%rd88, 0, %rd21, %p23;
	shl.b32 	%r374, %r203, 8;
	add.s32 	%r375, %r374, -256;
	selp.b32 	%r376, 0, %r375, %p23;
	add.s32 	%r377, %r27, %r376;
	shl.b32 	%r378, %r377, 3;
	add.s32 	%r379, %r450, %r378;
	ld.shared.u64 	%rd89, [%r379];
	xor.b64  	%rd90, %rd88, %rd89;
	st.shared.u64 	[%r379], %rd90;
	// begin inline asm
	bfe.u32 %r207, %r248, %r209, %r250; 
	
	// end inline asm
	setp.eq.s32 	%p24, %r207, 0;
	selp.b64 	%rd91, 0, %rd21, %p24;
	shl.b32 	%r380, %r207, 8;
	add.s32 	%r381, %r380, -256;
	selp.b32 	%r382, 0, %r381, %p24;
	add.s32 	%r383, %r28, %r382;
	shl.b32 	%r384, %r383, 3;
	add.s32 	%r385, %r450, %r384;
	ld.shared.u64 	%rd92, [%r385];
	xor.b64  	%rd93, %rd91, %rd92;
	st.shared.u64 	[%r385], %rd93;
	// begin inline asm
	bfe.u32 %r211, %r248, %r213, %r250; 
	
	// end inline asm
	setp.eq.s32 	%p25, %r211, 0;
	selp.b64 	%rd94, 0, %rd21, %p25;
	shl.b32 	%r386, %r211, 8;
	add.s32 	%r387, %r386, -256;
	selp.b32 	%r388, 0, %r387, %p25;
	add.s32 	%r389, %r29, %r388;
	shl.b32 	%r390, %r389, 3;
	add.s32 	%r391, %r450, %r390;
	ld.shared.u64 	%rd95, [%r391];
	xor.b64  	%rd96, %rd94, %rd95;
	st.shared.u64 	[%r391], %rd96;
	// begin inline asm
	bfe.u32 %r215, %r248, %r217, %r250; 
	
	// end inline asm
	setp.eq.s32 	%p26, %r215, 0;
	selp.b64 	%rd97, 0, %rd21, %p26;
	shl.b32 	%r392, %r215, 8;
	add.s32 	%r393, %r392, -256;
	selp.b32 	%r394, 0, %r393, %p26;
	add.s32 	%r395, %r30, %r394;
	shl.b32 	%r396, %r395, 3;
	add.s32 	%r397, %r450, %r396;
	ld.shared.u64 	%rd98, [%r397];
	xor.b64  	%rd99, %rd97, %rd98;
	st.shared.u64 	[%r397], %rd99;
	// begin inline asm
	bfe.u32 %r219, %r248, %r221, %r250; 
	
	// end inline asm
	setp.eq.s32 	%p27, %r219, 0;
	selp.b64 	%rd100, 0, %rd21, %p27;
	shl.b32 	%r398, %r219, 8;
	add.s32 	%r399, %r398, -256;
	selp.b32 	%r400, 0, %r399, %p27;
	add.s32 	%r401, %r31, %r400;
	shl.b32 	%r402, %r401, 3;
	add.s32 	%r403, %r450, %r402;
	ld.shared.u64 	%rd101, [%r403];
	xor.b64  	%rd102, %rd100, %rd101;
	st.shared.u64 	[%r403], %rd102;
	// begin inline asm
	bfe.u32 %r223, %r248, %r225, %r250; 
	
	// end inline asm
	setp.eq.s32 	%p28, %r223, 0;
	selp.b64 	%rd103, 0, %rd21, %p28;
	shl.b32 	%r404, %r223, 8;
	add.s32 	%r405, %r404, -256;
	selp.b32 	%r406, 0, %r405, %p28;
	add.s32 	%r407, %r32, %r406;
	shl.b32 	%r408, %r407, 3;
	add.s32 	%r409, %r450, %r408;
	ld.shared.u64 	%rd104, [%r409];
	xor.b64  	%rd105, %rd103, %rd104;
	st.shared.u64 	[%r409], %rd105;
	// begin inline asm
	bfe.u32 %r227, %r248, %r229, %r250; 
	
	// end inline asm
	setp.eq.s32 	%p29, %r227, 0;
	selp.b64 	%rd106, 0, %rd21, %p29;
	shl.b32 	%r410, %r227, 8;
	add.s32 	%r411, %r410, -256;
	selp.b32 	%r412, 0, %r411, %p29;
	add.s32 	%r413, %r33, %r412;
	shl.b32 	%r414, %r413, 3;
	add.s32 	%r415, %r450, %r414;
	ld.shared.u64 	%rd107, [%r415];
	xor.b64  	%rd108, %rd106, %rd107;
	st.shared.u64 	[%r415], %rd108;
	// begin inline asm
	bfe.u32 %r231, %r248, %r233, %r250; 
	
	// end inline asm
	setp.eq.s32 	%p30, %r231, 0;
	selp.b64 	%rd109, 0, %rd21, %p30;
	shl.b32 	%r416, %r231, 8;
	add.s32 	%r417, %r416, -256;
	selp.b32 	%r418, 0, %r417, %p30;
	add.s32 	%r419, %r34, %r418;
	shl.b32 	%r420, %r419, 3;
	add.s32 	%r421, %r450, %r420;
	ld.shared.u64 	%rd110, [%r421];
	xor.b64  	%rd111, %rd109, %rd110;
	st.shared.u64 	[%r421], %rd111;
	// begin inline asm
	bfe.u32 %r235, %r248, %r237, %r250; 
	
	// end inline asm
	setp.eq.s32 	%p31, %r235, 0;
	selp.b64 	%rd112, 0, %rd21, %p31;
	shl.b32 	%r422, %r235, 8;
	add.s32 	%r423, %r422, -256;
	selp.b32 	%r424, 0, %r423, %p31;
	add.s32 	%r425, %r35, %r424;
	shl.b32 	%r426, %r425, 3;
	add.s32 	%r427, %r450, %r426;
	ld.shared.u64 	%rd113, [%r427];
	xor.b64  	%rd114, %rd112, %rd113;
	st.shared.u64 	[%r427], %rd114;
	// begin inline asm
	bfe.u32 %r239, %r248, %r241, %r250; 
	
	// end inline asm
	setp.eq.s32 	%p32, %r239, 0;
	selp.b64 	%rd115, 0, %rd21, %p32;
	shl.b32 	%r428, %r239, 8;
	add.s32 	%r429, %r428, -256;
	selp.b32 	%r430, 0, %r429, %p32;
	add.s32 	%r431, %r36, %r430;
	shl.b32 	%r432, %r431, 3;
	add.s32 	%r433, %r450, %r432;
	ld.shared.u64 	%rd116, [%r433];
	xor.b64  	%rd117, %rd115, %rd116;
	st.shared.u64 	[%r433], %rd117;
	// begin inline asm
	bfe.u32 %r243, %r248, %r245, %r250; 
	
	// end inline asm
	setp.eq.s32 	%p33, %r243, 0;
	selp.b64 	%rd118, 0, %rd21, %p33;
	shl.b32 	%r434, %r243, 8;
	add.s32 	%r435, %r434, -256;
	selp.b32 	%r436, 0, %r435, %p33;
	add.s32 	%r437, %r37, %r436;
	shl.b32 	%r438, %r437, 3;
	add.s32 	%r439, %r450, %r438;
	ld.shared.u64 	%rd119, [%r439];
	xor.b64  	%rd120, %rd118, %rd119;
	st.shared.u64 	[%r439], %rd120;
	// begin inline asm
	bfe.u32 %r247, %r248, %r249, %r250; 
	
	// end inline asm
	setp.eq.s32 	%p34, %r247, 0;
	selp.b64 	%rd121, 0, %rd21, %p34;
	shl.b32 	%r440, %r247, 8;
	add.s32 	%r441, %r440, -256;
	selp.b32 	%r442, 0, %r441, %p34;
	add.s32 	%r443, %r38, %r442;
	shl.b32 	%r444, %r443, 3;
	add.s32 	%r445, %r450, %r444;
	ld.shared.u64 	%rd122, [%r445];
	xor.b64  	%rd123, %rd121, %rd122;
	st.shared.u64 	[%r445], %rd123;
	add.s32 	%r453, %r453, %r1;
	setp.lt.u32 	%p35, %r453, %r46;
	@%p35 bra 	$L__BB3_4;

$L__BB3_5:
	setp.gt.u32 	%p36, %r2, 127;
	bar.sync 	0;
	ld.shared.u64 	%rd124, [%r4+4096];
	ld.shared.u64 	%rd125, [%r4];
	xor.b64  	%rd126, %rd125, %rd124;
	st.shared.u64 	[%r4], %rd126;
	ld.shared.u64 	%rd127, [%r4+2048];
	xor.b64  	%rd128, %rd127, %rd124;
	st.shared.u64 	[%r4+2048], %rd128;
	bar.sync 	0;
	@%p36 bra 	$L__BB3_7;

	ld.shared.u64 	%rd129, [%r4];
	ld.shared.u64 	%rd130, [%r4+1024];
	xor.b64  	%rd131, %rd129, %rd130;
	st.shared.u64 	[%r4], %rd131;
	ld.shared.u64 	%rd132, [%r4+2048];
	ld.shared.u64 	%rd133, [%r4+3072];
	xor.b64  	%rd134, %rd132, %rd133;
	st.shared.u64 	[%r4+2048], %rd134;

$L__BB3_7:
	setp.gt.u32 	%p37, %r2, 63;
	bar.sync 	0;
	@%p37 bra 	$L__BB3_9;

	ld.shared.u64 	%rd135, [%r4];
	ld.shared.u64 	%rd136, [%r4+512];
	xor.b64  	%rd137, %rd135, %rd136;
	st.shared.u64 	[%r4], %rd137;
	ld.shared.u64 	%rd138, [%r4+2048];
	ld.shared.u64 	%rd139, [%r4+2560];
	xor.b64  	%rd140, %rd138, %rd139;
	st.shared.u64 	[%r4+2048], %rd140;

$L__BB3_9:
	setp.gt.u32 	%p38, %r2, 31;
	bar.sync 	0;
	@%p38 bra 	$L__BB3_11;

	ld.shared.u64 	%rd141, [%r4];
	ld.shared.u64 	%rd142, [%r4+256];
	xor.b64  	%rd143, %rd141, %rd142;
	st.shared.u64 	[%r4], %rd143;
	ld.shared.u64 	%rd144, [%r4+2048];
	ld.shared.u64 	%rd145, [%r4+2304];
	xor.b64  	%rd146, %rd144, %rd145;
	st.shared.u64 	[%r4+2048], %rd146;

$L__BB3_11:
	setp.lt.u32 	%p39, %r2, 32;
	bar.sync 	0;
	@%p39 bra 	$L__BB3_14;
	bra.uni 	$L__BB3_12;

$L__BB3_14:
	ld.shared.u64 	%rd153, [%r4];
	shl.b64 	%rd154, %rd7, 5;
	add.s64 	%rd155, %rd3, %rd154;
	mul.wide.u32 	%rd156, %r452, 8;
	add.s64 	%rd157, %rd155, %rd156;
	atom.global.xor.b64 	%rd158, [%rd157], %rd153;
	bra.uni 	$L__BB3_15;

$L__BB3_12:
	@%p37 bra 	$L__BB3_15;

	ld.shared.u64 	%rd147, [%r4+1792];
	shl.b64 	%rd148, %rd8, 5;
	add.s64 	%rd149, %rd3, %rd148;
	mul.wide.u32 	%rd150, %r452, 8;
	add.s64 	%rd151, %rd149, %rd150;
	atom.global.xor.b64 	%rd152, [%rd151], %rd147;

$L__BB3_15:
	bar.sync 	0;
	add.s32 	%r452, %r452, 1;
	setp.lt.u32 	%p41, %r452, 4;
	@%p41 bra 	$L__BB3_2;

	cvt.u32.u64 	%r446, %rd6;
	add.s32 	%r451, %r446, 1;
	setp.lt.u32 	%p42, %r451, 4;
	@%p42 bra 	$L__BB3_1;

	ret;

}