//
// 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<14>;
	.reg .b16 	%rs<19>;
	.reg .f32 	%f<41>;
	.reg .b32 	%r<56>;
	.reg .b64 	%rd<27>;


	mov.u32 	%r25, %ctaid.x;
	mov.u32 	%r1, %ntid.x;
	mov.u32 	%r26, %tid.x;
	mad.lo.s32 	%r2, %r25, %r1, %r26;
	mov.u32 	%r3, %ntid.y;
	mov.u32 	%r27, %ctaid.y;
	mov.u32 	%r28, %tid.y;
	mad.lo.s32 	%r4, %r27, %r3, %r28;
	ld.const.u64 	%rd10, [params+144];
	cvta.to.global.u64 	%rd1, %rd10;
	ld.const.u32 	%r5, [params+136];
	mul.lo.s32 	%r6, %r5, %r4;
	add.s32 	%r29, %r6, %r2;
	mul.wide.u32 	%rd11, %r29, 2;
	add.s64 	%rd12, %rd1, %rd11;
	add.s64 	%rd2, %rd12, 1;
	ld.global.u8 	%rs1, [%rd12+1];
	setp.ne.s16 	%p1, %rs1, 0;
	@%p1 bra 	$L__BB0_17;

	ld.global.u8 	%rs2, [%rd2+-1];
	cvt.rn.f32.u16 	%f16, %rs2;
	div.rn.f32 	%f35, %f16, 0f437F0000;
	add.s32 	%r31, %r2, -1;
	setp.eq.s32 	%p2, %r2, 0;
	mov.u32 	%r50, 0;
	selp.b32 	%r7, 0, %r31, %p2;
	add.s32 	%r32, %r4, -1;
	setp.eq.s32 	%p3, %r4, 0;
	selp.b32 	%r33, 0, %r32, %p3;
	mov.u32 	%r34, %nctaid.x;
	mad.lo.s32 	%r35, %r34, %r1, -1;
	setp.eq.s32 	%p4, %r2, %r35;
	add.s32 	%r36, %r2, 1;
	selp.b32 	%r8, %r35, %r36, %p4;
	mov.u32 	%r37, %nctaid.y;
	mad.lo.s32 	%r38, %r37, %r3, -1;
	setp.eq.s32 	%p5, %r4, %r38;
	add.s32 	%r39, %r4, 1;
	selp.b32 	%r9, %r38, %r39, %p5;
	mul.lo.s32 	%r10, %r5, %r33;
	add.s32 	%r40, %r10, %r7;
	mul.wide.u32 	%rd13, %r40, 2;
	add.s64 	%rd14, %rd1, %rd13;
	add.s64 	%rd3, %rd14, 1;
	ld.global.u8 	%rs3, [%rd14+1];
	setp.eq.s16 	%p6, %rs3, 0;
	@%p6 bra 	$L__BB0_3;

	ld.global.u8 	%rs4, [%rd3+-1];
	cvt.rn.f32.u16 	%f17, %rs4;
	div.rn.f32 	%f18, %f17, 0f437F0000;
	add.f32 	%f35, %f35, %f18;
	mov.u32 	%r50, 1;

$L__BB0_3:
	add.s32 	%r42, %r10, %r2;
	mul.wide.u32 	%rd15, %r42, 2;
	add.s64 	%rd16, %rd1, %rd15;
	add.s64 	%rd4, %rd16, 1;
	ld.global.u8 	%rs5, [%rd16+1];
	setp.eq.s16 	%p7, %rs5, 0;
	@%p7 bra 	$L__BB0_5;

	ld.global.u8 	%rs6, [%rd4+-1];
	cvt.rn.f32.u16 	%f19, %rs6;
	div.rn.f32 	%f20, %f19, 0f437F0000;
	add.f32 	%f35, %f35, %f20;
	add.s32 	%r50, %r50, 1;

