map.ptx

//
// Generated by LLVM NVPTX Back-End
//

.version 3.1
.target sm_20
.address_size 64

	// .globl	map
                                        // @map
.visible .entry map(
	.param .u64 map_param_0,
	.param .u64 map_param_1,
	.param .u64 map_param_2,
	.param .u64 map_param_3
)
{
	.reg .pred 	%p<3>;
	.reg .s32 	%r<9>;
	.reg .s64 	%rl<8>;

// BB#0:                                // %entry
	ld.param.u32 	%r1, [map_param_1];
	mov.u32 	%r8, %tid.x;
	setp.ge.s32	%p1, %r8, %r1;
	@%p1 bra 	BB0_3;
// BB#1:
	ld.param.u64 	%rl1, [map_param_0];
	mov.u32 	%r6, %nctaid.x;
	ld.param.u64 	%rl2, [map_param_2];
	mov.u32 	%r7, %ntid.x;
	mul.lo.s32 	%r2, %r7, %r6;
BB0_2:                                  // %loop1.top
                                        // =>This Inner Loop Header: Depth=1
	mul.wide.s32 	%rl3, %r8, 8;
	add.s64 	%rl4, %rl2, %rl3;
	ld.u64 	%rl5, [%rl4];
	add.s64 	%rl6, %rl5, 1;
	add.s64 	%rl7, %rl1, %rl3;
	st.u64 	[%rl7], %rl6;
	add.s32 	%r8, %r8, %r2;
	setp.lt.s32	%p2, %r8, %r1;
	@%p2 bra 	BB0_2;
BB0_3:                                  // %loop1.exit
	ret;
}