m_0+16];
	ld.param.u64 	%rd3, [mvn_normalize_fp16_param_0+8];
	ld.param.u64 	%rd2, [mvn_normalize_fp16_param_0];
	mov.u32 	%r15, %ctaid.x;
	shl.b32 	%r16, %r15, 8;
	mov.u32 	%r17, %tid.x;
	add.s32 	%r18, %r16, %r17;
	shl.b32 	%r1, %r18, 2;
	setp.ge.u32 	%p1, %r1, %r10;
	@%p1 bra 	$L__BB0_11;

	cvta.to.global.u64 	%rd8, %rd4;
	cvta.to.global.u64 	%rd9, %rd3;
	cvta.to.global.u64 	%rd10, %rd2;
	mov.u32 	%r19, %ctaid.y;
	mad.lo.s32 	%r20, %r10, %r19, %r1;
	rem.u32 	%r3, %r1, %r12;
	mad.lo.s32 	%r21, %r12, %r19, %r3;
	cvt.u64.u32 	%rd1, %r20;
	mul.wide.u32 	%rd11, %r20, 2;
	add.s64 	%rd12, %rd10, %rd11;
	ld.global.v2.u32 	{%r22, %r23}, [%rd12];
	mul.wide.u32 	%rd13, %r21, 4;
	add.s64 	%rd14, %rd9, %rd13;
	ld.global.v4.f32 	{%f17, %f18, %f19, %f20}, [%rd14];
	add.s64 	%rd15, %rd8, %rd13;
	ld.global.v4.f32 	{%f21, %f22, %f23, %f24}, [%rd15];
	setp.eq.s16 	%p2, %rs57, 0;
	@%p2 bra 	$L__BB0_3;

	cvta.to.global.u64 	%rd16, %rd6;
	cvta.to.global.u64 	%rd17, %rd5;
	mul.wide.u32 	%rd18, %r3, 2;
	add.s64 	%rd19, %rd17, %rd18;
	ld.global.v2.u32 	{%r24, %r25}, [%rd19];
	add.s64 	%rd20, %rd16, %rd18;
	ld.global.v2.u32 	{%r28, %r29}, [%rd20];
	mov.b32 	{%rs118, %rs119}, %r24;
	mov.b32 	{%rs122, %rs123}, %r28;
	mov.b32 	{%rs116, %rs117}, %r25;
	mov.b32 	{%rs120, %rs121}, %r29;

$L__BB0_3:
	mov.b32 	{%rs64, %rs20}, %r22;
	// begin inline asm
	{  cvt.f32.f16 %f25, %rs64;}

	// end inline asm
	div.rn.f32 	%f26, %f17, %f15;
	sub.f32 	%f11, %f25, %f26;
	div.rn.f32 	%f12, %f18, %f15;
	mov.b32 	{%rs21, %rs22}, %r23;
	div.rn.f32 	%f13, %f19, %f15;
	div.rn.f32 	%f14, %f20, %f15;
	@%p2 bra 	$L__BB0_7;

	setp.eq.s16 	%p4, %rs55, 0;
	@%p4 bra 	$L__BB0_6;

	setp.ne.s16 	%p5, %rs58, 0;
	// begin inline asm
	{  cvt.f32.f16 %f27, %rs118;}

	// end inline asm
	// begin inline asm
	{  cvt.f32.f16 %f28, %rs122;}

	// end inline asm
	fma.rn.f32 	%f42, %f11, %f27, %f28;
	setp.lt.f32 	%p6, %f42, 0f00000000;
	and.pred  	%p7, %p5, %p6;
	selp.f32 	%f43, 0f00000000, %f42, %p7;
	abs.f32 	%f44, %f43;
	setp.le.f32 	%p8, %f44, 0f7F800000;
	selp.f32 	%f29, %f43, 0f00000000, %p8;
	// begin inline asm
	{  cvt.rn.f16.f32 %rs126, %f29;}

	// end inline asm
	// begin inline asm
	{  cvt.f32.f16 %f30, %rs20;}

	// end inline asm
	sub.f32 	%f45, %f30, %f12;
	// begin inline asm
	{  cvt.f32.f16 %f31, %rs119;}

	// end inline asm
	// begin inline asm
	{  cvt.f32.f16 %f32, %rs123;}

	// end inline asm
	fma.rn.f32 	%f46, %f45, %f31, %f32;
	setp.lt.f32 	%p9, %f46, 0f00000000;
	and.pred  	%p10, %p5, %p9;
	selp.f32 	%f47, 0f00000000, %f46, %p10;
	abs.f32 	%f48, %f47;
	setp.le.f32 	%p11, %f48, 0f7F800000;
	selp.f32 	%f33, %f47, 0f00000000, %p11;
	// begin inline asm
	{  cvt.rn.f16.f32 %rs127, %f33;}

	// end inline asm
	// begin inline asm
	{  cvt.f32.f16 %f34, %rs21;}

	// end inline asm
	sub.f32 	%f49, %f34, %f13;
	// begin inline asm
	{  cvt.f32.f16 %f35, %rs116;}

	// end inline asm
	// begin inline asm
	{  cvt.f32.f16 %f36, %rs120;}

	// end inline asm
	fma.rn.f32 	%f50, %f49, %f35, %f36;
	setp.lt.f32 	%p12, %f50, 0f00000000;
	and.pred  	%p13, %p5, %p12;
	selp.f32 	%f51, 0f00000000, %f50, %p13;
	abs.f32 	%f52, %f51;
	setp.le.f32 	%p14, %f52, 0f7F800000;
	selp.f32 	%f37, %f51, 0f00000000, %p14;
	// begin inline asm
	{  cvt.rn.f16.f32 %rs124, %f37;}

	// end inline asm
	// begin inline asm
	{  cvt.f32.f16 %f38, %rs22;}

	// end inline asm
	sub.f32 	%f53, %f38, %f14;
	// begin inline asm
	{  cvt.f32.f16 %f39, %rs117;}

	// end inline asm
	// begin inline asm
	{  cvt.f32.f16 %f40, %rs121;}

	// end inline asm
	fma.rn.f32 	%f54, %f53, %f39, %f40;
	setp.lt.f32 	%p15, %f54, 0f00000000;
	and.pred  	%p16, %p5, %p15;
	selp.f32 	%f55, 0f00000000, %f54, %p16;
	abs.f32 	%f56, %f55;
	setp.le.f32 	%p17, %f56, 0f7F800000;
	selp.f32 	%f41, %f55, 0f00000000, %p17;
	// begin inline asm
	{  cvt.rn.f16.f32 %rs125, %f41;}

	// end inline asm
	bra.uni 	$L__BB0_10;

