//
// Generated by NVIDIA NVVM Compiler
// Compiler built on Sat Mar 15 02:52:22 2014 (1394848342)
// Cuda compilation tools, release 6.0, V6.0.1
//

.version 4.0
.target sm_20
.address_size 32


.visible .entry BenchMarkDRAMKernel(
	.param .u32 BenchMarkDRAMKernel_param_0
)
{
	.reg .pred 	%p<2>;
	.reg .s32 	%r<9>;
	.reg .f32 	%f<18>;


	ld.param.u32 	%r2, [BenchMarkDRAMKernel_param_0];
	cvta.to.global.u32 	%r1, %r2;
	mov.u32 	%r3, %ctaid.x;
	mov.u32 	%r4, %ntid.x;
	mov.u32 	%r5, %tid.x;
	mad.lo.s32 	%r6, %r3, %r4, %r5;
	shl.b32 	%r7, %r6, 4;
	add.s32 	%r8, %r1, %r7;
	ld.global.v4.f32 	{%f5, %f6, %f7, %f8}, [%r8];
	add.f32 	%f1, %f5, 0f3F800000;
	add.f32 	%f2, %f6, 0f3F800000;
	add.f32 	%f3, %f7, 0f3F800000;
	add.f32 	%f4, %f8, 0f3F800000;
	mul.f32 	%f13, %f2, %f2;
	fma.rn.f32 	%f14, %f1, %f1, %f13;
	fma.rn.f32 	%f15, %f3, %f3, %f14;
	fma.rn.f32 	%f16, %f4, %f4, %f15;
	sqrt.rn.f32 	%f17, %f16;
	setp.neu.f32	%p1, %f17, 0fC6410800;
	@%p1 bra 	BB0_2;

	st.global.v4.f32 	[%r1], {%f1, %f2, %f3, %f4};

BB0_2:
	ret;
}