//
// Generated by NVIDIA NVVM Compiler
//
// Compiler Build ID: CL-21124049
// Cuda compilation tools, release 8.0, V8.0.44
// Based on LLVM 3.4svn
//

.version 5.0
.target sm_20
.address_size 64

	// .globl	_Z1KPii

.visible .entry _Z1KPii(
	.param .u64 _Z1KPii_param_0,
	.param .u32 _Z1KPii_param_1
)
{
	.reg .pred 	%p<6>;
	.reg .b32 	%r<27>;
	.reg .b64 	%rd<9>;


	ld.param.u64 	%rd2, [_Z1KPii_param_0];
	ld.param.u32 	%r14, [_Z1KPii_param_1];
	cvta.to.global.u64 	%rd1, %rd2;
	setp.eq.s32	%p1, %r14, 0;
	@%p1 bra 	BB0_7;

	shr.u32 	%r16, %r14, 31;
	add.s32 	%r17, %r14, %r16;
	shr.s32 	%r1, %r17, 1;
	and.b32  	%r2, %r17, -2;
	mov.u32 	%r15, 0;
	setp.lt.s32	%p2, %r2, 1;
	mov.u32 	%r26, %r15;
	@%p2 bra 	BB0_5;

	shl.b32 	%r19, %r1, 1;
	neg.s32 	%r21, %r19;
	mov.u32 	%r22, 0;

BB0_3:
	.pragma "nounroll";
	mul.wide.u32 	%rd3, %r22, 4;
	add.s64 	%rd4, %rd1, %rd3;
	add.s32 	%r20, %r22, 1;
	st.global.u32 	[%rd4], %r20;
	mul.wide.u32 	%rd5, %r20, 4;
	add.s64 	%rd6, %rd1, %rd5;
	add.s32 	%r22, %r22, 2;
	st.global.u32 	[%rd6], %r22;
	add.s32 	%r21, %r21, 2;
	setp.ne.s32	%p3, %r21, 0;
	@%p3 bra 	BB0_3;

	setp.eq.s32	%p4, %r2, %r14;
	mov.u32 	%r26, %r2;
	@%p4 bra 	BB0_7;

BB0_5:
	mov.u32 	%r25, %r26;
	sub.s32 	%r23, %r25, %r14;

BB0_6:
	.pragma "nounroll";
	mul.wide.u32 	%rd7, %r25, 4;
	add.s64 	%rd8, %rd1, %rd7;
	add.s32 	%r25, %r25, 1;
	st.global.u32 	[%rd8], %r25;
	add.s32 	%r23, %r23, 1;
	setp.ne.s32	%p5, %r23, 0;
	@%p5 bra 	BB0_6;

BB0_7:
	ret;
}