$L__BB0_7:
	setp.eq.s16 	%p31, %rs55, 0;
	@%p31 bra 	$L__BB0_9;

	setp.ne.s16 	%p32, %rs58, 0;
	setp.lt.f32 	%p33, %f11, 0f00000000;
	and.pred  	%p34, %p32, %p33;
	selp.f32 	%f122, 0f00000000, %f11, %p34;
	abs.f32 	%f123, %f122;
	setp.le.f32 	%p35, %f123, 0f7F800000;
	selp.f32 	%f115, %f122, 0f00000000, %p35;
	// begin inline asm
	{  cvt.rn.f16.f32 %rs126, %f115;}

	// end inline asm
	// begin inline asm
	{  cvt.f32.f16 %f116, %rs20;}

	// end inline asm
	sub.f32 	%f124, %f116, %f12;
	setp.lt.f32 	%p36, %f124, 0f00000000;
	and.pred  	%p37, %p32, %p36;
	selp.f32 	%f125, 0f00000000, %f124, %p37;
	abs.f32 	%f126, %f125;
	setp.le.f32 	%p38, %f126, 0f7F800000;
	selp.f32 	%f117, %f125, 0f00000000, %p38;
	// begin inline asm
	{  cvt.rn.f16.f32 %rs127, %f117;}

	// end inline asm
	// begin inline asm
	{  cvt.f32.f16 %f118, %rs21;}

	// end inline asm
	sub.f32 	%f127, %f118, %f13;
	setp.lt.f32 	%p39, %f127, 0f00000000;
	and.pred  	%p40, %p32, %p39;
	selp.f32 	%f128, 0f00000000, %f127, %p40;
	abs.f32 	%f129, %f128;
	setp.le.f32 	%p41, %f129, 0f7F800000;
	selp.f32 	%f119, %f128, 0f00000000, %p41;
	// begin inline asm
	{  cvt.rn.f16.f32 %rs124, %f119;}

	// end inline asm
	// begin inline asm
	{  cvt.f32.f16 %f120, %rs22;}

	// end inline asm
	sub.f32 	%f130, %f120, %f14;
	setp.lt.f32 	%p42, %f130, 0f00000000;
	and.pred  	%p43, %p32, %p42;
	selp.f32 	%f131, 0f00000000, %f130, %p43;
	abs.f32 	%f132, %f131;
	setp.le.f32 	%p44, %f132, 0f7F800000;
	selp.f32 	%f121, %f131, 0f00000000, %p44;
	// begin inline asm
	{  cvt.rn.f16.f32 %rs125, %f121;}

	// end inline asm
	bra.uni 	$L__BB0_10;

