//
// Generated by NVIDIA NVVM Compiler
//
// Compiler Build ID: CL-31968024
// Cuda compilation tools, release 12.0, V12.0.76
// Based on NVVM 7.0.1
//

.version 8.0
.target sm_52
.address_size 64

	// .globl	oxMain
.const .align 16 .b8 params[1184];

.visible .entry oxMain()
{
	.reg .b16 	%rs<93>;
	.reg .f32 	%f<120>;
	.reg .b32 	%r<45>;
	.reg .b64 	%rd<73>;


	mov.u32 	%r1, %ctaid.x;
	mov.u32 	%r2, %ntid.x;
	mov.u32 	%r3, %tid.x;
	mad.lo.s32 	%r4, %r1, %r2, %r3;
	mov.u32 	%r5, %ctaid.y;
	mov.u32 	%r6, %ntid.y;
	mov.u32 	%r7, %tid.y;
	mad.lo.s32 	%r8, %r5, %r6, %r7;
	ld.const.u64 	%rd1, [params+784];
	cvta.to.global.u64 	%rd2, %rd1;
	ld.const.u32 	%r9, [params+776];
	mad.lo.s32 	%r10, %r9, %r8, %r4;
	mul.wide.u32 	%rd3, %r10, 8;
	add.s64 	%rd4, %rd2, %rd3;
	ld.global.u16 	%rs1, [%rd4];
	// begin inline asm
	{  cvt.f32.f16 %f1, %rs1;}

	// end inline asm
	ld.global.u16 	%rs2, [%rd4+2];
	// begin inline asm
	{  cvt.f32.f16 %f2, %rs2;}

	// end inline asm
	ld.global.u16 	%rs3, [%rd4+4];
	// begin inline asm
	{  cvt.f32.f16 %f3, %rs3;}

	// end inline asm
	ld.const.u64 	%rd5, [params+800];
	cvta.to.global.u64 	%rd6, %rd5;
	ld.const.u32 	%r11, [params+792];
	mad.lo.s32 	%r12, %r11, %r8, %r4;
	mul.wide.u32 	%rd7, %r12, 8;
	add.s64 	%rd8, %rd6, %rd7;
	ld.global.u16 	%rs4, [%rd8];
	// begin inline asm
	{  cvt.f32.f16 %f4, %rs4;}

	// end inline asm
	ld.global.u16 	%rs5, [%rd8+2];
	// begin inline asm
	{  cvt.f32.f16 %f5, %rs5;}

	// end inline asm
	ld.global.u16 	%rs6, [%rd8+4];
	// begin inline asm
	{  cvt.f32.f16 %f6, %rs6;}

	// end inline asm
	ld.const.u64 	%rd9, [params+816];
	cvta.to.global.u64 	%rd10, %rd9;
	ld.const.u32 	%r13, [params+808];
	mad.lo.s32 	%r14, %r13, %r8, %r4;
	mul.wide.u32 	%rd11, %r14, 8;
	add.s64 	%rd12, %rd10, %rd11;
	ld.global.u16 	%rs7, [%rd12];
	// begin inline asm
	{  cvt.f32.f16 %f7, %rs7;}

	// end inline asm
	ld.global.u16 	%rs8, [%rd12+2];
	// begin inline asm
	{  cvt.f32.f16 %f8, %rs8;}

	// end inline asm
	ld.global.u16 	%rs9, [%rd12+4];
	// begin inline asm
	{  cvt.f32.f16 %f9, %rs9;}

	// end inline asm
	ld.const.u64 	%rd13, [params+832];
	cvta.to.global.u64 	%rd14, %rd13;
	ld.const.u32 	%r15, [params+824];
	mad.lo.s32 	%r16, %r15, %r8, %r4;
	mul.wide.u32 	%rd15, %r16, 8;
	add.s64 	%rd16, %rd14, %rd15;
	ld.global.u16 	%rs10, [%rd16];
	// begin inline asm
	{  cvt.f32.f16 %f10, %rs10;}

	// end inline asm
	ld.global.u16 	%rs11, [%rd16+2];
	// begin inline asm
	{  cvt.f32.f16 %f11, %rs11;}

	// end inline asm
	ld.global.u16 	%rs12, [%rd16+4];
	// begin inline asm
	{  cvt.f32.f16 %f12, %rs12;}

	// end inline asm
	ld.const.u64 	%rd17, [params+848];
	cvta.to.global.u64 	%rd18, %rd17;
	ld.const.u32 	%r17, [params+840];
	mad.lo.s32 	%r18, %r17, %r8, %r4;
	mul.wide.u32 	%rd19, %r18, 8;
	add.s64 	%rd20, %rd18, %rd19;
	ld.global.u16 	%rs13, [%rd20];
	// begin inline asm
	{  cvt.f32.f16 %f13, %rs13;}

	// end inline asm
	ld.global.u16 	%rs14, [%rd20+2];
	// begin inline asm
	{  cvt.f32.f16 %f14, %rs14;}

	// end inline asm
	ld.global.u16 	%rs15, [%rd20+4];
	// begin inline asm
	{  cvt.f32.f16 %f15, %rs15;}

	// end inline asm
	ld.const.u64 	%rd21, [params+864];
	cvta.to.global.u64 	%rd22, %rd21;
	ld.const.u32 	%r19, [params+856];
	mad.lo.s32 	%r20, %r19, %r8, %r4;
	mul.wide.u32 	%rd23, %r20, 8;
	add.s64 	%rd24, %rd22, %rd23;
	ld.global.u16 	%rs16, [%rd24];
	// begin inline asm
	{  cvt.f32.f16 %f16, %rs16;}

	// end inline asm
	ld.global.u16 	%rs17, [%rd24+2];
	// begin inline asm
	{  cvt.f32.f16 %f17, %rs17;}

	// end inline asm
	ld.global.u16 	%rs18, [%rd24+4];
	// begin inline asm
	{  cvt.f32.f16 %f18, %rs18;}

	// end inline asm
	ld.const.u64 	%rd25, [params+880];
	cvta.to.global.u64 	%rd26, %rd25;
	ld.const.u32 	%r21, [params+872];
	mad.lo.s32 	%r22, %r21, %r8, %r4;
	mul.wide.u32 	%rd27, %r22, 8;
	add.s64 	%rd28, %rd26, %rd27;
	ld.global.u16 	%rs19, [%rd28];
	// begin inline asm
	{  cvt.f32.f16 %f19, %rs19;}

	// end inline asm
	ld.global.u16 	%rs20, [%rd28+2];
	// begin inline asm
	{  cvt.f32.f16 %f20, %rs20;}

	// end inline asm
	ld.global.u16 	%rs21, [%rd28+4];
	// begin inline asm
	{  cvt.f32.f16 %f21, %rs21;}

	// end inline asm
	ld.const.u64 	%rd29, [params+896];
	cvta.to.global.u64 	%rd30, %rd29;
	ld.const.u32 	%r23, [params+888];
	mad.lo.s32 	%r24, %r23, %r8, %r4;
	mul.wide.u32 	%rd31, %r24, 8;
	add.s64 	%rd32, %rd30, %rd31;
	ld.global.u16 	%rs22, [%rd32];
	// begin inline asm
	{  cvt.f32.f16 %f22, %rs22;}

	// end inline asm
	ld.global.u16 	%rs23, [%rd32+2];
	// begin inline asm
	{  cvt.f32.f16 %f23, %rs23;}

	// end inline asm
	ld.global.u16 	%rs24, [%rd32+4];
	// begin inline asm
	{  cvt.f32.f16 %f24, %rs24;}

	// end inline asm
	ld.const.u64 	%rd33, [params+912];
	cvta.to.global.u64 	%rd34, %rd33;
	ld.const.u32 	%r25, [params+904];
	mad.lo.s32 	%r26, %r25, %r8, %r4;
	mul.wide.u32 	%rd35, %r26, 8;
	add.s64 	%rd36, %rd34, %rd35;
	ld.global.u16 	%rs25, [%rd36];
	// begin inline asm
	{  cvt.f32.f16 %f25, %rs25;}

	// end inline asm
	ld.global.u16 	%rs26, [%rd36+2];
	// begin inline asm
	{  cvt.f32.f16 %f26, %rs26;}

	// end inline asm
	ld.global.u16 	%rs27, [%rd36+4];
	// begin inline asm
	{  cvt.f32.f16 %f27, %rs27;}

	// end inline asm
	ld.const.u64 	%rd37, [params+928];
	cvta.to.global.u64 	%rd38, %rd37;
	ld.const.u32 	%r27, [params+920];
	mad.lo.s32 	%r28, %r27, %r8, %r4;
	mul.wide.u32 	%rd39, %r28, 8;
	add.s64 	%rd40, %rd38, %rd39;
	ld.global.u16 	%rs28, [%rd40];
	// begin inline asm
	{  cvt.f32.f16 %f28, %rs28;}

	// end inline asm
	ld.global.u16 	%rs29, [%rd40+2];
	// begin inline asm
	{  cvt.f32.f16 %f29, %rs29;}

	// end inline asm
	ld.global.u16 	%rs30, [%rd40+4];
	// begin inline asm
	{  cvt.f32.f16 %f30, %rs30;}

	// end inline asm
	ld.const.u64 	%rd41, [params+944];
	cvta.to.global.u64 	%rd42, %rd41;
	ld.const.u32 	%r29, [params+936];
	mad.lo.s32 	%r30, %r29, %r8, %r4;
	mul.wide.u32 	%rd43, %r30, 8;
	add.s64 	%rd44, %rd42, %rd43;
	ld.global.u16 	%rs31, [%rd44];
	// begin inline asm
	{  cvt.f32.f16 %f31, %rs31;}

	// end inline asm
	ld.global.u16 	%rs32, [%rd44+2];
	// begin inline asm
	{  cvt.f32.f16 %f32, %rs32;}

	// end inline asm
	ld.global.u16 	%rs33, [%rd44+4];
	// begin inline asm
	{  cvt.f32.f16 %f33, %rs33;}

	// end inline asm
	ld.const.u64 	%rd45, [params+960];
	cvta.to.global.u64 	%rd46, %rd45;
	ld.const.u32 	%r31, [params+952];
	mad.lo.s32 	%r32, %r31, %r8, %r4;
	mul.wide.u32 	%rd47, %r32, 8;
	add.s64 	%rd48, %rd46, %rd47;
	ld.global.u16 	%rs34, [%rd48];
	// begin inline asm
	{  cvt.f32.f16 %f34, %rs34;}

	// end inline asm
	ld.global.u16 	%rs35, [%rd48+2];
	// begin inline asm
	{  cvt.f32.f16 %f35, %rs35;}

	// end inline asm
	ld.global.u16 	%rs36, [%rd48+4];
	// begin inline asm
	{  cvt.f32.f16 %f36, %rs36;}

	// end inline asm
	ld.const.u64 	%rd49, [params+976];
	cvta.to.global.u64 	%rd50, %rd49;
	ld.const.u32 	%r33, [params+968];
	mad.lo.s32 	%r34, %r33, %r8, %r4;
	mul.wide.u32 	%rd51, %r34, 8;
	add.s64 	%rd52, %rd50, %rd51;
	ld.global.u16 	%rs37, [%rd52];
	// begin inline asm
	{  cvt.f32.f16 %f37, %rs37;}

	// end inline asm
	ld.global.u16 	%rs38, [%rd52+2];
	// begin inline asm
	{  cvt.f32.f16 %f38, %rs38;}

	// end inline asm
	ld.global.u16 	%rs39, [%rd52+4];
	// begin inline asm
	{  cvt.f32.f16 %f39, %rs39;}

	// end inline asm
	ld.const.u64 	%rd53, [params+992];
	cvta.to.global.u64 	%rd54, %rd53;
	ld.const.u32 	%r35, [params+984];
	mad.lo.s32 	%r36, %r35, %r8, %r4;
	mul.wide.u32 	%rd55, %r36, 8;
	add.s64 	%rd56, %rd54, %rd55;
	ld.global.u16 	%rs40, [%rd56];
	// begin inline asm
	{  cvt.f32.f16 %f40, %rs40;}

	// end inline asm
	ld.global.u16 	%rs41, [%rd56+2];
	// begin inline asm
	{  cvt.f32.f16 %f41, %rs41;}

	// end inline asm
	ld.global.u16 	%rs42, [%rd56+4];
	// begin inline asm
	{  cvt.f32.f16 %f42, %rs42;}

	// end inline asm
	ld.const.u64 	%rd57, [params+1008];
	cvta.to.global.u64 	%rd58, %rd57;
	ld.const.u32 	%r37, [params+1000];
	mad.lo.s32 	%r38, %r37, %r8, %r4;
	mul.wide.u32 	%rd59, %r38, 8;
	add.s64 	%rd60, %rd58, %rd59;
	ld.global.u16 	%rs43, [%rd60];
	// begin inline asm
	{  cvt.f32.f16 %f43, %rs43;}

	// end inline asm
	ld.global.u16 	%rs44, [%rd60+2];
	// begin inline asm
	{  cvt.f32.f16 %f44, %rs44;}

	// end inline asm
	ld.global.u16 	%rs45, [%rd60+4];
	// begin inline asm
	{  cvt.f32.f16 %f45, %rs45;}

	// end inline asm
	ld.const.u64 	%rd61, [params+1024];
	cvta.to.global.u64 	%rd62, %rd61;
	ld.const.u32 	%r39, [params+1016];
	mad.lo.s32 	%r40, %r39, %r8, %r4;
	mul.wide.u32 	%rd63, %r40, 8;
	add.s64 	%rd64, %rd62, %rd63;
	ld.global.u16 	%rs46, [%rd64];
	// begin inline asm
	{  cvt.f32.f16 %f46, %rs46;}

	// end inline asm
	ld.global.u16 	%rs47, [%rd64+2];
	// begin inline asm
	{  cvt.f32.f16 %f47, %rs47;}

	// end inline asm
	ld.global.u16 	%rs48, [%rd64+4];
	// begin inline asm
	{  cvt.f32.f16 %f48, %rs48;}

	// end inline asm
	ld.const.u64 	%rd65, [params+1040];
	cvta.to.global.u64 	%rd66, %rd65;
	ld.const.u32 	%r41, [params+1032];
	mad.lo.s32 	%r42, %r41, %r8, %r4;
	mul.wide.u32 	%rd67, %r42, 8;
	add.s64 	%rd68, %rd66, %rd67;
	ld.global.u16 	%rs49, [%rd68];
	// begin inline asm
	{  cvt.f32.f16 %f49, %rs49;}

	// end inline asm
	ld.global.u16 	%rs50, [%rd68+2];
	// begin inline asm
	{  cvt.f32.f16 %f50, %rs50;}

	// end inline asm
	ld.global.u16 	%rs51, [%rd68+4];
	// begin inline asm
	{  cvt.f32.f16 %f51, %rs51;}

	// end inline asm
	ld.const.u64 	%rd69, [params+1056];
	cvta.to.global.u64 	%rd70, %rd69;
	ld.const.u32 	%r43, [params+1048];
	mad.lo.s32 	%r44, %r43, %r8, %r4;
	mul.wide.u32 	%rd71, %r44, 8;
	add.s64 	%rd72, %rd70, %rd71;
	ld.global.u16 	%rs52, [%rd72];
	// begin inline asm
	{  cvt.f32.f16 %f52, %rs52;}

	// end inline asm
	ld.global.u16 	%rs53, [%rd72+2];
	// begin inline asm
	{  cvt.f32.f16 %f53, %rs53;}

	// end inline asm
	ld.global.u16 	%rs54, [%rd72+4];
	// begin inline asm
	{  cvt.f32.f16 %f54, %rs54;}

	// end inline asm
	ld.global.u16 	%rs55, [%rd4+6];
	// begin inline asm
	{  cvt.f32.f16 %f55, %rs55;}

	// end inline asm
	ld.global.u16 	%rs56, [%rd40+6];
	// begin inline asm
	{  cvt.f32.f16 %f56, %rs56;}

	// end inline asm
	min.f32 	%f92, %f55, %f56;
	add.f32 	%f93, %f1, %f28;
	add.f32 	%f94, %f2, %f29;
	add.f32 	%f95, %f3, %f30;
	add.f32 	%f96, %f4, %f31;
	add.f32 	%f97, %f5, %f32;
	add.f32 	%f98, %f6, %f33;
	add.f32 	%f99, %f7, %f34;
	add.f32 	%f100, %f8, %f35;
	add.f32 	%f101, %f9, %f36;
	add.f32 	%f102, %f10, %f37;
	add.f32 	%f103, %f11, %f38;
	add.f32 	%f104, %f12, %f39;
	add.f32 	%f105, %f13, %f40;
	add.f32 	%f106, %f14, %f41;
	add.f32 	%f107, %f15, %f42;
	add.f32 	%f108, %f16, %f43;
	add.f32 	%f109, %f17, %f44;
	add.f32 	%f110, %f18, %f45;
	add.f32 	%f111, %f19, %f46;
	add.f32 	%f112, %f20, %f47;
	add.f32 	%f113, %f21, %f48;
	add.f32 	%f114, %f22, %f49;
	add.f32 	%f115, %f23, %f50;
	add.f32 	%f116, %f24, %f51;
	add.f32 	%f117, %f25, %f52;
	add.f32 	%f118, %f26, %f53;
	add.f32 	%f119, %f27, %f54;
	mul.f32 	%f57, %f93, %f92;
	mul.f32 	%f58, %f94, %f92;
	mul.f32 	%f59, %f95, %f92;
	mul.f32 	%f61, %f96, %f92;
	mul.f32 	%f62, %f97, %f92;
	mul.f32 	%f63, %f98, %f92;
	mul.f32 	%f65, %f99, %f92;
	mul.f32 	%f66, %f100, %f92;
	mul.f32 	%f67, %f101, %f92;
	mul.f32 	%f69, %f102, %f92;
	mul.f32 	%f70, %f103, %f92;
	mul.f32 	%f71, %f104, %f92;
	mul.f32 	%f73, %f105, %f92;
	mul.f32 	%f74, %f106, %f92;
	mul.f32 	%f75, %f107, %f92;
	mul.f32 	%f77, %f108, %f92;
	mul.f32 	%f78, %f109, %f92;
	mul.f32 	%f79, %f110, %f92;
	mul.f32 	%f81, %f111, %f92;
	mul.f32 	%f82, %f112, %f92;
	mul.f32 	%f83, %f113, %f92;
	mul.f32 	%f85, %f114, %f92;
	mul.f32 	%f86, %f115, %f92;
	mul.f32 	%f87, %f116, %f92;
	mul.f32 	%f89, %f117, %f92;
	mul.f32 	%f90, %f118, %f92;
	mul.f32 	%f91, %f119, %f92;
	// begin inline asm
	{  cvt.rn.f16.f32 %rs59, %f59;}

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

	// end inline asm
	// begin inline asm
	{  cvt.rn.f16.f32 %rs57, %f57;}

	// end inline asm
	// begin inline asm
	{  cvt.rn.f16.f32 %rs60, %f92;}

	// end inline asm
	st.global.v4.u16 	[%rd40], {%rs57, %rs58, %rs59, %rs60};
	// begin inline asm
	{  cvt.rn.f16.f32 %rs63, %f63;}

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

	// end inline asm
	// begin inline asm
	{  cvt.rn.f16.f32 %rs61, %f61;}

	// end inline asm
	// begin inline asm
	{  cvt.rn.f16.f32 %rs64, %f92;}

	// end inline asm
	st.global.v4.u16 	[%rd44], {%rs61, %rs62, %rs63, %rs64};
	// begin inline asm
	{  cvt.rn.f16.f32 %rs67, %f67;}

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

	// end inline asm
	// begin inline asm
	{  cvt.rn.f16.f32 %rs65, %f65;}

	// end inline asm
	// begin inline asm
	{  cvt.rn.f16.f32 %rs68, %f92;}

	// end inline asm
	st.global.v4.u16 	[%rd48], {%rs65, %rs66, %rs67, %rs68};
	// begin inline asm
	{  cvt.rn.f16.f32 %rs71, %f71;}

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

	// end inline asm
	// begin inline asm
	{  cvt.rn.f16.f32 %rs69, %f69;}

	// end inline asm
	// begin inline asm
	{  cvt.rn.f16.f32 %rs72, %f92;}

	// end inline asm
	st.global.v4.u16 	[%rd52], {%rs69, %rs70, %rs71, %rs72};
	// begin inline asm
	{  cvt.rn.f16.f32 %rs75, %f75;}

	// end inline asm
	// begin inline asm
	{  cvt.rn.f16.f32 %rs74, %f74;}

	// end inline asm
	// begin inline asm
	{  cvt.rn.f16.f32 %rs73, %f73;}

	// end inline asm
	// begin inline asm
	{  cvt.rn.f16.f32 %rs76, %f92;}

	// end inline asm
	st.global.v4.u16 	[%rd56], {%rs73, %rs74, %rs75, %rs76};
	// begin inline asm
	{  cvt.rn.f16.f32 %rs79, %f79;}

	// end inline asm
	// begin inline asm
	{  cvt.rn.f16.f32 %rs78, %f78;}

	// end inline asm
	// begin inline asm
	{  cvt.rn.f16.f32 %rs77, %f77;}

	// end inline asm
	// begin inline asm
	{  cvt.rn.f16.f32 %rs80, %f92;}

	// end inline asm
	st.global.v4.u16 	[%rd60], {%rs77, %rs78, %rs79, %rs80};
	// begin inline asm
	{  cvt.rn.f16.f32 %rs83, %f83;}

	// end inline asm
	// begin inline asm
	{  cvt.rn.f16.f32 %rs82, %f82;}

	// end inline asm
	// begin inline asm
	{  cvt.rn.f16.f32 %rs81, %f81;}

	// end inline asm
	// begin inline asm
	{  cvt.rn.f16.f32 %rs84, %f92;}

	// end inline asm
	st.global.v4.u16 	[%rd64], {%rs81, %rs82, %rs83, %rs84};
	// begin inline asm
	{  cvt.rn.f16.f32 %rs87, %f87;}

	// end inline asm
	// begin inline asm
	{  cvt.rn.f16.f32 %rs86, %f86;}

	// end inline asm
	// begin inline asm
	{  cvt.rn.f16.f32 %rs85, %f85;}

	// end inline asm
	// begin inline asm
	{  cvt.rn.f16.f32 %rs88, %f92;}

	// end inline asm
	st.global.v4.u16 	[%rd68], {%rs85, %rs86, %rs87, %rs88};
	// begin inline asm
	{  cvt.rn.f16.f32 %rs91, %f91;}

	// end inline asm
	// begin inline asm
	{  cvt.rn.f16.f32 %rs90, %f90;}

	// end inline asm
	// begin inline asm
	{  cvt.rn.f16.f32 %rs89, %f89;}

	// end inline asm
	// begin inline asm
	{  cvt.rn.f16.f32 %rs92, %f92;}

	// end inline asm
	st.global.v4.u16 	[%rd72], {%rs89, %rs90, %rs91, %rs92};
	ret;

}