$L__BB0_5:
	add.s32 	%r43, %r10, %r8;
	mul.wide.u32 	%rd17, %r43, 2;
	add.s64 	%rd18, %rd1, %rd17;
	add.s64 	%rd5, %rd18, 1;
	ld.global.u8 	%rs7, [%rd18+1];
	setp.eq.s16 	%p8, %rs7, 0;
	@%p8 bra 	$L__BB0_7;

	ld.global.u8 	%rs8, [%rd5+-1];
	cvt.rn.f32.u16 	%f21, %rs8;
	div.rn.f32 	%f22, %f21, 0f437F0000;
	add.f32 	%f35, %f35, %f22;
	add.s32 	%r50, %r50, 1;

$L__BB0_7:
	add.s32 	%r44, %r6, %r7;
	mul.wide.u32 	%rd19, %r44, 2;
	add.s64 	%rd20, %rd1, %rd19;
	add.s64 	%rd6, %rd20, 1;
	ld.global.u8 	%rs9, [%rd20+1];
	setp.eq.s16 	%p9, %rs9, 0;
	@%p9 bra 	$L__BB0_9;

	ld.global.u8 	%rs10, [%rd6+-1];
	cvt.rn.f32.u16 	%f23, %rs10;
	div.rn.f32 	%f24, %f23, 0f437F0000;
	add.f32 	%f35, %f35, %f24;
	add.s32 	%r50, %r50, 1;

$L__BB0_9:
	add.s32 	%r45, %r6, %r8;
	mul.wide.u32 	%rd21, %r45, 2;
	add.s64 	%rd22, %rd1, %rd21;
	add.s64 	%rd7, %rd22, 1;
	ld.global.u8 	%rs11, [%rd22+1];
	setp.eq.s16 	%p10, %rs11, 0;
	@%p10 bra 	$L__BB0_11;

	ld.global.u8 	%rs12, [%rd7+-1];
	cvt.rn.f32.u16 	%f25, %rs12;
	div.rn.f32 	%f26, %f25, 0f437F0000;
	add.f32 	%f35, %f35, %f26;
	add.s32 	%r50, %r50, 1;

$L__BB0_11:
	mul.lo.s32 	%r20, %r5, %r9;
	add.s32 	%r46, %r20, %r7;
	mul.wide.u32 	%rd23, %r46, 2;
	add.s64 	%rd24, %rd1, %rd23;
	add.s64 	%rd8, %rd24, 1;
	ld.global.u8 	%rs13, [%rd24+1];
	setp.eq.s16 	%p11, %rs13, 0;
	@%p11 bra 	$L__BB0_13;

	ld.global.u8 	%rs14, [%rd8+-1];
	cvt.rn.f32.u16 	%f27, %rs14;
	div.rn.f32 	%f28, %f27, 0f437F0000;
	add.f32 	%f35, %f35, %f28;
	add.s32 	%r50, %r50, 1;

$L__BB0_13:
	add.s32 	%r47, %r20, %r8;
	mul.wide.u32 	%rd25, %r47, 2;
	add.s64 	%rd26, %rd1, %rd25;
	add.s64 	%rd9, %rd26, 1;
	ld.global.u8 	%rs15, [%rd26+1];
	setp.eq.s16 	%p12, %rs15, 0;
	@%p12 bra 	$L__BB0_15;

	ld.global.u8 	%rs16, [%rd9+-1];
	cvt.rn.f32.u16 	%f29, %rs16;
	div.rn.f32 	%f30, %f29, 0f437F0000;
	add.f32 	%f35, %f35, %f30;
	add.s32 	%r50, %r50, 1;

$L__BB0_15:
	setp.eq.s32 	%p13, %r50, 0;
	@%p13 bra 	$L__BB0_17;

	cvt.rn.f32.u32 	%f31, %r50;
	div.rn.f32 	%f32, %f35, %f31;
	mul.f32 	%f33, %f32, 0f437F0000;
	cvt.rzi.u32.f32 	%r48, %f33;
	cvt.u16.u32 	%rs17, %r48;
	mov.u16 	%rs18, 255;
	st.global.v2.u8 	[%rd2+-1], {%rs17, %rs18};

$L__BB0_17:
	ret;

}