$L__BB0_6:
	setp.ne.s16 	%p18, %rs58, 0;
	mul.f32 	%f72, %f17, %f17;
	div.rn.f32 	%f73, %f72, %f15;
	sub.f32 	%f74, %f21, %f73;
	div.rn.f32 	%f75, %f74, %f15;
	add.f32 	%f76, %f16, %f75;
	rsqrt.approx.f32 	%f77, %f76;
	mul.f32 	%f78, %f77, %f11;
	// begin inline asm
	{  cvt.f32.f16 %f57, %rs118;}

	// end inline asm
	// begin inline asm
	{  cvt.f32.f16 %f58, %rs122;}

	// end inline asm
	fma.rn.f32 	%f79, %f78, %f57, %f58;
	setp.lt.f32 	%p19, %f79, 0f00000000;
	and.pred  	%p20, %p18, %p19;
	selp.f32 	%f80, 0f00000000, %f79, %p20;
	abs.f32 	%f81, %f80;
	setp.le.f32 	%p21, %f81, 0f7F800000;
	selp.f32 	%f59, %f80, 0f00000000, %p21;
	// begin inline asm
	{  cvt.rn.f16.f32 %rs126, %f59;}

	// end inline asm
	// begin inline asm
	{  cvt.f32.f16 %f60, %rs20;}

	// end inline asm
	mul.f32 	%f82, %f18, %f18;
	div.rn.f32 	%f83, %f82, %f15;
	sub.f32 	%f84, %f22, %f83;
	div.rn.f32 	%f85, %f84, %f15;
	add.f32 	%f86, %f16, %f85;
	rsqrt.approx.f32 	%f87, %f86;
	sub.f32 	%f88, %f60, %f12;
	mul.f32 	%f89, %f87, %f88;
	// begin inline asm
	{  cvt.f32.f16 %f61, %rs119;}

	// end inline asm
	// begin inline asm
	{  cvt.f32.f16 %f62, %rs123;}

	// end inline asm
	fma.rn.f32 	%f90, %f89, %f61, %f62;
	setp.lt.f32 	%p22, %f90, 0f00000000;
	and.pred  	%p23, %p18, %p22;
	selp.f32 	%f91, 0f00000000, %f90, %p23;
	abs.f32 	%f92, %f91;
	setp.le.f32 	%p24, %f92, 0f7F800000;
	selp.f32 	%f63, %f91, 0f00000000, %p24;
	// begin inline asm
	{  cvt.rn.f16.f32 %rs127, %f63;}

	// end inline asm
	// begin inline asm
	{  cvt.f32.f16 %f64, %rs21;}

	// end inline asm
	mul.f32 	%f93, %f19, %f19;
	div.rn.f32 	%f94, %f93, %f15;
	sub.f32 	%f95, %f23, %f94;
	div.rn.f32 	%f96, %f95, %f15;
	add.f32 	%f97, %f16, %f96;
	rsqrt.approx.f32 	%f98, %f97;
	sub.f32 	%f99, %f64, %f13;
	mul.f32 	%f100, %f98, %f99;
	// begin inline asm
	{  cvt.f32.f16 %f65, %rs116;}

	// end inline asm
	// begin inline asm
	{  cvt.f32.f16 %f66, %rs120;}

	// end inline asm
	fma.rn.f32 	%f101, %f100, %f65, %f66;
	setp.lt.f32 	%p25, %f101, 0f00000000;
	and.pred  	%p26, %p18, %p25;
	selp.f32 	%f102, 0f00000000, %f101, %p26;
	abs.f32 	%f103, %f102;
	setp.le.f32 	%p27, %f103, 0f7F800000;
	selp.f32 	%f67, %f102, 0f00000000, %p27;
	// begin inline asm
	{  cvt.rn.f16.f32 %rs124, %f67;}

	// end inline asm
	// begin inline asm
	{  cvt.f32.f16 %f68, %rs22;}

	// end inline asm
	mul.f32 	%f104, %f20, %f20;
	div.rn.f32 	%f105, %f104, %f15;
	sub.f32 	%f106, %f24, %f105;
	div.rn.f32 	%f107, %f106, %f15;
	add.f32 	%f108, %f16, %f107;
	rsqrt.approx.f32 	%f109, %f108;
	sub.f32 	%f110, %f68, %f14;
	mul.f32 	%f111, %f109, %f110;
	// begin inline asm
	{  cvt.f32.f16 %f69, %rs117;}

	// end inline asm
	// begin inline asm
	{  cvt.f32.f16 %f70, %rs121;}

	// end inline asm
	fma.rn.f32 	%f112, %f111, %f69, %f70;
	setp.lt.f32 	%p28, %f112, 0f00000000;
	and.pred  	%p29, %p18, %p28;
	selp.f32 	%f113, 0f00000000, %f112, %p29;
	abs.f32 	%f114, %f113;
	setp.le.f32 	%p30, %f114, 0f7F800000;
	selp.f32 	%f71, %f113, 0f00000000, %p30;
	// begin inline asm
	{  cvt.rn.f16.f32 %rs125, %f71;}

	// end inline asm
	bra.uni 	$L__BB0_10;

