//
// 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 .pred 	%p<5>;
	.reg .b16 	%rs<29>;
	.reg .f32 	%f<91>;
	.reg .b32 	%r<31>;
	.reg .b64 	%rd<21>;


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

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

	// end inline asm
	ld.global.u16 	%rs10, [%rd6+4];
	// begin inline asm
	{  cvt.f32.f16 %f19, %rs10;}

	// end inline asm
	mul.f32 	%f20, %f18, 0f3F372474;
	fma.rn.f32 	%f21, %f17, 0f3E59999A, %f20;
	fma.rn.f32 	%f22, %f19, 0f3D93A92A, %f21;
	ld.const.u64 	%rd7, [params+160];
	cvta.to.global.u64 	%rd8, %rd7;
	ld.const.u32 	%r11, [params+152];
	mad.lo.s32 	%r12, %r11, %r2, %r1;
	mul.wide.u32 	%rd9, %r12, 16;
	add.s64 	%rd1, %rd8, %rd9;
	ld.global.v4.f32 	{%f23, %f24, %f25, %f26}, [%rd1];
	ld.const.u64 	%rd10, [params+144];
	cvta.to.global.u64 	%rd11, %rd10;
	ld.const.u32 	%r13, [params+136];
	mad.lo.s32 	%r14, %r13, %r2, %r1;
	mul.wide.u32 	%rd12, %r14, 4;
	add.s64 	%rd13, %rd11, %rd12;
	ld.global.v4.u8 	{%rs11, %rs12, %rs13, %rs14}, [%rd13];
	cvt.rn.f32.u16 	%f30, %rs11;
	div.rn.f32 	%f31, %f30, 0f437F0000;
	fma.rn.f32 	%f32, %f31, 0f40000000, 0fBF800000;
	cvt.rn.f32.u16 	%f33, %rs12;
	div.rn.f32 	%f34, %f33, 0f437F0000;
	fma.rn.f32 	%f35, %f34, 0f40000000, 0fBF800000;
	cvt.rn.f32.u16 	%f36, %rs13;
	div.rn.f32 	%f37, %f36, 0f437F0000;
	fma.rn.f32 	%f38, %f37, 0f40000000, 0fBF800000;
	mul.f32 	%f39, %f35, %f35;
	fma.rn.f32 	%f40, %f32, %f32, %f39;
	fma.rn.f32 	%f41, %f38, %f38, %f40;
	sqrt.rn.f32 	%f42, %f41;
	rcp.rn.f32 	%f43, %f42;
	mul.f32 	%f44, %f32, %f43;
	mul.f32 	%f45, %f35, %f43;
	mul.f32 	%f46, %f43, %f38;
	fma.rn.f32 	%f47, %f22, %f44, %f23;
	fma.rn.f32 	%f48, %f22, %f45, %f24;
	fma.rn.f32 	%f49, %f22, %f46, %f25;
	cvt.rn.f32.u16 	%f50, %rs14;
	div.rn.f32 	%f51, %f50, 0f437F0000;
	ld.global.f32 	%f52, [%rd1+12];
	min.f32 	%f90, %f51, %f52;
	mul.f32 	%f2, %f47, %f90;
	mul.f32 	%f3, %f48, %f90;
	mul.f32 	%f4, %f49, %f90;
	ld.const.f32 	%f53, [params+1168];
	setp.gt.f32 	%p1, %f53, 0f3F000000;
	@%p1 bra 	$L__BB0_2;
	bra.uni 	$L__BB0_1;

$L__BB0_2:
	mul.f32 	%f54, %f3, %f3;
	fma.rn.f32 	%f55, %f2, %f2, %f54;
	fma.rn.f32 	%f56, %f4, %f4, %f55;
	sqrt.rn.f32 	%f57, %f56;
	rcp.rn.f32 	%f58, %f57;
	mul.f32 	%f5, %f2, %f58;
	mul.f32 	%f6, %f3, %f58;
	mul.f32 	%f7, %f4, %f58;
	ld.const.u64 	%rd14, [params+400];
	cvta.to.global.u64 	%rd15, %rd14;
	ld.const.u32 	%r15, [params+392];
	mad.lo.s32 	%r16, %r15, %r2, %r1;
	mul.wide.u32 	%rd16, %r16, 4;
	add.s64 	%rd2, %rd15, %rd16;
	ld.global.v2.u8 	{%rs15, %rs28}, [%rd2];
	or.b16  	%rs17, %rs15, %rs28;
	and.b16  	%rs18, %rs17, 255;
	setp.eq.s16 	%p2, %rs18, 0;
	@%p2 bra 	$L__BB0_4;

	ld.global.u8 	%rs27, [%rd2+2];
	bra.uni 	$L__BB0_5;

