//
// 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<2>;
	.reg .b16 	%rs<23>;
	.reg .b32 	%r<15>;
	.reg .b64 	%rd<13>;


	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, 4;
	add.s64 	%rd4, %rd2, %rd3;
	ld.global.v4.u8 	{%rs1, %rs2, %rs3, %rs4}, [%rd4];
	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, 4;
	add.s64 	%rd8, %rd6, %rd7;
	ld.global.v4.u8 	{%rs8, %rs9, %rs10, %rs11}, [%rd8];
	ld.const.u64 	%rd9, [params+176];
	cvta.to.global.u64 	%rd10, %rd9;
	ld.const.u32 	%r13, [params+168];
	mad.lo.s32 	%r14, %r13, %r8, %r4;
	mul.wide.u32 	%rd11, %r14, 4;
	add.s64 	%rd12, %rd10, %rd11;
	ld.global.u8 	%rs15, [%rd12+3];
	setp.eq.s16 	%p1, %rs15, 0;
	selp.b16 	%rs16, 0, %rs1, %p1;
	selp.b16 	%rs17, 0, %rs2, %p1;
	selp.b16 	%rs18, 0, %rs3, %p1;
	max.u16 	%rs19, %rs11, %rs15;
	add.s16 	%rs20, %rs18, %rs10;
	add.s16 	%rs21, %rs17, %rs9;
	add.s16 	%rs22, %rs16, %rs8;
	st.global.v4.u8 	[%rd8], {%rs22, %rs21, %rs20, %rs19};
	ret;

}