//
// 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<13>;
	.reg .f32 	%f<16>;
	.reg .b32 	%r<13>;
	.reg .b64 	%rd<9>;


	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+144];
	cvta.to.global.u64 	%rd2, %rd1;
	ld.const.u32 	%r9, [params+136];
	mad.lo.s32 	%r10, %r9, %r8, %r4;
	mul.wide.u32 	%rd3, %r10, 8;
	add.s64 	%rd4, %rd2, %rd3;
	ld.global.u16 	%rs1, [%rd4+6];
	// begin inline asm
	{  cvt.f32.f16 %f1, %rs1;}

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

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

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

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

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

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

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

	// end inline asm
	add.f32 	%f13, %f2, %f6;
	add.f32 	%f14, %f3, %f7;
	add.f32 	%f15, %f4, %f8;
	mul.f32 	%f9, %f1, %f13;
	mul.f32 	%f10, %f1, %f14;
	mul.f32 	%f11, %f1, %f15;
	min.f32 	%f12, %f1, %f5;
	// begin inline asm
	{  cvt.rn.f16.f32 %rs12, %f12;}

	// end inline asm
	// begin inline asm
	{  cvt.rn.f16.f32 %rs11, %f11;}

	// end inline asm
	// begin inline asm
	{  cvt.rn.f16.f32 %rs10, %f10;}

	// end inline asm
	// begin inline asm
	{  cvt.rn.f16.f32 %rs9, %f9;}

	// end inline asm
	st.global.v4.u16 	[%rd8], {%rs9, %rs10, %rs11, %rs12};
	ret;

}