//
// Generated by NVIDIA NVVM Compiler
//
// Compiler Build ID: CL-19324607
// Cuda compilation tools, release 7.0, V7.0.27
// Based on LLVM 3.4svn
//

.version 4.2
.target sm_20
.address_size 64

	// .globl	amp
.const .align 4 .b8 c_bfs[4096];

.visible .entry amp(
	.param .u64 amp_param_0,
	.param .u64 amp_param_1,
	.param .u64 amp_param_2,
	.param .u64 amp_param_3,
	.param .u64 amp_param_4,
	.param .u32 amp_param_5,
	.param .u32 amp_param_6
)
.maxntid 256, 1, 1
.minnctapersm 1
{
	.reg .pred 	%p<2>;
	.reg .s32 	%r<25>;
	.reg .s64 	%rd<8>;


	ld.param.u64 	%rd1, [amp_param_0];
	ld.param.u64 	%rd2, [amp_param_1];
	ld.param.u32 	%r2, [amp_param_6];
	mov.u32 	%r3, %ntid.x;
	mov.u32 	%r4, %ctaid.x;
	mov.u32 	%r5, %tid.x;
	mad.lo.s32 	%r1, %r3, %r4, %r5;
	setp.ge.u32	%p1, %r1, %r2;
	@%p1 bra 	BB0_2;

	cvta.to.global.u64 	%rd3, %rd2;
	cvta.to.global.u64 	%rd4, %rd1;
	mul.wide.u32 	%rd5, %r1, 272;
	add.s64 	%rd6, %rd4, %rd5;
	ld.global.u32 	%r6, [%rd6+256];
	ld.global.u32 	%r7, [%rd6+4];
	ld.global.u32 	%r8, [%rd6+8];
	ld.global.u32 	%r9, [%rd6+12];
	ld.global.u32 	%r10, [%rd6+16];
	ld.global.u32 	%r11, [%rd6+20];
	ld.global.u32 	%r12, [%rd6+24];
	ld.global.u32 	%r13, [%rd6+28];
	ld.global.u32 	%r14, [%rd6+32];
	ld.global.u32 	%r15, [%rd6+36];
	ld.global.u32 	%r16, [%rd6+40];
	ld.global.u32 	%r17, [%rd6+44];
	ld.global.u32 	%r18, [%rd6+48];
	ld.global.u32 	%r19, [%rd6+52];
	ld.global.u32 	%r20, [%rd6+56];
	ld.global.u32 	%r21, [%rd6+60];
	ld.const.u32 	%r22, [c_bfs];
	ld.global.u32 	%r23, [%rd6];
	or.b32  	%r24, %r22, %r23;
	add.s64 	%rd7, %rd3, %rd5;
	st.global.u32 	[%rd7], %r24;
	st.global.u32 	[%rd7+4], %r7;
	st.global.u32 	[%rd7+8], %r8;
	st.global.u32 	[%rd7+12], %r9;
	st.global.u32 	[%rd7+16], %r10;
	st.global.u32 	[%rd7+20], %r11;
	st.global.u32 	[%rd7+24], %r12;
	st.global.u32 	[%rd7+28], %r13;
	st.global.u32 	[%rd7+32], %r14;
	st.global.u32 	[%rd7+36], %r15;
	st.global.u32 	[%rd7+40], %r16;
	st.global.u32 	[%rd7+44], %r17;
	st.global.u32 	[%rd7+48], %r18;
	st.global.u32 	[%rd7+52], %r19;
	st.global.u32 	[%rd7+56], %r20;
	st.global.u32 	[%rd7+60], %r21;
	st.global.u32 	[%rd7+256], %r6;

BB0_2:
	ret;
}


