;
	.reg .b16 	%rs<9>;
	.reg .f32 	%f<65>;
	.reg .b32 	%r<39>;
	.reg .b64 	%rd<26>;


	ld.param.v2.u32 	{%r20, %r21}, [mvn_compute_sums_fp16_param_0+48];
	ld.param.u64 	%rd2, [mvn_compute_sums_fp16_param_0+8];
	ld.param.u64 	%rd3, [mvn_compute_sums_fp16_param_0+16];
	ld.param.u32 	%r5, [mvn_compute_sums_fp16_param_0+64];
	ld.param.u64 	%rd4, [mvn_compute_sums_fp16_param_0];
	cvta.to.global.u64 	%rd1, %rd4;
	mov.u32 	%r22, %ctaid.x;
	shl.b32 	%r23, %r22, 8;
	mov.u32 	%r24, %tid.x;
	add.s32 	%r25, %r23, %r24;
	div.u32 	%r26, %r25, %r21;
	mul.lo.s32 	%r27, %r26, %r21;
	sub.s32 	%r2, %r25, %r27;
	shl.b32 	%r38, %r26, 6;
	mov.u32 	%r4, %ctaid.y;
	mul.lo.s32 	%r6, %r5, %r4;
	mov.f32 	%f50, 0f00000000;
	mov.u32 	%r37, 0;
	mov.f32 	%f49, %f50;

$L__BB0_1:
	mad.lo.s32 	%r9, %r38, %r21, %r2;
	setp.ge.u32 	%p1, %r9, %r5;
	@%p1 bra 	$L__BB0_3;

	add.s32 	%r28, %r9, %r6;
	mul.wide.u32 	%rd5, %r28, 2;
	add.s64 	%rd6, %rd1, %rd5;
	ld.global.u16 	%rs1, [%rd6];
	// begin inline asm
	{  cvt.f32.f16 %f37, %rs1;}

	// end inline asm
	add.f32 	%f49, %f49, %f37;
	fma.rn.f32 	%f50, %f37, %f37, %f50;

$L__BB0_3:
	add.s32 	%r10, %r9, %r21;
	setp.ge.u32 	%p2, %r10, %r5;
	@%p2 bra 	$L__BB0_5;

	add.s32 	%r29, %r10, %r6;
	mul.wide.u32 	%rd7, %r29, 2;
	add.s64 	%rd8, %rd1, %rd7;
	ld.global.u16 	%rs2, [%rd8];
	// begin inline asm
	{  cvt.f32.f16 %f38, %rs2;}

	// end inline asm
	add.f32 	%f49, %f49, %f38;
	fma.rn.f32 	%f50, %f38, %f38, %f50;

$L__BB0_5:
	add.s32 	%r11, %r10, %r21;
	setp.ge.u32 	%p3, %r11, %r5;
	@%p3 bra 	$L__BB0_7;

	add.s32 	%r30, %r11, %r6;
	mul.wide.u32 	%rd9, %r30, 2;
	add.s64 	%rd10, %rd1, %rd9;
	ld.global.u16 	%rs3, [%rd10];
	// begin inline asm
	{  cvt.f32.f16 %f39, %rs3;}

	// end inline asm
	add.f32 	%f49, %f49, %f39;
	fma.rn.f32 	%f50, %f39, %f39, %f50;

$L__BB0_7:
	add.s32 	%r12, %r11, %r21;
	setp.ge.u32 	%p4, %r12, %r5;
	@%p4 bra 	$L__BB0_9;

	add.s32 	%r31, %r12, %r6;
	mul.wide.u32 	%rd11, %r31, 2;
	add.s64 	%rd12, %rd1, %rd11;
	ld.global.u16 	%rs4, [%rd12];
	// begin inline asm
	{  cvt.f32.f16 %f40, %rs4;}

	// end inline asm
	add.f32 	%f49, %f49, %f40;
	fma.rn.f32 	%f50, %f40, %f40, %f50;

$L__BB0_9:
	add.s32 	%r13, %r12, %r21;
	setp.ge.u32 	%p5, %r13, %r5;
	@%p5 bra 	$L__BB0_11;

	add.s32 	%r32, %r13, %r6;
	mul.wide.u32 	%rd13, %r32, 2;
	add.s64 	%rd14, %rd1, %rd13;
	ld.global.u16 	%rs5, [%rd14];
	// begin inline asm
	{  cvt.f32.f16 %f41, %rs5;}

	// end inline asm
	add.f32 	%f49, %f49, %f41;
	fma.rn.f32 	%f50, %f41, %f41, %f50;

$L__BB0_11:
	add.s32 	%r14, %r13, %r21;
	setp.ge.u32 	%p6, %r14, %r5;
	@%p6 bra 	$L__BB0_13;

	add.s32 	%r33, %r14, %r6;
	mul.wide.u32 	%rd15, %r33, 2;
	add.s64 	%rd16, %rd1, %rd15;
	ld.global.u16 	%rs6, [%rd16];
	// begin inline asm
	{  cvt.f32.f16 %f42, %rs6;}

	// end inline asm
	add.f32 	%f49, %f49, %f42;
	fma.rn.f32 	%f50, %f42, %f42, %f50;

$L__BB0_13:
	add.s32 	%r15, %r14, %r21;
	setp.ge.u32 	%p7, %r15, %r5;
	@%p7 bra 	$L__BB0_15;

	add.s32 	%r34, %r15, %r6;
	mul.wide.u32 	%rd17, %r34, 2;
	add.s64 	%rd18, %rd1, %rd17;
	ld.global.u16 	%rs7, [%rd18];
	// begin inline asm
	{  cvt.f32.f16 %f43, %rs7;}

	// end inline asm
	add.f32 	%f49, %f49, %f43;
	fma.rn.f32 	%f50, %f43, %f43, %f50;

$L__BB0_15:
	add.s32 	%r16, %r15, %r21;
	setp.ge.u32 	%p8, %r16, %r5;
	@%p8 bra 	$L__BB0_17;

	add.s32 	%r35, %r16, %r6;
	mul.wide.u32 	%rd19, %r35, 2;
	add.s64 	%rd20, %rd1, %rd19;
	ld.global.u16 	%rs8, [%rd20];
	// begin inline asm
	{  cvt.f32.f16 %f44, %rs8;}

	// end inline asm
	add.f32 	%f49, %f49, %f44;
	fma.rn.f32 	%f50, %f44, %f44, %f50;

$L__BB0_17:
	add.s32 	%r38, %r38, 8;
	add.s32 	%r37, %r37, 8;
	setp.ne.s32 	%p9, %r37, 64;
	@%p9 bra 	$L__BB0_1;

	cvta.to.global.u64 	%rd21, %rd2;
	mad.lo.s32 	%r36, %r21, %r4, %r2;
	mul.wide.u32 	%rd22, %r36, 4;
	add.s64 	%rd23, %rd21, %rd22;
	atom.global.add.f32 	%f45, [%rd23], %f49;
	cvta.to.global.u64 	%rd24, %rd3;
	add.s64 	%rd25, %rd24, %rd22;
	atom.global.add.f32 	%f46, [%rd25], %f50;
	ret;