$L__BB0_9:
	setp.ne.s16 	%p45, %rs58, 0;
	mul.f32 	%f140, %f17, %f17;
	div.rn.f32 	%f141, %f140, %f15;
	sub.f32 	%f142, %f21, %f141;
	div.rn.f32 	%f143, %f142, %f15;
	add.f32 	%f144, %f16, %f143;
	rsqrt.approx.f32 	%f145, %f144;
	mul.f32 	%f146, %f145, %f11;
	setp.lt.f32 	%p46, %f146, 0f00000000;
	and.pred  	%p47, %p45, %p46;
	selp.f32 	%f147, 0f00000000, %f146, %p47;
	abs.f32 	%f148, %f147;
	setp.le.f32 	%p48, %f148, 0f7F800000;
	selp.f32 	%f133, %f147, 0f00000000, %p48;
	// begin inline asm
	{  cvt.rn.f16.f32 %rs126, %f133;}

	// end inline asm
	// begin inline asm
	{  cvt.f32.f16 %f134, %rs20;}

	// end inline asm
	mul.f32 	%f149, %f18, %f18;
	div.rn.f32 	%f150, %f149, %f15;
	sub.f32 	%f151, %f22, %f150;
	div.rn.f32 	%f152, %f151, %f15;
	add.f32 	%f153, %f16, %f152;
	rsqrt.approx.f32 	%f154, %f153;
	sub.f32 	%f155, %f134, %f12;
	mul.f32 	%f156, %f154, %f155;
	setp.lt.f32 	%p49, %f156, 0f00000000;
	and.pred  	%p50, %p45, %p49;
	selp.f32 	%f157, 0f00000000, %f156, %p50;
	abs.f32 	%f158, %f157;
	setp.le.f32 	%p51, %f158, 0f7F800000;
	selp.f32 	%f135, %f157, 0f00000000, %p51;
	// begin inline asm
	{  cvt.rn.f16.f32 %rs127, %f135;}

	// end inline asm
	// begin inline asm
	{  cvt.f32.f16 %f136, %rs21;}

	// end inline asm
	mul.f32 	%f159, %f19, %f19;
	div.rn.f32 	%f160, %f159, %f15;
	sub.f32 	%f161, %f23, %f160;
	div.rn.f32 	%f162, %f161, %f15;
	add.f32 	%f163, %f16, %f162;
	rsqrt.approx.f32 	%f164, %f163;
	sub.f32 	%f165, %f136, %f13;
	mul.f32 	%f166, %f164, %f165;
	setp.lt.f32 	%p52, %f166, 0f00000000;
	and.pred  	%p53, %p45, %p52;
	selp.f32 	%f167, 0f00000000, %f166, %p53;
	abs.f32 	%f168, %f167;
	setp.le.f32 	%p54, %f168, 0f7F800000;
	selp.f32 	%f137, %f167, 0f00000000, %p54;
	// begin inline asm
	{  cvt.rn.f16.f32 %rs124, %f137;}

	// end inline asm
	// begin inline asm
	{  cvt.f32.f16 %f138, %rs22;}

	// end inline asm
	mul.f32 	%f169, %f20, %f20;
	div.rn.f32 	%f170, %f169, %f15;
	sub.f32 	%f171, %f24, %f170;
	div.rn.f32 	%f172, %f171, %f15;
	add.f32 	%f173, %f16, %f172;
	rsqrt.approx.f32 	%f174, %f173;
	sub.f32 	%f175, %f138, %f14;
	mul.f32 	%f176, %f174, %f175;
	setp.lt.f32 	%p55, %f176, 0f00000000;
	and.pred  	%p56, %p45, %p55;
	selp.f32 	%f177, 0f00000000, %f176, %p56;
	abs.f32 	%f178, %f177;
	setp.le.f32 	%p57, %f178, 0f7F800000;
	selp.f32 	%f139, %f177, 0f00000000, %p57;
	// begin inline asm
	{  cvt.rn.f16.f32 %rs125, %f139;}

	// end inline asm

$L__BB0_10:
	cvta.to.global.u64 	%rd21, %rd7;
	shl.b64 	%rd22, %rd1, 1;
	add.s64 	%rd23, %rd21, %rd22;
	st.global.v4.u16 	[%rd23], {%rs126, %rs127, %rs124, %rs125};

$L__BB0_11:
	ret;

}