$L__BB0_1:
	st.global.v4.f32 	[%rd1], {%f2, %f3, %f4, %f90};
	bra.uni 	$L__BB0_9;

$L__BB0_4:
	ld.global.u8 	%rs27, [%rd2+2];
	setp.eq.s16 	%p3, %rs27, 0;
	mov.f32 	%f87, 0f00000000;
	mov.u16 	%rs28, 0;
	mov.f32 	%f88, %f87;
	mov.f32 	%f89, %f87;
	@%p3 bra 	$L__BB0_6;

$L__BB0_5:
	cvt.rn.f32.u16 	%f62, %rs15;
	div.rn.f32 	%f63, %f62, 0f437F0000;
	fma.rn.f32 	%f64, %f63, 0f40000000, 0fBF800000;
	and.b16  	%rs21, %rs28, 255;
	cvt.rn.f32.u16 	%f65, %rs21;
	div.rn.f32 	%f66, %f65, 0f437F0000;
	fma.rn.f32 	%f67, %f66, 0f40000000, 0fBF800000;
	cvt.rn.f32.u16 	%f68, %rs27;
	div.rn.f32 	%f69, %f68, 0f437F0000;
	fma.rn.f32 	%f70, %f69, 0f40000000, 0fBF800000;
	mul.f32 	%f71, %f67, %f67;
	fma.rn.f32 	%f72, %f64, %f64, %f71;
	fma.rn.f32 	%f73, %f70, %f70, %f72;
	sqrt.rn.f32 	%f74, %f73;
	rcp.rn.f32 	%f75, %f74;
	mul.f32 	%f89, %f75, %f70;
	mul.f32 	%f88, %f75, %f67;
	mul.f32 	%f87, %f64, %f75;

$L__BB0_6:
	mul.f32 	%f76, %f6, %f88;
	fma.rn.f32 	%f77, %f5, %f87, %f76;
	fma.rn.f32 	%f14, %f7, %f89, %f77;
	setp.leu.f32 	%p4, %f90, 0f00000000;
	@%p4 bra 	$L__BB0_8;

	fma.rn.f32 	%f78, %f14, 0f3F000000, 0f3F000000;
	mov.f32 	%f79, 0f3B808081;
	max.f32 	%f90, %f78, %f79;

$L__BB0_8:
	ld.const.u64 	%rd17, [params+1160];
	cvta.to.global.u64 	%rd18, %rd17;
	ld.const.u32 	%r21, [params+1152];
	mad.lo.s32 	%r26, %r21, %r2, %r1;
	fma.rn.f32 	%f80, %f5, 0f3F000000, 0f3F000000;
	mul.f32 	%f81, %f80, 0f437F0000;
	cvt.rzi.u32.f32 	%r27, %f81;
	fma.rn.f32 	%f82, %f6, 0f3F000000, 0f3F000000;
	mul.f32 	%f83, %f82, 0f437F0000;
	cvt.rzi.u32.f32 	%r28, %f83;
	fma.rn.f32 	%f84, %f7, 0f3F000000, 0f3F000000;
	mul.f32 	%f85, %f84, 0f437F0000;
	cvt.rzi.u32.f32 	%r29, %f85;
	mul.f32 	%f86, %f90, 0f437F0000;
	cvt.rzi.u32.f32 	%r30, %f86;
	mul.wide.u32 	%rd19, %r26, 4;
	add.s64 	%rd20, %rd18, %rd19;
	cvt.u16.u32 	%rs23, %r30;
	cvt.u16.u32 	%rs24, %r29;
	cvt.u16.u32 	%rs25, %r28;
	cvt.u16.u32 	%rs26, %r27;
	st.global.v4.u8 	[%rd20], {%rs26, %rs25, %rs24, %rs23};

$L__BB0_9:
	ret;

}