//
// Generated by NVIDIA NVVM Compiler
//
// Compiler Build ID: CL-23083092
// Cuda compilation tools, release 9.1, V9.1.85
// Based on LLVM 3.4svn
//

.version 6.1
.target sm_30
.address_size 64

	// .globl	_Z6oxMainv
.global .align 8 .b8 pixelID[8];
.global .align 8 .b8 resolution[8];
.global .align 4 .b8 normal[12];
.global .align 4 .b8 camPos[12];
.global .align 4 .b8 root[4];
.global .align 4 .u32 imageEnabled;
.global .texref lightmap;
.global .align 16 .b8 tileInfo[16];
.global .align 4 .u32 additive;
.global .align 1 .b8 image[1];
.global .align 1 .b8 image_HDR[1];
.global .align 1 .b8 image_HDR2[1];
.global .align 1 .b8 image_Mask[1];
.global .align 1 .b8 image_RNM0[1];
.global .align 1 .b8 image_RNM1[1];
.global .align 1 .b8 image_RNM2[1];
.global .align 1 .b8 image_RNM3[1];
.global .align 1 .b8 uvpos[1];
.global .align 1 .b8 uvnormal[1];
.global .align 1 .b8 lightMeshBuffer[1];
.global .align 4 .u32 lightMeshBufferSize;
.global .align 4 .f32 lightInvCutoff;
.global .align 4 .f32 lightPointSize;
.global .align 4 .b8 lightColor[12];
.global .align 1 .b8 rnd_seeds[1];
.global .align 4 .u32 samples;
.global .align 4 .b8 _ZN21rti_internal_typeinfo7pixelIDE[8] = {82, 97, 121, 0, 8, 0, 0, 0};
.global .align 4 .b8 _ZN21rti_internal_typeinfo10resolutionE[8] = {82, 97, 121, 0, 8, 0, 0, 0};
.global .align 4 .b8 _ZN21rti_internal_typeinfo6normalE[8] = {82, 97, 121, 0, 12, 0, 0, 0};
.global .align 4 .b8 _ZN21rti_internal_typeinfo6camPosE[8] = {82, 97, 121, 0, 12, 0, 0, 0};
.global .align 4 .b8 _ZN21rti_internal_typeinfo4rootE[8] = {82, 97, 121, 0, 4, 0, 0, 0};
.global .align 4 .b8 _ZN21rti_internal_typeinfo12imageEnabledE[8] = {82, 97, 121, 0, 4, 0, 0, 0};
.global .align 4 .b8 _ZN21rti_internal_typeinfo8tileInfoE[8] = {82, 97, 121, 0, 16, 0, 0, 0};
.global .align 4 .b8 _ZN21rti_internal_typeinfo8additiveE[8] = {82, 97, 121, 0, 4, 0, 0, 0};
.global .align 4 .b8 _ZN21rti_internal_typeinfo19lightMeshBufferSizeE[8] = {82, 97, 121, 0, 4, 0, 0, 0};
.global .align 4 .b8 _ZN21rti_internal_typeinfo14lightInvCutoffE[8] = {82, 97, 121, 0, 4, 0, 0, 0};
.global .align 4 .b8 _ZN21rti_internal_typeinfo14lightPointSizeE[8] = {82, 97, 121, 0, 4, 0, 0, 0};
.global .align 4 .b8 _ZN21rti_internal_typeinfo10lightColorE[8] = {82, 97, 121, 0, 12, 0, 0, 0};
.global .align 4 .b8 _ZN21rti_internal_typeinfo7samplesE[8] = {82, 97, 121, 0, 4, 0, 0, 0};
.global .align 8 .u64 _ZN21rti_internal_register20reg_bitness_detectorE;
.global .align 8 .u64 _ZN21rti_internal_register24reg_exception_64_detail0E;
.global .align 8 .u64 _ZN21rti_internal_register24reg_exception_64_detail1E;
.global .align 8 .u64 _ZN21rti_internal_register24reg_exception_64_detail2E;
.global .align 8 .u64 _ZN21rti_internal_register24reg_exception_64_detail3E;
.global .align 8 .u64 _ZN21rti_internal_register24reg_exception_64_detail4E;
.global .align 8 .u64 _ZN21rti_internal_register24reg_exception_64_detail5E;
.global .align 8 .u64 _ZN21rti_internal_register24reg_exception_64_detail6E;
.global .align 8 .u64 _ZN21rti_internal_register24reg_exception_64_detail7E;
.global .align 8 .u64 _ZN21rti_internal_register24reg_exception_64_detail8E;
.global .align 8 .u64 _ZN21rti_internal_register24reg_exception_64_detail9E;
.global .align 4 .u32 _ZN21rti_internal_register21reg_exception_detail0E;
.global .align 4 .u32 _ZN21rti_internal_register21reg_exception_detail1E;
.global .align 4 .u32 _ZN21rti_internal_register21reg_exception_detail2E;
.global .align 4 .u32 _ZN21rti_internal_register21reg_exception_detail3E;
.global .align 4 .u32 _ZN21rti_internal_register21reg_exception_detail4E;
.global .align 4 .u32 _ZN21rti_internal_register21reg_exception_detail5E;
.global .align 4 .u32 _ZN21rti_internal_register21reg_exception_detail6E;
.global .align 4 .u32 _ZN21rti_internal_register21reg_exception_detail7E;
.global .align 4 .u32 _ZN21rti_internal_register21reg_exception_detail8E;
.global .align 4 .u32 _ZN21rti_internal_register21reg_exception_detail9E;
.global .align 4 .u32 _ZN21rti_internal_register14reg_rayIndex_xE;
.global .align 4 .u32 _ZN21rti_internal_register14reg_rayIndex_yE;
.global .align 4 .u32 _ZN21rti_internal_register14reg_rayIndex_zE;
.global .align 8 .b8 _ZN21rti_internal_typename7pixelIDE[6] = {117, 105, 110, 116, 50, 0};
.global .align 8 .b8 _ZN21rti_internal_typename10resolutionE[6] = {117, 105, 110, 116, 50, 0};
.global .align 8 .b8 _ZN21rti_internal_typename6normalE[7] = {102, 108, 111, 97, 116, 51, 0};
.global .align 8 .b8 _ZN21rti_internal_typename6camPosE[7] = {102, 108, 111, 97, 116, 51, 0};
.global .align 16 .b8 _ZN21rti_internal_typename4rootE[9] = {114, 116, 79, 98, 106, 101, 99, 116, 0};
.global .align 4 .b8 _ZN21rti_internal_typename12imageEnabledE[4] = {105, 110, 116, 0};
.global .align 8 .b8 _ZN21rti_internal_typename8tileInfoE[6] = {117, 105, 110, 116, 52, 0};
.global .align 4 .b8 _ZN21rti_internal_typename8additiveE[4] = {105, 110, 116, 0};
.global .align 16 .b8 _ZN21rti_internal_typename19lightMeshBufferSizeE[13] = {117, 110, 115, 105, 103, 110, 101, 100, 32, 105, 110, 116, 0};
.global .align 8 .b8 _ZN21rti_internal_typename14lightInvCutoffE[6] = {102, 108, 111, 97, 116, 0};
.global .align 8 .b8 _ZN21rti_internal_typename14lightPointSizeE[6] = {102, 108, 111, 97, 116, 0};
.global .align 8 .b8 _ZN21rti_internal_typename10lightColorE[7] = {102, 108, 111, 97, 116, 51, 0};
.global .align 4 .b8 _ZN21rti_internal_typename7samplesE[4] = {105, 110, 116, 0};
.global .align 4 .u32 _ZN21rti_internal_typeenum7pixelIDE = 4919;
.global .align 4 .u32 _ZN21rti_internal_typeenum10resolutionE = 4919;
.global .align 4 .u32 _ZN21rti_internal_typeenum6normalE = 4919;
.global .align 4 .u32 _ZN21rti_internal_typeenum6camPosE = 4919;
.global .align 4 .u32 _ZN21rti_internal_typeenum4rootE = 4919;
.global .align 4 .u32 _ZN21rti_internal_typeenum12imageEnabledE = 4919;
.global .align 4 .u32 _ZN21rti_internal_typeenum8tileInfoE = 4919;
.global .align 4 .u32 _ZN21rti_internal_typeenum8additiveE = 4919;
.global .align 4 .u32 _ZN21rti_internal_typeenum19lightMeshBufferSizeE = 4919;
.global .align 4 .u32 _ZN21rti_internal_typeenum14lightInvCutoffE = 4919;
.global .align 4 .u32 _ZN21rti_internal_typeenum14lightPointSizeE = 4919;
.global .align 4 .u32 _ZN21rti_internal_typeenum10lightColorE = 4919;
.global .align 4 .u32 _ZN21rti_internal_typeenum7samplesE = 4919;
.global .align 16 .b8 _ZN21rti_internal_semantic7pixelIDE[14] = {114, 116, 76, 97, 117, 110, 99, 104, 73, 110, 100, 101, 120, 0};
.global .align 16 .b8 _ZN21rti_internal_semantic10resolutionE[12] = {114, 116, 76, 97, 117, 110, 99, 104, 68, 105, 109, 0};
.global .align 16 .b8 _ZN21rti_internal_semantic6normalE[17] = {97, 116, 116, 114, 105, 98, 117, 116, 101, 32, 110, 111, 114, 109, 97, 108, 0};
.global .align 1 .b8 _ZN21rti_internal_semantic6camPosE[1];
.global .align 1 .b8 _ZN21rti_internal_semantic4rootE[1];
.global .align 1 .b8 _ZN21rti_internal_semantic12imageEnabledE[1];
.global .align 1 .b8 _ZN21rti_internal_semantic8tileInfoE[1];
.global .align 1 .b8 _ZN21rti_internal_semantic8additiveE[1];
.global .align 1 .b8 _ZN21rti_internal_semantic19lightMeshBufferSizeE[1];
.global .align 1 .b8 _ZN21rti_internal_semantic14lightInvCutoffE[1];
.global .align 1 .b8 _ZN21rti_internal_semantic14lightPointSizeE[1];
.global .align 1 .b8 _ZN21rti_internal_semantic10lightColorE[1];
.global .align 1 .b8 _ZN21rti_internal_semantic7samplesE[1];
.global .align 1 .b8 _ZN23rti_internal_annotation7pixelIDE[1];
.global .align 1 .b8 _ZN23rti_internal_annotation10resolutionE[1];
.global .align 1 .b8 _ZN23rti_internal_annotation6normalE[1];
.global .align 1 .b8 _ZN23rti_internal_annotation6camPosE[1];
.global .align 1 .b8 _ZN23rti_internal_annotation4rootE[1];
.global .align 1 .b8 _ZN23rti_internal_annotation12imageEnabledE[1];
.global .align 1 .b8 _ZN23rti_internal_annotation8tileInfoE[1];
.global .align 1 .b8 _ZN23rti_internal_annotation8additiveE[1];
.global .align 1 .b8 _ZN23rti_internal_annotation19lightMeshBufferSizeE[1];
.global .align 1 .b8 _ZN23rti_internal_annotation14lightInvCutoffE[1];
.global .align 1 .b8 _ZN23rti_internal_annotation14lightPointSizeE[1];
.global .align 1 .b8 _ZN23rti_internal_annotation10lightColorE[1];
.global .align 1 .b8 _ZN23rti_internal_annotation7samplesE[1];
.const .align 4 .b8 __cudart_i2opi_f[24] = {65, 144, 67, 60, 153, 149, 98, 219, 192, 221, 52, 245, 209, 87, 39, 252, 41, 21, 68, 78, 110, 131, 249, 162};

.visible .entry _Z6oxMainv(

)
{
	.local .align 4 .b8 	__local_depot0[36];
	.reg .b64 	%SP;
	.reg .b64 	%SPL;
	.reg .pred 	%p<149>;
	.reg .b16 	%rs<152>;
	.reg .f32 	%f<988>;
	.reg .b32 	%r<411>;
	.reg .b64 	%rd<298>;


	mov.u64 	%rd297, __local_depot0;
	cvta.local.u64 	%SP, %rd297;
	ld.global.u32 	%r1, [samples];
	shl.b32 	%r2, %r1, 1;
	ld.global.v2.u32 	{%r100, %r101}, [pixelID];
	cvt.u64.u32	%rd21, %r100;
	cvt.u64.u32	%rd22, %r101;
	mov.u64 	%rd25, uvnormal;
	cvta.global.u64 	%rd20, %rd25;
	mov.u32 	%r98, 2;
	mov.u32 	%r99, 4;
	mov.u64 	%rd24, 0;
	// inline asm
	call (%rd19), _rt_buffer_get_64, (%rd20, %r98, %r99, %rd21, %rd22, %rd24, %rd24);
	// inline asm
	ld.u32 	%r3, [%rd19];
	shr.u32 	%r104, %r3, 16;
	cvt.u16.u32	%rs1, %r104;
	and.b16  	%rs7, %rs1, 255;
	cvt.u16.u32	%rs8, %r3;
	or.b16  	%rs9, %rs8, %rs7;
	setp.eq.s16	%p7, %rs9, 0;
	mov.f32 	%f916, 0f00000000;
	mov.f32 	%f917, %f916;
	mov.f32 	%f918, %f916;
	@%p7 bra 	BB0_2;

	ld.u8 	%rs10, [%rd19+1];
	and.b16  	%rs12, %rs8, 255;
	cvt.rn.f32.u16	%f217, %rs12;
	div.rn.f32 	%f218, %f217, 0f437F0000;
	fma.rn.f32 	%f219, %f218, 0f40000000, 0fBF800000;
	cvt.rn.f32.u16	%f220, %rs10;
	div.rn.f32 	%f221, %f220, 0f437F0000;
	fma.rn.f32 	%f222, %f221, 0f40000000, 0fBF800000;
	cvt.rn.f32.u16	%f223, %rs7;
	div.rn.f32 	%f224, %f223, 0f437F0000;
	fma.rn.f32 	%f225, %f224, 0f40000000, 0fBF800000;
	mul.f32 	%f226, %f222, %f222;
	fma.rn.f32 	%f227, %f219, %f219, %f226;
	fma.rn.f32 	%f228, %f225, %f225, %f227;
	sqrt.rn.f32 	%f229, %f228;
	rcp.rn.f32 	%f230, %f229;
	mul.f32 	%f916, %f219, %f230;
	mul.f32 	%f917, %f222, %f230;
	mul.f32 	%f918, %f225, %f230;

BB0_2:
	ld.global.v2.u32 	{%r105, %r106}, [pixelID];
	ld.global.v2.u32 	{%r108, %r109}, [tileInfo];
	add.s32 	%r4, %r105, %r108;
	add.s32 	%r5, %r106, %r109;
	setp.eq.f32	%p8, %f917, 0f00000000;
	setp.eq.f32	%p9, %f916, 0f00000000;
	and.pred  	%p10, %p9, %p8;
	setp.eq.f32	%p11, %f918, 0f00000000;
	and.pred  	%p12, %p10, %p11;
	@%p12 bra 	BB0_139;
	bra.uni 	BB0_3;

BB0_139:
	ld.global.u32 	%r410, [imageEnabled];
	and.b32  	%r334, %r410, 1;
	setp.eq.b32	%p140, %r334, 1;
	@!%p140 bra 	BB0_141;
	bra.uni 	BB0_140;

BB0_140:
	cvt.u64.u32	%rd179, %r4;
	cvt.u64.u32	%rd180, %r5;
	mov.u64 	%rd183, image;
	cvta.global.u64 	%rd178, %rd183;
	// inline asm
	call (%rd177), _rt_buffer_get_64, (%rd178, %r98, %r99, %rd179, %rd180, %rd24, %rd24);
	// inline asm
	mov.u16 	%rs84, 0;
	st.v4.u8 	[%rd177], {%rs84, %rs84, %rs84, %rs84};
	ld.global.u32 	%r410, [imageEnabled];

BB0_141:
	and.b32  	%r337, %r410, 8;
	setp.eq.s32	%p141, %r337, 0;
	@%p141 bra 	BB0_143;

	cvt.u64.u32	%rd187, %r5;
	cvt.u64.u32	%rd186, %r4;
	mov.u64 	%rd190, image_Mask;
	cvta.global.u64 	%rd185, %rd190;
	// inline asm
	call (%rd184), _rt_buffer_get_64, (%rd185, %r98, %r98, %rd186, %rd187, %rd24, %rd24);
	// inline asm
	mov.f32 	%f873, 0f00000000;
	cvt.rzi.u32.f32	%r340, %f873;
	cvt.u16.u32	%rs85, %r340;
	mov.u16 	%rs86, 0;
	st.v2.u8 	[%rd184], {%rs85, %rs86};
	ld.global.u32 	%r410, [imageEnabled];

BB0_143:
	cvt.u64.u32	%rd17, %r4;
	cvt.u64.u32	%rd18, %r5;
	and.b32  	%r341, %r410, 4;
	setp.eq.s32	%p142, %r341, 0;
	@%p142 bra 	BB0_147;

	ld.global.u32 	%r342, [additive];
	setp.eq.s32	%p143, %r342, 0;
	@%p143 bra 	BB0_146;

	mov.u64 	%rd203, image_HDR;
	cvta.global.u64 	%rd192, %rd203;
	mov.u32 	%r346, 8;
	// inline asm
	call (%rd191), _rt_buffer_get_64, (%rd192, %r98, %r346, %rd17, %rd18, %rd24, %rd24);
	// inline asm
	ld.v4.u16 	{%rs93, %rs94, %rs95, %rs96}, [%rd191];
	// inline asm
	{  cvt.f32.f16 %f874, %rs93;}

	// inline asm
	// inline asm
	{  cvt.f32.f16 %f875, %rs94;}

	// inline asm
	// inline asm
	{  cvt.f32.f16 %f876, %rs95;}

	// inline asm
	// inline asm
	call (%rd197), _rt_buffer_get_64, (%rd192, %r98, %r346, %rd17, %rd18, %rd24, %rd24);
	// inline asm
	add.f32 	%f877, %f874, 0f00000000;
	add.f32 	%f878, %f875, 0f00000000;
	add.f32 	%f879, %f876, 0f00000000;
	// inline asm
	{  cvt.rn.f16.f32 %rs92, %f879;}

	// inline asm
	// inline asm
	{  cvt.rn.f16.f32 %rs91, %f878;}

	// inline asm
	// inline asm
	{  cvt.rn.f16.f32 %rs90, %f877;}

	// inline asm
	mov.u16 	%rs97, 0;
	st.v4.u16 	[%rd197], {%rs90, %rs91, %rs92, %rs97};
	bra.uni 	BB0_147;

BB0_3:
	ld.global.v2.u32 	{%r118, %r119}, [pixelID];
	cvt.u64.u32	%rd28, %r118;
	cvt.u64.u32	%rd29, %r119;
	mov.u64 	%rd38, uvpos;
	cvta.global.u64 	%rd27, %rd38;
	mov.u32 	%r114, 12;
	// inline asm
	call (%rd26), _rt_buffer_get_64, (%rd27, %r98, %r114, %rd28, %rd29, %rd24, %rd24);
	// inline asm
	ld.f32 	%f9, [%rd26+8];
	ld.f32 	%f8, [%rd26+4];
	ld.f32 	%f7, [%rd26];
	mul.f32 	%f237, %f7, 0f3456BF95;
	mul.f32 	%f238, %f8, 0f3456BF95;
	mul.f32 	%f239, %f9, 0f3456BF95;
	abs.f32 	%f240, %f916;
	div.rn.f32 	%f241, %f237, %f240;
	abs.f32 	%f242, %f917;
	div.rn.f32 	%f243, %f238, %f242;
	abs.f32 	%f244, %f918;
	div.rn.f32 	%f245, %f239, %f244;
	abs.f32 	%f246, %f241;
	abs.f32 	%f247, %f243;
	abs.f32 	%f248, %f245;
	mov.f32 	%f249, 0f38D1B717;
	max.f32 	%f250, %f246, %f249;
	max.f32 	%f251, %f247, %f249;
	max.f32 	%f252, %f248, %f249;
	fma.rn.f32 	%f10, %f916, %f250, %f7;
	fma.rn.f32 	%f11, %f917, %f251, %f8;
	fma.rn.f32 	%f12, %f918, %f252, %f9;
	ld.global.v2.u32 	{%r122, %r123}, [pixelID];
	cvt.u64.u32	%rd34, %r122;
	cvt.u64.u32	%rd35, %r123;
	mov.u64 	%rd39, rnd_seeds;
	cvta.global.u64 	%rd33, %rd39;
	// inline asm
	call (%rd32), _rt_buffer_get_64, (%rd33, %r98, %r99, %rd34, %rd35, %rd24, %rd24);
	// inline asm
	ld.u32 	%r126, [%rd32];
	mad.lo.s32 	%r6, %r126, 1664525, 1013904223;
	ld.global.u32 	%r127, [lightMeshBufferSize];
	setp.eq.s32	%p14, %r127, 0;
	mov.pred 	%p13, 0;
	mov.f32 	%f18, 0f00000000;
	mov.u32 	%r382, 0;
	@%p14 bra 	BB0_4;

	ld.global.f32 	%f13, [lightPointSize];
	mul.f32 	%f14, %f10, 0f3456BF95;
	mul.f32 	%f15, %f11, 0f3456BF95;
	mul.f32 	%f16, %f12, 0f3456BF95;
	and.b32  	%r129, %r6, 16777215;
	cvt.rn.f32.u32	%f259, %r129;
	mul.f32 	%f260, %f259, 0fB3800000;
	fma.rn.f32 	%f17, %f260, 0f3F333333, 0f3F800000;
	mov.f32 	%f18, 0f00000000;
	mov.u32 	%r382, 0;
	abs.f32 	%f382, %f15;
	abs.f32 	%f383, %f14;
	max.f32 	%f384, %f383, %f382;
	abs.f32 	%f385, %f16;
	max.f32 	%f386, %f384, %f385;
	mov.f32 	%f19, %f18;
	mov.f32 	%f20, %f18;
	mov.f32 	%f21, %f18;
	mov.f32 	%f22, %f18;
	mov.f32 	%f23, %f18;

BB0_6:
	shl.b32 	%r132, %r382, 1;
	cvt.s64.s32	%rd42, %r132;
	mov.u64 	%rd46, lightMeshBuffer;
	cvta.global.u64 	%rd41, %rd46;
	mov.u32 	%r130, 1;
	// inline asm
	call (%rd40), _rt_buffer_get_64, (%rd41, %r130, %r114, %rd42, %rd24, %rd24, %rd24);
	// inline asm
	ld.f32 	%f261, [%rd40];
	sub.f32 	%f262, %f261, %f7;
	ld.f32 	%f263, [%rd40+4];
	sub.f32 	%f264, %f263, %f8;
	ld.f32 	%f265, [%rd40+8];
	sub.f32 	%f266, %f265, %f9;
	mul.f32 	%f267, %f264, %f264;
	fma.rn.f32 	%f268, %f262, %f262, %f267;
	fma.rn.f32 	%f269, %f266, %f266, %f268;
	sqrt.rn.f32 	%f24, %f269;
	rcp.rn.f32 	%f270, %f24;
	mul.f32 	%f25, %f262, %f270;
	mul.f32 	%f26, %f264, %f270;
	mul.f32 	%f27, %f266, %f270;
	mul.f32 	%f271, %f24, %f24;
	mul.f32 	%f272, %f271, 0f40C90FDB;
	div.rn.f32 	%f273, %f13, %f272;
	add.f32 	%f28, %f273, %f273;
	setp.gt.f32	%p16, %f28, %f17;
	setp.ne.s32	%p17, %r1, 0;
	and.pred  	%p18, %p17, %p16;
	mov.pred 	%p148, -1;
	@%p18 bra 	BB0_22;

	ld.global.f32 	%f276, [lightInvCutoff];
	mul.f32 	%f29, %f24, %f276;
	mov.f32 	%f280, 0f40800000;
	abs.f32 	%f31, %f29;
	setp.lt.f32	%p19, %f31, 0f00800000;
	mul.f32 	%f282, %f31, 0f4B800000;
	selp.f32	%f283, 0fC3170000, 0fC2FE0000, %p19;
	selp.f32	%f284, %f282, %f31, %p19;
	mov.b32 	 %r133, %f284;
	and.b32  	%r134, %r133, 8388607;
	or.b32  	%r135, %r134, 1065353216;
	mov.b32 	 %f285, %r135;
	shr.u32 	%r136, %r133, 23;
	cvt.rn.f32.u32	%f286, %r136;
	add.f32 	%f287, %f283, %f286;
	setp.gt.f32	%p20, %f285, 0f3FB504F3;
	mul.f32 	%f288, %f285, 0f3F000000;
	add.f32 	%f289, %f287, 0f3F800000;
	selp.f32	%f290, %f288, %f285, %p20;
	selp.f32	%f291, %f289, %f287, %p20;
	add.f32 	%f292, %f290, 0fBF800000;
	add.f32 	%f275, %f290, 0f3F800000;
	// inline asm
	rcp.approx.ftz.f32 %f274,%f275;
	// inline asm
	add.f32 	%f293, %f292, %f292;
	mul.f32 	%f294, %f274, %f293;
	mul.f32 	%f295, %f294, %f294;
	mov.f32 	%f296, 0f3C4CAF63;
	mov.f32 	%f297, 0f3B18F0FE;
	fma.rn.f32 	%f298, %f297, %f295, %f296;
	mov.f32 	%f299, 0f3DAAAABD;
	fma.rn.f32 	%f300, %f298, %f295, %f299;
	mul.rn.f32 	%f301, %f300, %f295;
	mul.rn.f32 	%f302, %f301, %f294;
	sub.f32 	%f303, %f292, %f294;
	neg.f32 	%f304, %f294;
	add.f32 	%f305, %f303, %f303;
	fma.rn.f32 	%f306, %f304, %f292, %f305;
	mul.rn.f32 	%f307, %f274, %f306;
	add.f32 	%f308, %f302, %f294;
	sub.f32 	%f309, %f294, %f308;
	add.f32 	%f310, %f302, %f309;
	add.f32 	%f311, %f307, %f310;
	add.f32 	%f312, %f308, %f311;
	sub.f32 	%f313, %f308, %f312;
	add.f32 	%f314, %f311, %f313;
	mov.f32 	%f315, 0f3F317200;
	mul.rn.f32 	%f316, %f291, %f315;
	mov.f32 	%f317, 0f35BFBE8E;
	mul.rn.f32 	%f318, %f291, %f317;
	add.f32 	%f319, %f316, %f312;
	sub.f32 	%f320, %f316, %f319;
	add.f32 	%f321, %f312, %f320;
	add.f32 	%f322, %f314, %f321;
	add.f32 	%f323, %f318, %f322;
	add.f32 	%f324, %f319, %f323;
	sub.f32 	%f325, %f319, %f324;
	add.f32 	%f326, %f323, %f325;
	mul.rn.f32 	%f327, %f280, %f324;
	neg.f32 	%f328, %f327;
	fma.rn.f32 	%f329, %f280, %f324, %f328;
	fma.rn.f32 	%f330, %f280, %f326, %f329;
	mov.f32 	%f331, 0f00000000;
	fma.rn.f32 	%f332, %f331, %f324, %f330;
	add.rn.f32 	%f333, %f327, %f332;
	neg.f32 	%f334, %f333;
	add.rn.f32 	%f335, %f327, %f334;
	add.rn.f32 	%f336, %f335, %f332;
	mov.b32 	 %r137, %f333;
	setp.eq.s32	%p21, %r137, 1118925336;
	add.s32 	%r138, %r137, -1;
	mov.b32 	 %f337, %r138;
	add.f32 	%f338, %f336, 0f37000000;
	selp.f32	%f339, %f337, %f333, %p21;
	selp.f32	%f32, %f338, %f336, %p21;
	mul.f32 	%f340, %f339, 0f3FB8AA3B;
	cvt.rzi.f32.f32	%f341, %f340;
	mov.f32 	%f342, 0fBF317200;
	fma.rn.f32 	%f343, %f341, %f342, %f339;
	mov.f32 	%f344, 0fB5BFBE8E;
	fma.rn.f32 	%f345, %f341, %f344, %f343;
	mul.f32 	%f346, %f345, 0f3FB8AA3B;
	ex2.approx.ftz.f32 	%f347, %f346;
	add.f32 	%f348, %f341, 0f00000000;
	ex2.approx.f32 	%f349, %f348;
	mul.f32 	%f350, %f347, %f349;
	setp.lt.f32	%p22, %f339, 0fC2D20000;
	selp.f32	%f351, 0f00000000, %f350, %p22;
	setp.gt.f32	%p23, %f339, 0f42D20000;
	selp.f32	%f925, 0f7F800000, %f351, %p23;
	setp.eq.f32	%p24, %f925, 0f7F800000;
	@%p24 bra 	BB0_9;

	fma.rn.f32 	%f925, %f925, %f32, %f925;

BB0_9:
	mov.f32 	%f914, 0f40000000;
	cvt.rzi.f32.f32	%f913, %f914;
	add.f32 	%f912, %f913, %f913;
	mov.f32 	%f911, 0f40800000;
	sub.f32 	%f910, %f911, %f912;
	abs.f32 	%f909, %f910;
	setp.lt.f32	%p25, %f29, 0f00000000;
	setp.eq.f32	%p26, %f909, 0f3F800000;
	and.pred  	%p1, %p25, %p26;
	mov.b32 	 %r139, %f925;
	xor.b32  	%r140, %r139, -2147483648;
	mov.b32 	 %f352, %r140;
	selp.f32	%f927, %f352, %f925, %p1;
	setp.eq.f32	%p27, %f29, 0f00000000;
	@%p27 bra 	BB0_12;
	bra.uni 	BB0_10;

BB0_12:
	add.f32 	%f355, %f29, %f29;
	selp.f32	%f927, %f355, 0f00000000, %p26;
	bra.uni 	BB0_13;

BB0_10:
	setp.geu.f32	%p28, %f29, 0f00000000;
	@%p28 bra 	BB0_13;

	mov.f32 	%f915, 0f40800000;
	cvt.rzi.f32.f32	%f354, %f915;
	setp.neu.f32	%p29, %f354, 0f40800000;
	selp.f32	%f927, 0f7FFFFFFF, %f927, %p29;

BB0_13:
	add.f32 	%f356, %f31, 0f40800000;
	mov.b32 	 %r141, %f356;
	setp.lt.s32	%p31, %r141, 2139095040;
	@%p31 bra 	BB0_18;

	setp.gtu.f32	%p32, %f31, 0f7F800000;
	@%p32 bra 	BB0_17;
	bra.uni 	BB0_15;

BB0_17:
	add.f32 	%f927, %f29, 0f40800000;
	bra.uni 	BB0_18;

BB0_15:
	setp.neu.f32	%p33, %f31, 0f7F800000;
	@%p33 bra 	BB0_18;

	selp.f32	%f927, 0fFF800000, 0f7F800000, %p1;

BB0_18:
	mov.u32 	%r378, 1;
	mov.u64 	%rd292, lightMeshBuffer;
	cvta.global.u64 	%rd291, %rd292;
	shl.b32 	%r377, %r382, 1;
	mov.f32 	%f357, 0f3F800000;
	sub.f32 	%f358, %f357, %f927;
	setp.eq.f32	%p34, %f29, 0f3F800000;
	selp.f32	%f359, 0f00000000, %f358, %p34;
	cvt.sat.f32.f32	%f360, %f359;
	mul.f32 	%f361, %f28, %f360;
	add.s32 	%r145, %r377, 1;
	cvt.s64.s32	%rd49, %r145;
	// inline asm
	call (%rd47), _rt_buffer_get_64, (%rd291, %r378, %r114, %rd49, %rd24, %rd24, %rd24);
	// inline asm
	ld.f32 	%f362, [%rd47];
	mul.f32 	%f363, %f25, %f362;
	ld.f32 	%f364, [%rd47+4];
	mul.f32 	%f365, %f26, %f364;
	neg.f32 	%f366, %f365;
	sub.f32 	%f367, %f366, %f363;
	ld.f32 	%f368, [%rd47+8];
	mul.f32 	%f369, %f27, %f368;
	sub.f32 	%f370, %f367, %f369;
	cvt.sat.f32.f32	%f371, %f370;
	mul.f32 	%f43, %f361, %f371;
	mul.f32 	%f372, %f917, %f26;
	fma.rn.f32 	%f373, %f916, %f25, %f372;
	fma.rn.f32 	%f44, %f918, %f27, %f373;
	setp.leu.f32	%p35, %f43, 0f3727C5AC;
	@%p35 bra 	BB0_20;

	mov.u32 	%r380, 1;
	add.u64 	%rd54, %SP, 4;
	cvta.to.local.u64 	%rd55, %rd54;
	max.f32 	%f380, %f386, %f249;
	sub.f32 	%f381, %f24, %f380;
	mov.u32 	%r149, 1065353216;
	st.local.u32 	[%rd55], %r149;
	ld.global.u32 	%r146, [root];
	// inline asm
	call _rt_trace_64, (%r146, %f10, %f11, %f12, %f25, %f26, %f27, %r380, %f380, %f381, %rd54, %r99);
	// inline asm
	ld.local.f32 	%f388, [%rd55];
	mul.f32 	%f389, %f43, %f388;
	cvt.sat.f32.f32	%f390, %f44;
	fma.rn.f32 	%f23, %f390, %f389, %f23;
	fma.rn.f32 	%f21, %f25, %f389, %f21;
	fma.rn.f32 	%f20, %f26, %f389, %f20;
	fma.rn.f32 	%f19, %f27, %f389, %f19;
	add.f32 	%f22, %f22, %f389;
	add.f32 	%f18, %f18, %f388;

BB0_20:
	ld.global.u32 	%r150, [lightMeshBufferSize];
	add.s32 	%r382, %r382, 1;
	setp.lt.u32	%p37, %r382, %r150;
	@%p37 bra 	BB0_6;
	bra.uni 	BB0_21;

BB0_4:
	mov.f32 	%f19, %f18;
	mov.f32 	%f20, %f18;
	mov.f32 	%f21, %f18;
	mov.f32 	%f22, %f18;
	mov.f32 	%f23, %f18;

BB0_21:
	mov.pred 	%p148, %p13;

BB0_22:
	cvt.rn.f32.s32	%f391, %r382;
	mov.f32 	%f392, 0f3F800000;
	max.f32 	%f393, %f391, %f392;
	div.rn.f32 	%f970, %f23, %f393;
	div.rn.f32 	%f975, %f18, %f393;
	div.rn.f32 	%f971, %f22, %f393;
	div.rn.f32 	%f972, %f21, %f393;
	div.rn.f32 	%f973, %f20, %f393;
	div.rn.f32 	%f974, %f19, %f393;
	@!%p148 bra 	BB0_75;
	bra.uni 	BB0_23;

BB0_23:
	mov.f32 	%f946, 0f00000000;
	setp.lt.s32	%p38, %r1, 1;
	mov.f32 	%f947, %f946;
	mov.f32 	%f948, %f946;
	mov.f32 	%f949, %f946;
	mov.f32 	%f950, %f946;
	mov.f32 	%f951, %f946;
	@%p38 bra 	BB0_74;

	mad.lo.s32 	%r386, %r126, 1664525, 1013904223;
	cvt.rn.f32.s32	%f406, %r2;
	rcp.rn.f32 	%f69, %f406;
	add.u64 	%rd56, %SP, 8;
	cvta.to.local.u64 	%rd2, %rd56;
	mul.f32 	%f70, %f10, 0f3456BF95;
	mul.f32 	%f71, %f11, 0f3456BF95;
	mul.f32 	%f72, %f12, 0f3456BF95;
	add.u64 	%rd57, %SP, 0;
	cvta.to.local.u64 	%rd3, %rd57;
	mov.f32 	%f946, 0f00000000;
	mov.u32 	%r151, 0;
	abs.f32 	%f407, %f71;
	abs.f32 	%f408, %f70;
	max.f32 	%f409, %f408, %f407;
	abs.f32 	%f410, %f72;
	max.f32 	%f411, %f409, %f410;
	mov.u32 	%r383, %r151;
	mov.f32 	%f947, %f946;
	mov.f32 	%f948, %f946;
	mov.f32 	%f949, %f946;
	mov.f32 	%f950, %f946;
	mov.f32 	%f951, %f946;

BB0_25:
	cvt.rn.f32.s32	%f79, %r383;
	max.f32 	%f80, %f411, %f249;
	mov.u32 	%r385, %r151;

BB0_26:
	mad.lo.s32 	%r153, %r386, 1664525, 1013904223;
	and.b32  	%r154, %r153, 16777215;
	cvt.rn.f32.u32	%f413, %r154;
	fma.rn.f32 	%f414, %f413, 0f33800000, %f79;
	mul.f32 	%f415, %f69, %f414;
	mad.lo.s32 	%r386, %r153, 1664525, 1013904223;
	and.b32  	%r155, %r386, 16777215;
	cvt.rn.f32.u32	%f416, %r155;
	cvt.rn.f32.s32	%f417, %r385;
	fma.rn.f32 	%f418, %f416, 0f33800000, %f417;
	mul.f32 	%f419, %f69, %f418;
	fma.rn.f32 	%f87, %f415, 0fC0000000, 0f3F800000;
	mul.f32 	%f420, %f87, %f87;
	sub.f32 	%f422, %f392, %f420;
	mov.f32 	%f423, 0f00000000;
	max.f32 	%f424, %f423, %f422;
	sqrt.rn.f32 	%f88, %f424;
	mul.f32 	%f958, %f419, 0f40C90FDB;
	abs.f32 	%f90, %f958;
	setp.neu.f32	%p39, %f90, 0f7F800000;
	mov.f32 	%f952, %f958;
	@%p39 bra 	BB0_28;

	mul.rn.f32 	%f952, %f958, %f423;

BB0_28:
	mul.f32 	%f426, %f952, 0f3F22F983;
	cvt.rni.s32.f32	%r396, %f426;
	cvt.rn.f32.s32	%f427, %r396;
	neg.f32 	%f428, %f427;
	mov.f32 	%f429, 0f3FC90FDA;
	fma.rn.f32 	%f430, %f428, %f429, %f952;
	mov.f32 	%f431, 0f33A22168;
	fma.rn.f32 	%f432, %f428, %f431, %f430;
	mov.f32 	%f433, 0f27C234C5;
	fma.rn.f32 	%f953, %f428, %f433, %f432;
	abs.f32 	%f434, %f952;
	setp.leu.f32	%p40, %f434, 0f47CE4780;
	@%p40 bra 	BB0_39;

	mov.b32 	 %r16, %f952;
	shr.u32 	%r17, %r16, 23;
	shl.b32 	%r158, %r16, 8;
	or.b32  	%r18, %r158, -2147483648;
	mov.u32 	%r388, 0;
	mov.u64 	%rd293, __cudart_i2opi_f;
	mov.u32 	%r387, -6;
	mov.u64 	%rd294, %rd2;

BB0_30:
	.pragma "nounroll";
	ld.const.u32 	%r161, [%rd293];
	// inline asm
	{
	mad.lo.cc.u32   %r159, %r161, %r18, %r388;
	madc.hi.u32     %r388, %r161, %r18,  0;
	}
	// inline asm
	st.local.u32 	[%rd294], %r159;
	add.s64 	%rd294, %rd294, 4;
	add.s64 	%rd293, %rd293, 4;
	add.s32 	%r387, %r387, 1;
	setp.ne.s32	%p41, %r387, 0;
	@%p41 bra 	BB0_30;

	and.b32  	%r164, %r17, 255;
	add.s32 	%r165, %r164, -128;
	shr.u32 	%r166, %r165, 5;
	and.b32  	%r23, %r16, -2147483648;
	st.local.u32 	[%rd2+24], %r388;
	mov.u32 	%r167, 6;
	sub.s32 	%r168, %r167, %r166;
	mul.wide.s32 	%rd59, %r168, 4;
	add.s64 	%rd8, %rd2, %rd59;
	ld.local.u32 	%r389, [%rd8];
	ld.local.u32 	%r390, [%rd8+-4];
	and.b32  	%r26, %r17, 31;
	setp.eq.s32	%p42, %r26, 0;
	@%p42 bra 	BB0_33;

	mov.u32 	%r169, 32;
	sub.s32 	%r170, %r169, %r26;
	shr.u32 	%r171, %r390, %r170;
	shl.b32 	%r172, %r389, %r26;
	add.s32 	%r389, %r171, %r172;
	ld.local.u32 	%r173, [%rd8+-8];
	shr.u32 	%r174, %r173, %r170;
	shl.b32 	%r175, %r390, %r26;
	add.s32 	%r390, %r174, %r175;

BB0_33:
	shr.u32 	%r176, %r390, 30;
	shl.b32 	%r177, %r389, 2;
	add.s32 	%r391, %r176, %r177;
	shl.b32 	%r32, %r390, 2;
	shr.u32 	%r178, %r391, 31;
	shr.u32 	%r179, %r389, 30;
	add.s32 	%r33, %r178, %r179;
	setp.eq.s32	%p43, %r178, 0;
	@%p43 bra 	BB0_34;
	bra.uni 	BB0_35;

BB0_34:
	mov.u32 	%r392, %r23;
	mov.u32 	%r393, %r32;
	bra.uni 	BB0_36;

BB0_35:
	not.b32 	%r180, %r391;
	neg.s32 	%r393, %r32;
	setp.eq.s32	%p44, %r32, 0;
	selp.u32	%r181, 1, 0, %p44;
	add.s32 	%r391, %r181, %r180;
	xor.b32  	%r392, %r23, -2147483648;

BB0_36:
	clz.b32 	%r395, %r391;
	setp.eq.s32	%p45, %r395, 0;
	shl.b32 	%r182, %r391, %r395;
	mov.u32 	%r183, 32;
	sub.s32 	%r184, %r183, %r395;
	shr.u32 	%r185, %r393, %r184;
	add.s32 	%r186, %r185, %r182;
	selp.b32	%r41, %r391, %r186, %p45;
	mov.u32 	%r187, -921707870;
	mul.hi.u32 	%r394, %r41, %r187;
	setp.eq.s32	%p46, %r23, 0;
	neg.s32 	%r188, %r33;
	selp.b32	%r396, %r33, %r188, %p46;
	setp.lt.s32	%p47, %r394, 1;
	@%p47 bra 	BB0_38;

	mul.lo.s32 	%r189, %r41, -921707870;
	shr.u32 	%r190, %r189, 31;
	shl.b32 	%r191, %r394, 1;
	add.s32 	%r394, %r190, %r191;
	add.s32 	%r395, %r395, 1;

BB0_38:
	mov.u32 	%r192, 126;
	sub.s32 	%r193, %r192, %r395;
	shl.b32 	%r194, %r193, 23;
	add.s32 	%r195, %r394, 1;
	shr.u32 	%r196, %r195, 7;
	add.s32 	%r197, %r196, 1;
	shr.u32 	%r198, %r197, 1;
	add.s32 	%r199, %r198, %r194;
	or.b32  	%r200, %r199, %r392;
	mov.b32 	 %f953, %r200;

BB0_39:
	mul.rn.f32 	%f96, %f953, %f953;
	add.s32 	%r49, %r396, 1;
	and.b32  	%r50, %r49, 1;
	setp.eq.s32	%p48, %r50, 0;
	@%p48 bra 	BB0_41;
	bra.uni 	BB0_40;

BB0_41:
	mov.f32 	%f437, 0f3C08839E;
	mov.f32 	%f438, 0fB94CA1F9;
	fma.rn.f32 	%f954, %f438, %f96, %f437;
	bra.uni 	BB0_42;

BB0_40:
	mov.f32 	%f435, 0fBAB6061A;
	mov.f32 	%f436, 0f37CCF5CE;
	fma.rn.f32 	%f954, %f436, %f96, %f435;

BB0_42:
	@%p48 bra 	BB0_44;
	bra.uni 	BB0_43;

BB0_44:
	mov.f32 	%f442, 0fBE2AAAA3;
	fma.rn.f32 	%f443, %f954, %f96, %f442;
	fma.rn.f32 	%f955, %f443, %f96, %f423;
	bra.uni 	BB0_45;

BB0_43:
	mov.f32 	%f439, 0f3D2AAAA5;
	fma.rn.f32 	%f440, %f954, %f96, %f439;
	mov.f32 	%f441, 0fBF000000;
	fma.rn.f32 	%f955, %f440, %f96, %f441;

BB0_45:
	fma.rn.f32 	%f956, %f955, %f953, %f953;
	@%p48 bra 	BB0_47;

	fma.rn.f32 	%f956, %f955, %f96, %f392;

BB0_47:
	and.b32  	%r201, %r49, 2;
	setp.eq.s32	%p51, %r201, 0;
	@%p51 bra 	BB0_49;

	mov.f32 	%f447, 0fBF800000;
	fma.rn.f32 	%f956, %f956, %f447, %f423;

BB0_49:
	@%p39 bra 	BB0_51;

	mul.rn.f32 	%f958, %f958, %f423;

BB0_51:
	mul.f32 	%f449, %f958, 0f3F22F983;
	cvt.rni.s32.f32	%r406, %f449;
	cvt.rn.f32.s32	%f450, %r406;
	neg.f32 	%f451, %f450;
	fma.rn.f32 	%f453, %f451, %f429, %f958;
	fma.rn.f32 	%f455, %f451, %f431, %f453;
	fma.rn.f32 	%f959, %f451, %f433, %f455;
	abs.f32 	%f457, %f958;
	setp.leu.f32	%p53, %f457, 0f47CE4780;
	@%p53 bra 	BB0_62;

	mov.b32 	 %r52, %f958;
	shr.u32 	%r53, %r52, 23;
	shl.b32 	%r204, %r52, 8;
	or.b32  	%r54, %r204, -2147483648;
	mov.u32 	%r398, 0;
	mov.u64 	%rd295, __cudart_i2opi_f;
	mov.u32 	%r397, -6;
	mov.u64 	%rd296, %rd2;

BB0_53:
	.pragma "nounroll";
	ld.const.u32 	%r207, [%rd295];
	// inline asm
	{
	mad.lo.cc.u32   %r205, %r207, %r54, %r398;
	madc.hi.u32     %r398, %r207, %r54,  0;
	}
	// inline asm
	st.local.u32 	[%rd296], %r205;
	add.s64 	%rd296, %rd296, 4;
	add.s64 	%rd295, %rd295, 4;
	add.s32 	%r397, %r397, 1;
	setp.ne.s32	%p54, %r397, 0;
	@%p54 bra 	BB0_53;

	and.b32  	%r210, %r53, 255;
	add.s32 	%r211, %r210, -128;
	shr.u32 	%r212, %r211, 5;
	and.b32  	%r59, %r52, -2147483648;
	st.local.u32 	[%rd2+24], %r398;
	mov.u32 	%r213, 6;
	sub.s32 	%r214, %r213, %r212;
	mul.wide.s32 	%rd61, %r214, 4;
	add.s64 	%rd13, %rd2, %rd61;
	ld.local.u32 	%r399, [%rd13];
	ld.local.u32 	%r400, [%rd13+-4];
	and.b32  	%r62, %r53, 31;
	setp.eq.s32	%p55, %r62, 0;
	@%p55 bra 	BB0_56;

	mov.u32 	%r215, 32;
	sub.s32 	%r216, %r215, %r62;
	shr.u32 	%r217, %r400, %r216;
	shl.b32 	%r218, %r399, %r62;
	add.s32 	%r399, %r217, %r218;
	ld.local.u32 	%r219, [%rd13+-8];
	shr.u32 	%r220, %r219, %r216;
	shl.b32 	%r221, %r400, %r62;
	add.s32 	%r400, %r220, %r221;

BB0_56:
	shr.u32 	%r222, %r400, 30;
	shl.b32 	%r223, %r399, 2;
	add.s32 	%r401, %r222, %r223;
	shl.b32 	%r68, %r400, 2;
	shr.u32 	%r224, %r401, 31;
	shr.u32 	%r225, %r399, 30;
	add.s32 	%r69, %r224, %r225;
	setp.eq.s32	%p56, %r224, 0;
	@%p56 bra 	BB0_57;
	bra.uni 	BB0_58;

BB0_57:
	mov.u32 	%r402, %r59;
	mov.u32 	%r403, %r68;
	bra.uni 	BB0_59;

BB0_58:
	not.b32 	%r226, %r401;
	neg.s32 	%r403, %r68;
	setp.eq.s32	%p57, %r68, 0;
	selp.u32	%r227, 1, 0, %p57;
	add.s32 	%r401, %r227, %r226;
	xor.b32  	%r402, %r59, -2147483648;

BB0_59:
	clz.b32 	%r405, %r401;
	setp.eq.s32	%p58, %r405, 0;
	shl.b32 	%r228, %r401, %r405;
	mov.u32 	%r229, 32;
	sub.s32 	%r230, %r229, %r405;
	shr.u32 	%r231, %r403, %r230;
	add.s32 	%r232, %r231, %r228;
	selp.b32	%r77, %r401, %r232, %p58;
	mov.u32 	%r233, -921707870;
	mul.hi.u32 	%r404, %r77, %r233;
	setp.eq.s32	%p59, %r59, 0;
	neg.s32 	%r234, %r69;
	selp.b32	%r406, %r69, %r234, %p59;
	setp.lt.s32	%p60, %r404, 1;
	@%p60 bra 	BB0_61;

	mul.lo.s32 	%r235, %r77, -921707870;
	shr.u32 	%r236, %r235, 31;
	shl.b32 	%r237, %r404, 1;
	add.s32 	%r404, %r236, %r237;
	add.s32 	%r405, %r405, 1;

BB0_61:
	mov.u32 	%r238, 126;
	sub.s32 	%r239, %r238, %r405;
	shl.b32 	%r240, %r239, 23;
	add.s32 	%r241, %r404, 1;
	shr.u32 	%r242, %r241, 7;
	add.s32 	%r243, %r242, 1;
	shr.u32 	%r244, %r243, 1;
	add.s32 	%r245, %r244, %r240;
	or.b32  	%r246, %r245, %r402;
	mov.b32 	 %f959, %r246;

BB0_62:
	mul.rn.f32 	%f113, %f959, %f959;
	and.b32  	%r85, %r406, 1;
	setp.eq.s32	%p61, %r85, 0;
	@%p61 bra 	BB0_64;
	bra.uni 	BB0_63;

BB0_64:
	mov.f32 	%f460, 0f3C08839E;
	mov.f32 	%f461, 0fB94CA1F9;
	fma.rn.f32 	%f960, %f461, %f113, %f460;
	bra.uni 	BB0_65;

BB0_63:
	mov.f32 	%f458, 0fBAB6061A;
	mov.f32 	%f459, 0f37CCF5CE;
	fma.rn.f32 	%f960, %f459, %f113, %f458;

BB0_65:
	@%p61 bra 	BB0_67;
	bra.uni 	BB0_66;

BB0_67:
	mov.f32 	%f465, 0fBE2AAAA3;
	fma.rn.f32 	%f466, %f960, %f113, %f465;
	fma.rn.f32 	%f961, %f466, %f113, %f423;
	bra.uni 	BB0_68;

BB0_66:
	mov.f32 	%f462, 0f3D2AAAA5;
	fma.rn.f32 	%f463, %f960, %f113, %f462;
	mov.f32 	%f464, 0fBF000000;
	fma.rn.f32 	%f961, %f463, %f113, %f464;

BB0_68:
	fma.rn.f32 	%f962, %f961, %f959, %f959;
	@%p61 bra 	BB0_70;

	fma.rn.f32 	%f962, %f961, %f113, %f392;

BB0_70:
	and.b32  	%r247, %r406, 2;
	setp.eq.s32	%p64, %r247, 0;
	@%p64 bra 	BB0_72;

	mov.f32 	%f470, 0fBF800000;
	fma.rn.f32 	%f962, %f962, %f470, %f423;

BB0_72:
	mul.f32 	%f474, %f88, %f956;
	mov.u32 	%r249, 0;
	st.local.u32 	[%rd3], %r249;
	ld.global.u32 	%r248, [root];
	mul.f32 	%f475, %f88, %f962;
	mov.f32 	%f478, 0f6C4ECB8F;
	// inline asm
	call _rt_trace_64, (%r248, %f10, %f11, %f12, %f474, %f475, %f87, %r249, %f80, %f478, %rd57, %r99);
	// inline asm
	ld.local.f32 	%f479, [%rd3];
	setp.lt.f32	%p65, %f479, 0f00000000;
	selp.f32	%f480, 0f00000000, %f479, %p65;
	selp.f32	%f481, 0f00000000, 0f3F800000, %p65;
	add.f32 	%f946, %f946, %f481;
	fma.rn.f32 	%f949, %f474, %f480, %f949;
	fma.rn.f32 	%f948, %f475, %f480, %f948;
	fma.rn.f32 	%f947, %f87, %f480, %f947;
	add.f32 	%f950, %f950, %f480;
	mul.f32 	%f482, %f917, %f475;
	fma.rn.f32 	%f483, %f916, %f474, %f482;
	fma.rn.f32 	%f484, %f918, %f87, %f483;
	cvt.sat.f32.f32	%f485, %f484;
	fma.rn.f32 	%f951, %f480, %f485, %f951;
	add.s32 	%r385, %r385, 1;
	setp.lt.s32	%p66, %r385, %r2;
	@%p66 bra 	BB0_26;

	add.s32 	%r383, %r383, 1;
	setp.lt.s32	%p67, %r383, %r2;
	@%p67 bra 	BB0_25;

BB0_74:
	mul.lo.s32 	%r251, %r2, %r2;
	cvt.rn.f32.s32	%f486, %r251;
	div.rn.f32 	%f487, %f951, %f486;
	div.rn.f32 	%f975, %f946, %f486;
	div.rn.f32 	%f971, %f950, %f486;
	div.rn.f32 	%f972, %f949, %f486;
	div.rn.f32 	%f973, %f948, %f486;
	div.rn.f32 	%f974, %f947, %f486;
	add.f32 	%f970, %f487, %f487;

BB0_75:
	ld.global.u32 	%r408, [imageEnabled];
	and.b32  	%r252, %r408, 8;
	setp.eq.s32	%p68, %r252, 0;
	@%p68 bra 	BB0_88;

	cvt.u64.u32	%rd65, %r4;
	cvt.u64.u32	%rd66, %r5;
	mov.u64 	%rd69, image_Mask;
	cvta.global.u64 	%rd64, %rd69;
	// inline asm
	call (%rd63), _rt_buffer_get_64, (%rd64, %r98, %r98, %rd65, %rd66, %rd24, %rd24);
	// inline asm
	mov.f32 	%f490, 0f3E68BA2E;
	cvt.rzi.f32.f32	%f491, %f490;
	fma.rn.f32 	%f492, %f491, 0fC0000000, 0f3EE8BA2E;
	abs.f32 	%f149, %f492;
	abs.f32 	%f150, %f975;
	setp.lt.f32	%p69, %f150, 0f00800000;
	mul.f32 	%f493, %f150, 0f4B800000;
	selp.f32	%f494, 0fC3170000, 0fC2FE0000, %p69;
	selp.f32	%f495, %f493, %f150, %p69;
	mov.b32 	 %r255, %f495;
	and.b32  	%r256, %r255, 8388607;
	or.b32  	%r257, %r256, 1065353216;
	mov.b32 	 %f496, %r257;
	shr.u32 	%r258, %r255, 23;
	cvt.rn.f32.u32	%f497, %r258;
	add.f32 	%f498, %f494, %f497;
	setp.gt.f32	%p70, %f496, 0f3FB504F3;
	mul.f32 	%f499, %f496, 0f3F000000;
	add.f32 	%f500, %f498, 0f3F800000;
	selp.f32	%f501, %f499, %f496, %p70;
	selp.f32	%f502, %f500, %f498, %p70;
	add.f32 	%f503, %f501, 0fBF800000;
	add.f32 	%f489, %f501, 0f3F800000;
	// inline asm
	rcp.approx.ftz.f32 %f488,%f489;
	// inline asm
	add.f32 	%f504, %f503, %f503;
	mul.f32 	%f505, %f488, %f504;
	mul.f32 	%f506, %f505, %f505;
	mov.f32 	%f507, 0f3C4CAF63;
	mov.f32 	%f508, 0f3B18F0FE;
	fma.rn.f32 	%f509, %f508, %f506, %f507;
	mov.f32 	%f510, 0f3DAAAABD;
	fma.rn.f32 	%f511, %f509, %f506, %f510;
	mul.rn.f32 	%f512, %f511, %f506;
	mul.rn.f32 	%f513, %f512, %f505;
	sub.f32 	%f514, %f503, %f505;
	neg.f32 	%f515, %f505;
	add.f32 	%f516, %f514, %f514;
	fma.rn.f32 	%f517, %f515, %f503, %f516;
	mul.rn.f32 	%f518, %f488, %f517;
	add.f32 	%f519, %f513, %f505;
	sub.f32 	%f520, %f505, %f519;
	add.f32 	%f521, %f513, %f520;
	add.f32 	%f522, %f518, %f521;
	add.f32 	%f523, %f519, %f522;
	sub.f32 	%f524, %f519, %f523;
	add.f32 	%f525, %f522, %f524;
	mov.f32 	%f526, 0f3F317200;
	mul.rn.f32 	%f527, %f502, %f526;
	mov.f32 	%f528, 0f35BFBE8E;
	mul.rn.f32 	%f529, %f502, %f528;
	add.f32 	%f530, %f527, %f523;
	sub.f32 	%f531, %f527, %f530;
	add.f32 	%f532, %f523, %f531;
	add.f32 	%f533, %f525, %f532;
	add.f32 	%f534, %f529, %f533;
	add.f32 	%f535, %f530, %f534;
	sub.f32 	%f536, %f530, %f535;
	add.f32 	%f537, %f534, %f536;
	mov.f32 	%f538, 0f3EE8BA2E;
	mul.rn.f32 	%f539, %f538, %f535;
	neg.f32 	%f540, %f539;
	fma.rn.f32 	%f541, %f538, %f535, %f540;
	fma.rn.f32 	%f542, %f538, %f537, %f541;
	mov.f32 	%f543, 0f00000000;
	fma.rn.f32 	%f544, %f543, %f535, %f542;
	add.rn.f32 	%f545, %f539, %f544;
	neg.f32 	%f546, %f545;
	add.rn.f32 	%f547, %f539, %f546;
	add.rn.f32 	%f548, %f547, %f544;
	mov.b32 	 %r259, %f545;
	setp.eq.s32	%p71, %r259, 1118925336;
	add.s32 	%r260, %r259, -1;
	mov.b32 	 %f549, %r260;
	add.f32 	%f550, %f548, 0f37000000;
	selp.f32	%f551, %f549, %f545, %p71;
	selp.f32	%f151, %f550, %f548, %p71;
	mul.f32 	%f552, %f551, 0f3FB8AA3B;
	cvt.rzi.f32.f32	%f553, %f552;
	mov.f32 	%f554, 0fBF317200;
	fma.rn.f32 	%f555, %f553, %f554, %f551;
	mov.f32 	%f556, 0fB5BFBE8E;
	fma.rn.f32 	%f557, %f553, %f556, %f555;
	mul.f32 	%f558, %f557, 0f3FB8AA3B;
	ex2.approx.ftz.f32 	%f559, %f558;
	add.f32 	%f560, %f553, 0f00000000;
	ex2.approx.f32 	%f561, %f560;
	mul.f32 	%f562, %f559, %f561;
	setp.lt.f32	%p72, %f551, 0fC2D20000;
	selp.f32	%f563, 0f00000000, %f562, %p72;
	setp.gt.f32	%p73, %f551, 0f42D20000;
	selp.f32	%f976, 0f7F800000, %f563, %p73;
	setp.eq.f32	%p74, %f976, 0f7F800000;
	@%p74 bra 	BB0_78;

	fma.rn.f32 	%f976, %f976, %f151, %f976;

BB0_78:
	setp.lt.f32	%p75, %f975, 0f00000000;
	setp.eq.f32	%p76, %f149, 0f3F800000;
	and.pred  	%p3, %p75, %p76;
	mov.b32 	 %r261, %f976;
	xor.b32  	%r262, %r261, -2147483648;
	mov.b32 	 %f564, %r262;
	selp.f32	%f978, %f564, %f976, %p3;
	setp.eq.f32	%p77, %f975, 0f00000000;
	@%p77 bra 	BB0_81;
	bra.uni 	BB0_79;

BB0_81:
	add.f32 	%f567, %f975, %f975;
	selp.f32	%f978, %f567, 0f00000000, %p76;
	bra.uni 	BB0_82;

BB0_146:
	mov.u64 	%rd210, image_HDR;
	cvta.global.u64 	%rd205, %rd210;
	mov.u32 	%r348, 8;
	// inline asm
	call (%rd204), _rt_buffer_get_64, (%rd205, %r98, %r348, %rd17, %rd18, %rd24, %rd24);
	// inline asm
	mov.f32 	%f880, 0f00000000;
	// inline asm
	{  cvt.rn.f16.f32 %rs98, %f880;}

	// inline asm
	mov.u16 	%rs99, 0;
	st.v4.u16 	[%rd204], {%rs98, %rs98, %rs98, %rs99};

BB0_147:
	ld.global.u32 	%r349, [additive];
	setp.eq.s32	%p144, %r349, 0;
	@%p144 bra 	BB0_149;

	mov.u64 	%rd223, image_RNM0;
	cvta.global.u64 	%rd212, %rd223;
	mov.u32 	%r353, 8;
	// inline asm
	call (%rd211), _rt_buffer_get_64, (%rd212, %r98, %r353, %rd17, %rd18, %rd24, %rd24);
	// inline asm
	ld.v4.u16 	{%rs106, %rs107, %rs108, %rs109}, [%rd211];
	// inline asm
	{  cvt.f32.f16 %f881, %rs106;}

	// inline asm
	// inline asm
	{  cvt.f32.f16 %f882, %rs107;}

	// inline asm
	// inline asm
	{  cvt.f32.f16 %f883, %rs108;}

	// inline asm
	// inline asm
	call (%rd217), _rt_buffer_get_64, (%rd212, %r98, %r353, %rd17, %rd18, %rd24, %rd24);
	// inline asm
	add.f32 	%f884, %f881, 0f00000000;
	add.f32 	%f885, %f882, 0f00000000;
	add.f32 	%f886, %f883, 0f00000000;
	// inline asm
	{  cvt.rn.f16.f32 %rs105, %f886;}

	// inline asm
	// inline asm
	{  cvt.rn.f16.f32 %rs104, %f885;}

	// inline asm
	// inline asm
	{  cvt.rn.f16.f32 %rs103, %f884;}

	// inline asm
	mov.u16 	%rs110, 0;
	st.v4.u16 	[%rd217], {%rs103, %rs104, %rs105, %rs110};
	bra.uni 	BB0_150;

BB0_149:
	mov.u64 	%rd230, image_RNM0;
	cvta.global.u64 	%rd225, %rd230;
	mov.u32 	%r355, 8;
	// inline asm
	call (%rd224), _rt_buffer_get_64, (%rd225, %r98, %r355, %rd17, %rd18, %rd24, %rd24);
	// inline asm
	mov.f32 	%f887, 0f00000000;
	// inline asm
	{  cvt.rn.f16.f32 %rs111, %f887;}

	// inline asm
	mov.u16 	%rs112, 0;
	st.v4.u16 	[%rd224], {%rs111, %rs111, %rs111, %rs112};

BB0_150:
	ld.global.u32 	%r356, [additive];
	setp.eq.s32	%p145, %r356, 0;
	@%p145 bra 	BB0_152;

	mov.u64 	%rd243, image_RNM1;
	cvta.global.u64 	%rd232, %rd243;
	mov.u32 	%r360, 8;
	// inline asm
	call (%rd231), _rt_buffer_get_64, (%rd232, %r98, %r360, %rd17, %rd18, %rd24, %rd24);
	// inline asm
	ld.v4.u16 	{%rs119, %rs120, %rs121, %rs122}, [%rd231];
	// inline asm
	{  cvt.f32.f16 %f888, %rs119;}

	// inline asm
	// inline asm
	{  cvt.f32.f16 %f889, %rs120;}

	// inline asm
	// inline asm
	{  cvt.f32.f16 %f890, %rs121;}

	// inline asm
	// inline asm
	call (%rd237), _rt_buffer_get_64, (%rd232, %r98, %r360, %rd17, %rd18, %rd24, %rd24);
	// inline asm
	add.f32 	%f891, %f888, 0f00000000;
	add.f32 	%f892, %f889, 0f00000000;
	add.f32 	%f893, %f890, 0f00000000;
	// inline asm
	{  cvt.rn.f16.f32 %rs118, %f893;}

	// inline asm
	// inline asm
	{  cvt.rn.f16.f32 %rs117, %f892;}

	// inline asm
	// inline asm
	{  cvt.rn.f16.f32 %rs116, %f891;}

	// inline asm
	mov.u16 	%rs123, 0;
	st.v4.u16 	[%rd237], {%rs116, %rs117, %rs118, %rs123};
	bra.uni 	BB0_153;

BB0_152:
	mov.u64 	%rd250, image_RNM1;
	cvta.global.u64 	%rd245, %rd250;
	mov.u32 	%r362, 8;
	// inline asm
	call (%rd244), _rt_buffer_get_64, (%rd245, %r98, %r362, %rd17, %rd18, %rd24, %rd24);
	// inline asm
	mov.f32 	%f894, 0f00000000;
	// inline asm
	{  cvt.rn.f16.f32 %rs124, %f894;}

	// inline asm
	mov.u16 	%rs125, 0;
	st.v4.u16 	[%rd244], {%rs124, %rs124, %rs124, %rs125};

BB0_153:
	ld.global.u32 	%r363, [additive];
	setp.eq.s32	%p146, %r363, 0;
	@%p146 bra 	BB0_155;

	mov.u64 	%rd263, image_RNM2;
	cvta.global.u64 	%rd252, %rd263;
	mov.u32 	%r367, 8;
	// inline asm
	call (%rd251), _rt_buffer_get_64, (%rd252, %r98, %r367, %rd17, %rd18, %rd24, %rd24);
	// inline asm
	ld.v4.u16 	{%rs132, %rs133, %rs134, %rs135}, [%rd251];
	// inline asm
	{  cvt.f32.f16 %f895, %rs132;}

	// inline asm
	// inline asm
	{  cvt.f32.f16 %f896, %rs133;}

	// inline asm
	// inline asm
	{  cvt.f32.f16 %f897, %rs134;}

	// inline asm
	// inline asm
	call (%rd257), _rt_buffer_get_64, (%rd252, %r98, %r367, %rd17, %rd18, %rd24, %rd24);
	// inline asm
	add.f32 	%f898, %f895, 0f00000000;
	add.f32 	%f899, %f896, 0f00000000;
	add.f32 	%f900, %f897, 0f00000000;
	// inline asm
	{  cvt.rn.f16.f32 %rs131, %f900;}

	// inline asm
	// inline asm
	{  cvt.rn.f16.f32 %rs130, %f899;}

	// inline asm
	// inline asm
	{  cvt.rn.f16.f32 %rs129, %f898;}

	// inline asm
	mov.u16 	%rs136, 0;
	st.v4.u16 	[%rd257], {%rs129, %rs130, %rs131, %rs136};
	bra.uni 	BB0_156;

BB0_155:
	mov.u64 	%rd270, image_RNM2;
	cvta.global.u64 	%rd265, %rd270;
	mov.u32 	%r369, 8;
	// inline asm
	call (%rd264), _rt_buffer_get_64, (%rd265, %r98, %r369, %rd17, %rd18, %rd24, %rd24);
	// inline asm
	mov.f32 	%f901, 0f00000000;
	// inline asm
	{  cvt.rn.f16.f32 %rs137, %f901;}

	// inline asm
	mov.u16 	%rs138, 0;
	st.v4.u16 	[%rd264], {%rs137, %rs137, %rs137, %rs138};

BB0_156:
	ld.global.u32 	%r370, [additive];
	setp.eq.s32	%p147, %r370, 0;
	@%p147 bra 	BB0_158;

	mov.u64 	%rd283, image_RNM3;
	cvta.global.u64 	%rd272, %rd283;
	mov.u32 	%r374, 8;
	// inline asm
	call (%rd271), _rt_buffer_get_64, (%rd272, %r98, %r374, %rd17, %rd18, %rd24, %rd24);
	// inline asm
	ld.v4.u16 	{%rs145, %rs146, %rs147, %rs148}, [%rd271];
	// inline asm
	{  cvt.f32.f16 %f902, %rs145;}

	// inline asm
	// inline asm
	{  cvt.f32.f16 %f903, %rs146;}

	// inline asm
	// inline asm
	{  cvt.f32.f16 %f904, %rs147;}

	// inline asm
	// inline asm
	call (%rd277), _rt_buffer_get_64, (%rd272, %r98, %r374, %rd17, %rd18, %rd24, %rd24);
	// inline asm
	add.f32 	%f905, %f902, 0f00000000;
	add.f32 	%f906, %f903, 0f00000000;
	add.f32 	%f907, %f904, 0f00000000;
	// inline asm
	{  cvt.rn.f16.f32 %rs144, %f907;}

	// inline asm
	// inline asm
	{  cvt.rn.f16.f32 %rs143, %f906;}

	// inline asm
	// inline asm
	{  cvt.rn.f16.f32 %rs142, %f905;}

	// inline asm
	mov.u16 	%rs149, 0;
	st.v4.u16 	[%rd277], {%rs142, %rs143, %rs144, %rs149};
	bra.uni 	BB0_159;

BB0_158:
	mov.u64 	%rd290, image_RNM3;
	cvta.global.u64 	%rd285, %rd290;
	mov.u32 	%r376, 8;
	// inline asm
	call (%rd284), _rt_buffer_get_64, (%rd285, %r98, %r376, %rd17, %rd18, %rd24, %rd24);
	// inline asm
	mov.f32 	%f908, 0f00000000;
	// inline asm
	{  cvt.rn.f16.f32 %rs150, %f908;}

	// inline asm
	mov.u16 	%rs151, 0;
	st.v4.u16 	[%rd284], {%rs150, %rs150, %rs150, %rs151};
	bra.uni 	BB0_159;

BB0_79:
	setp.geu.f32	%p78, %f975, 0f00000000;
	@%p78 bra 	BB0_82;

	cvt.rzi.f32.f32	%f566, %f538;
	setp.neu.f32	%p79, %f566, 0f3EE8BA2E;
	selp.f32	%f978, 0f7FFFFFFF, %f978, %p79;

BB0_82:
	add.f32 	%f568, %f150, 0f3EE8BA2E;
	mov.b32 	 %r263, %f568;
	setp.lt.s32	%p81, %r263, 2139095040;
	@%p81 bra 	BB0_87;

	setp.gtu.f32	%p82, %f150, 0f7F800000;
	@%p82 bra 	BB0_86;
	bra.uni 	BB0_84;

BB0_86:
	add.f32 	%f978, %f975, 0f3EE8BA2E;
	bra.uni 	BB0_87;

BB0_84:
	setp.neu.f32	%p83, %f150, 0f7F800000;
	@%p83 bra 	BB0_87;

	selp.f32	%f978, 0fFF800000, 0f7F800000, %p3;

BB0_87:
	mul.f32 	%f569, %f978, 0f437F0000;
	setp.eq.f32	%p84, %f975, 0f3F800000;
	selp.f32	%f570, 0f437F0000, %f569, %p84;
	cvt.rzi.u32.f32	%r264, %f570;
	cvt.u16.u32	%rs14, %r264;
	mov.u16 	%rs15, 255;
	st.v2.u8 	[%rd63], {%rs14, %rs15};
	ld.global.u32 	%r408, [imageEnabled];

BB0_88:
	ld.global.f32 	%f571, [lightColor];
	mul.f32 	%f162, %f970, %f571;
	ld.global.f32 	%f572, [lightColor+4];
	mul.f32 	%f163, %f970, %f572;
	ld.global.f32 	%f573, [lightColor+8];
	mul.f32 	%f164, %f970, %f573;
	and.b32  	%r265, %r408, 1;
	setp.eq.b32	%p85, %r265, 1;
	@!%p85 bra 	BB0_123;
	bra.uni 	BB0_89;

BB0_89:
	mov.f32 	%f576, 0f3E666666;
	cvt.rzi.f32.f32	%f577, %f576;
	fma.rn.f32 	%f578, %f577, 0fC0000000, 0f3EE66666;
	abs.f32 	%f165, %f578;
	abs.f32 	%f166, %f162;
	setp.lt.f32	%p86, %f166, 0f00800000;
	mul.f32 	%f579, %f166, 0f4B800000;
	selp.f32	%f580, 0fC3170000, 0fC2FE0000, %p86;
	selp.f32	%f581, %f579, %f166, %p86;
	mov.b32 	 %r266, %f581;
	and.b32  	%r267, %r266, 8388607;
	or.b32  	%r268, %r267, 1065353216;
	mov.b32 	 %f582, %r268;
	shr.u32 	%r269, %r266, 23;
	cvt.rn.f32.u32	%f583, %r269;
	add.f32 	%f584, %f580, %f583;
	setp.gt.f32	%p87, %f582, 0f3FB504F3;
	mul.f32 	%f585, %f582, 0f3F000000;
	add.f32 	%f586, %f584, 0f3F800000;
	selp.f32	%f587, %f585, %f582, %p87;
	selp.f32	%f588, %f586, %f584, %p87;
	add.f32 	%f589, %f587, 0fBF800000;
	add.f32 	%f575, %f587, 0f3F800000;
	// inline asm
	rcp.approx.ftz.f32 %f574,%f575;
	// inline asm
	add.f32 	%f590, %f589, %f589;
	mul.f32 	%f591, %f574, %f590;
	mul.f32 	%f592, %f591, %f591;
	mov.f32 	%f593, 0f3C4CAF63;
	mov.f32 	%f594, 0f3B18F0FE;
	fma.rn.f32 	%f595, %f594, %f592, %f593;
	mov.f32 	%f596, 0f3DAAAABD;
	fma.rn.f32 	%f597, %f595, %f592, %f596;
	mul.rn.f32 	%f598, %f597, %f592;
	mul.rn.f32 	%f599, %f598, %f591;
	sub.f32 	%f600, %f589, %f591;
	neg.f32 	%f601, %f591;
	add.f32 	%f602, %f600, %f600;
	fma.rn.f32 	%f603, %f601, %f589, %f602;
	mul.rn.f32 	%f604, %f574, %f603;
	add.f32 	%f605, %f599, %f591;
	sub.f32 	%f606, %f591, %f605;
	add.f32 	%f607, %f599, %f606;
	add.f32 	%f608, %f604, %f607;
	add.f32 	%f609, %f605, %f608;
	sub.f32 	%f610, %f605, %f609;
	add.f32 	%f611, %f608, %f610;
	mov.f32 	%f612, 0f3F317200;
	mul.rn.f32 	%f613, %f588, %f612;
	mov.f32 	%f614, 0f35BFBE8E;
	mul.rn.f32 	%f615, %f588, %f614;
	add.f32 	%f616, %f613, %f609;
	sub.f32 	%f617, %f613, %f616;
	add.f32 	%f618, %f609, %f617;
	add.f32 	%f619, %f611, %f618;
	add.f32 	%f620, %f615, %f619;
	add.f32 	%f621, %f616, %f620;
	sub.f32 	%f622, %f616, %f621;
	add.f32 	%f623, %f620, %f622;
	mov.f32 	%f624, 0f3EE66666;
	mul.rn.f32 	%f625, %f624, %f621;
	neg.f32 	%f626, %f625;
	fma.rn.f32 	%f627, %f624, %f621, %f626;
	fma.rn.f32 	%f628, %f624, %f623, %f627;
	mov.f32 	%f629, 0f00000000;
	fma.rn.f32 	%f630, %f629, %f621, %f628;
	add.rn.f32 	%f631, %f625, %f630;
	neg.f32 	%f632, %f631;
	add.rn.f32 	%f633, %f625, %f632;
	add.rn.f32 	%f634, %f633, %f630;
	mov.b32 	 %r270, %f631;
	setp.eq.s32	%p88, %r270, 1118925336;
	add.s32 	%r271, %r270, -1;
	mov.b32 	 %f635, %r271;
	add.f32 	%f636, %f634, 0f37000000;
	selp.f32	%f637, %f635, %f631, %p88;
	selp.f32	%f167, %f636, %f634, %p88;
	mul.f32 	%f638, %f637, 0f3FB8AA3B;
	cvt.rzi.f32.f32	%f639, %f638;
	mov.f32 	%f640, 0fBF317200;
	fma.rn.f32 	%f641, %f639, %f640, %f637;
	mov.f32 	%f642, 0fB5BFBE8E;
	fma.rn.f32 	%f643, %f639, %f642, %f641;
	mul.f32 	%f644, %f643, 0f3FB8AA3B;
	ex2.approx.ftz.f32 	%f645, %f644;
	add.f32 	%f646, %f639, 0f00000000;
	ex2.approx.f32 	%f647, %f646;
	mul.f32 	%f648, %f645, %f647;
	setp.lt.f32	%p89, %f637, 0fC2D20000;
	selp.f32	%f649, 0f00000000, %f648, %p89;
	setp.gt.f32	%p90, %f637, 0f42D20000;
	selp.f32	%f979, 0f7F800000, %f649, %p90;
	setp.eq.f32	%p91, %f979, 0f7F800000;
	@%p91 bra 	BB0_91;

	fma.rn.f32 	%f979, %f979, %f167, %f979;

BB0_91:
	setp.lt.f32	%p92, %f162, 0f00000000;
	setp.eq.f32	%p93, %f165, 0f3F800000;
	and.pred  	%p4, %p92, %p93;
	mov.b32 	 %r272, %f979;
	xor.b32  	%r273, %r272, -2147483648;
	mov.b32 	 %f650, %r273;
	selp.f32	%f981, %f650, %f979, %p4;
	setp.eq.f32	%p94, %f162, 0f00000000;
	@%p94 bra 	BB0_94;
	bra.uni 	BB0_92;

BB0_94:
	add.f32 	%f653, %f162, %f162;
	selp.f32	%f981, %f653, 0f00000000, %p93;
	bra.uni 	BB0_95;

BB0_92:
	setp.geu.f32	%p95, %f162, 0f00000000;
	@%p95 bra 	BB0_95;

	cvt.rzi.f32.f32	%f652, %f624;
	setp.neu.f32	%p96, %f652, 0f3EE66666;
	selp.f32	%f981, 0f7FFFFFFF, %f981, %p96;

BB0_95:
	add.f32 	%f654, %f166, 0f3EE66666;
	mov.b32 	 %r274, %f654;
	setp.lt.s32	%p98, %r274, 2139095040;
	@%p98 bra 	BB0_100;

	setp.gtu.f32	%p99, %f166, 0f7F800000;
	@%p99 bra 	BB0_99;
	bra.uni 	BB0_97;

BB0_99:
	add.f32 	%f981, %f162, 0f3EE66666;
	bra.uni 	BB0_100;

BB0_97:
	setp.neu.f32	%p100, %f166, 0f7F800000;
	@%p100 bra 	BB0_100;

	selp.f32	%f981, 0fFF800000, 0f7F800000, %p4;

BB0_100:
	setp.eq.f32	%p101, %f162, 0f3F800000;
	selp.f32	%f178, 0f3F800000, %f981, %p101;
	abs.f32 	%f179, %f163;
	setp.lt.f32	%p102, %f179, 0f00800000;
	mul.f32 	%f657, %f179, 0f4B800000;
	selp.f32	%f658, 0fC3170000, 0fC2FE0000, %p102;
	selp.f32	%f659, %f657, %f179, %p102;
	mov.b32 	 %r275, %f659;
	and.b32  	%r276, %r275, 8388607;
	or.b32  	%r277, %r276, 1065353216;
	mov.b32 	 %f660, %r277;
	shr.u32 	%r278, %r275, 23;
	cvt.rn.f32.u32	%f661, %r278;
	add.f32 	%f662, %f658, %f661;
	setp.gt.f32	%p103, %f660, 0f3FB504F3;
	mul.f32 	%f663, %f660, 0f3F000000;
	add.f32 	%f664, %f662, 0f3F800000;
	selp.f32	%f665, %f663, %f660, %p103;
	selp.f32	%f666, %f664, %f662, %p103;
	add.f32 	%f667, %f665, 0fBF800000;
	add.f32 	%f656, %f665, 0f3F800000;
	// inline asm
	rcp.approx.ftz.f32 %f655,%f656;
	// inline asm
	add.f32 	%f668, %f667, %f667;
	mul.f32 	%f669, %f655, %f668;
	mul.f32 	%f670, %f669, %f669;
	fma.rn.f32 	%f673, %f594, %f670, %f593;
	fma.rn.f32 	%f675, %f673, %f670, %f596;
	mul.rn.f32 	%f676, %f675, %f670;
	mul.rn.f32 	%f677, %f676, %f669;
	sub.f32 	%f678, %f667, %f669;
	neg.f32 	%f679, %f669;
	add.f32 	%f680, %f678, %f678;
	fma.rn.f32 	%f681, %f679, %f667, %f680;
	mul.rn.f32 	%f682, %f655, %f681;
	add.f32 	%f683, %f677, %f669;
	sub.f32 	%f684, %f669, %f683;
	add.f32 	%f685, %f677, %f684;
	add.f32 	%f686, %f682, %f685;
	add.f32 	%f687, %f683, %f686;
	sub.f32 	%f688, %f683, %f687;
	add.f32 	%f689, %f686, %f688;
	mul.rn.f32 	%f691, %f666, %f612;
	mul.rn.f32 	%f693, %f666, %f614;
	add.f32 	%f694, %f691, %f687;
	sub.f32 	%f695, %f691, %f694;
	add.f32 	%f696, %f687, %f695;
	add.f32 	%f697, %f689, %f696;
	add.f32 	%f698, %f693, %f697;
	add.f32 	%f699, %f694, %f698;
	sub.f32 	%f700, %f694, %f699;
	add.f32 	%f701, %f698, %f700;
	mul.rn.f32 	%f703, %f624, %f699;
	neg.f32 	%f704, %f703;
	fma.rn.f32 	%f705, %f624, %f699, %f704;
	fma.rn.f32 	%f706, %f624, %f701, %f705;
	fma.rn.f32 	%f708, %f629, %f699, %f706;
	add.rn.f32 	%f709, %f703, %f708;
	neg.f32 	%f710, %f709;
	add.rn.f32 	%f711, %f703, %f710;
	add.rn.f32 	%f712, %f711, %f708;
	mov.b32 	 %r279, %f709;
	setp.eq.s32	%p104, %r279, 1118925336;
	add.s32 	%r280, %r279, -1;
	mov.b32 	 %f713, %r280;
	add.f32 	%f714, %f712, 0f37000000;
	selp.f32	%f715, %f713, %f709, %p104;
	selp.f32	%f180, %f714, %f712, %p104;
	mul.f32 	%f716, %f715, 0f3FB8AA3B;
	cvt.rzi.f32.f32	%f717, %f716;
	fma.rn.f32 	%f719, %f717, %f640, %f715;
	fma.rn.f32 	%f721, %f717, %f642, %f719;
	mul.f32 	%f722, %f721, 0f3FB8AA3B;
	ex2.approx.ftz.f32 	%f723, %f722;
	add.f32 	%f724, %f717, 0f00000000;
	ex2.approx.f32 	%f725, %f724;
	mul.f32 	%f726, %f723, %f725;
	setp.lt.f32	%p105, %f715, 0fC2D20000;
	selp.f32	%f727, 0f00000000, %f726, %p105;
	setp.gt.f32	%p106, %f715, 0f42D20000;
	selp.f32	%f982, 0f7F800000, %f727, %p106;
	setp.eq.f32	%p107, %f982, 0f7F800000;
	@%p107 bra 	BB0_102;

	fma.rn.f32 	%f982, %f982, %f180, %f982;

BB0_102:
	setp.lt.f32	%p108, %f163, 0f00000000;
	and.pred  	%p5, %p108, %p93;
	mov.b32 	 %r281, %f982;
	xor.b32  	%r282, %r281, -2147483648;
	mov.b32 	 %f728, %r282;
	selp.f32	%f984, %f728, %f982, %p5;
	setp.eq.f32	%p110, %f163, 0f00000000;
	@%p110 bra 	BB0_105;
	bra.uni 	BB0_103;

BB0_105:
	add.f32 	%f731, %f163, %f163;
	selp.f32	%f984, %f731, 0f00000000, %p93;
	bra.uni 	BB0_106;

BB0_103:
	setp.geu.f32	%p111, %f163, 0f00000000;
	@%p111 bra 	BB0_106;

	cvt.rzi.f32.f32	%f730, %f624;
	setp.neu.f32	%p112, %f730, 0f3EE66666;
	selp.f32	%f984, 0f7FFFFFFF, %f984, %p112;

BB0_106:
	add.f32 	%f732, %f179, 0f3EE66666;
	mov.b32 	 %r283, %f732;
	setp.lt.s32	%p114, %r283, 2139095040;
	@%p114 bra 	BB0_111;

	setp.gtu.f32	%p115, %f179, 0f7F800000;
	@%p115 bra 	BB0_110;
	bra.uni 	BB0_108;

BB0_110:
	add.f32 	%f984, %f163, 0f3EE66666;
	bra.uni 	BB0_111;

BB0_108:
	setp.neu.f32	%p116, %f179, 0f7F800000;
	@%p116 bra 	BB0_111;

	selp.f32	%f984, 0fFF800000, 0f7F800000, %p5;

BB0_111:
	setp.eq.f32	%p117, %f163, 0f3F800000;
	selp.f32	%f191, 0f3F800000, %f984, %p117;
	abs.f32 	%f192, %f164;
	setp.lt.f32	%p118, %f192, 0f00800000;
	mul.f32 	%f735, %f192, 0f4B800000;
	selp.f32	%f736, 0fC3170000, 0fC2FE0000, %p118;
	selp.f32	%f737, %f735, %f192, %p118;
	mov.b32 	 %r284, %f737;
	and.b32  	%r285, %r284, 8388607;
	or.b32  	%r286, %r285, 1065353216;
	mov.b32 	 %f738, %r286;
	shr.u32 	%r287, %r284, 23;
	cvt.rn.f32.u32	%f739, %r287;
	add.f32 	%f740, %f736, %f739;
	setp.gt.f32	%p119, %f738, 0f3FB504F3;
	mul.f32 	%f741, %f738, 0f3F000000;
	add.f32 	%f742, %f740, 0f3F800000;
	selp.f32	%f743, %f741, %f738, %p119;
	selp.f32	%f744, %f742, %f740, %p119;
	add.f32 	%f745, %f743, 0fBF800000;
	add.f32 	%f734, %f743, 0f3F800000;
	// inline asm
	rcp.approx.ftz.f32 %f733,%f734;
	// inline asm
	add.f32 	%f746, %f745, %f745;
	mul.f32 	%f747, %f733, %f746;
	mul.f32 	%f748, %f747, %f747;
	fma.rn.f32 	%f751, %f594, %f748, %f593;
	fma.rn.f32 	%f753, %f751, %f748, %f596;
	mul.rn.f32 	%f754, %f753, %f748;
	mul.rn.f32 	%f755, %f754, %f747;
	sub.f32 	%f756, %f745, %f747;
	neg.f32 	%f757, %f747;
	add.f32 	%f758, %f756, %f756;
	fma.rn.f32 	%f759, %f757, %f745, %f758;
	mul.rn.f32 	%f760, %f733, %f759;
	add.f32 	%f761, %f755, %f747;
	sub.f32 	%f762, %f747, %f761;
	add.f32 	%f763, %f755, %f762;
	add.f32 	%f764, %f760, %f763;
	add.f32 	%f765, %f761, %f764;
	sub.f32 	%f766, %f761, %f765;
	add.f32 	%f767, %f764, %f766;
	mul.rn.f32 	%f769, %f744, %f612;
	mul.rn.f32 	%f771, %f744, %f614;
	add.f32 	%f772, %f769, %f765;
	sub.f32 	%f773, %f769, %f772;
	add.f32 	%f774, %f765, %f773;
	add.f32 	%f775, %f767, %f774;
	add.f32 	%f776, %f771, %f775;
	add.f32 	%f777, %f772, %f776;
	sub.f32 	%f778, %f772, %f777;
	add.f32 	%f779, %f776, %f778;
	mul.rn.f32 	%f781, %f624, %f777;
	neg.f32 	%f782, %f781;
	fma.rn.f32 	%f783, %f624, %f777, %f782;
	fma.rn.f32 	%f784, %f624, %f779, %f783;
	fma.rn.f32 	%f786, %f629, %f777, %f784;
	add.rn.f32 	%f787, %f781, %f786;
	neg.f32 	%f788, %f787;
	add.rn.f32 	%f789, %f781, %f788;
	add.rn.f32 	%f790, %f789, %f786;
	mov.b32 	 %r288, %f787;
	setp.eq.s32	%p120, %r288, 1118925336;
	add.s32 	%r289, %r288, -1;
	mov.b32 	 %f791, %r289;
	add.f32 	%f792, %f790, 0f37000000;
	selp.f32	%f793, %f791, %f787, %p120;
	selp.f32	%f193, %f792, %f790, %p120;
	mul.f32 	%f794, %f793, 0f3FB8AA3B;
	cvt.rzi.f32.f32	%f795, %f794;
	fma.rn.f32 	%f797, %f795, %f640, %f793;
	fma.rn.f32 	%f799, %f795, %f642, %f797;
	mul.f32 	%f800, %f799, 0f3FB8AA3B;
	ex2.approx.ftz.f32 	%f801, %f800;
	add.f32 	%f802, %f795, 0f00000000;
	ex2.approx.f32 	%f803, %f802;
	mul.f32 	%f804, %f801, %f803;
	setp.lt.f32	%p121, %f793, 0fC2D20000;
	selp.f32	%f805, 0f00000000, %f804, %p121;
	setp.gt.f32	%p122, %f793, 0f42D20000;
	selp.f32	%f985, 0f7F800000, %f805, %p122;
	setp.eq.f32	%p123, %f985, 0f7F800000;
	@%p123 bra 	BB0_113;

	fma.rn.f32 	%f985, %f985, %f193, %f985;

BB0_113:
	setp.lt.f32	%p124, %f164, 0f00000000;
	and.pred  	%p6, %p124, %p93;
	mov.b32 	 %r290, %f985;
	xor.b32  	%r291, %r290, -2147483648;
	mov.b32 	 %f806, %r291;
	selp.f32	%f987, %f806, %f985, %p6;
	setp.eq.f32	%p126, %f164, 0f00000000;
	@%p126 bra 	BB0_116;
	bra.uni 	BB0_114;

BB0_116:
	add.f32 	%f809, %f164, %f164;
	selp.f32	%f987, %f809, 0f00000000, %p93;
	bra.uni 	BB0_117;

BB0_114:
	setp.geu.f32	%p127, %f164, 0f00000000;
	@%p127 bra 	BB0_117;

	cvt.rzi.f32.f32	%f808, %f624;
	setp.neu.f32	%p128, %f808, 0f3EE66666;
	selp.f32	%f987, 0f7FFFFFFF, %f987, %p128;

BB0_117:
	add.f32 	%f810, %f192, 0f3EE66666;
	mov.b32 	 %r292, %f810;
	setp.lt.s32	%p130, %r292, 2139095040;
	@%p130 bra 	BB0_122;

	setp.gtu.f32	%p131, %f192, 0f7F800000;
	@%p131 bra 	BB0_121;
	bra.uni 	BB0_119;

BB0_121:
	add.f32 	%f987, %f164, 0f3EE66666;
	bra.uni 	BB0_122;

BB0_119:
	setp.neu.f32	%p132, %f192, 0f7F800000;
	@%p132 bra 	BB0_122;

	selp.f32	%f987, 0fFF800000, 0f7F800000, %p6;

BB0_122:
	setp.eq.f32	%p133, %f164, 0f3F800000;
	selp.f32	%f811, 0f3F800000, %f987, %p133;
	cvt.u64.u32	%rd73, %r5;
	cvt.u64.u32	%rd72, %r4;
	mov.u64 	%rd76, image;
	cvta.global.u64 	%rd71, %rd76;
	// inline asm
	call (%rd70), _rt_buffer_get_64, (%rd71, %r98, %r99, %rd72, %rd73, %rd24, %rd24);
	// inline asm
	cvt.sat.f32.f32	%f812, %f811;
	mul.f32 	%f813, %f812, 0f437FFD71;
	cvt.rzi.u32.f32	%r295, %f813;
	cvt.sat.f32.f32	%f814, %f191;
	mul.f32 	%f815, %f814, 0f437FFD71;
	cvt.rzi.u32.f32	%r296, %f815;
	cvt.sat.f32.f32	%f816, %f178;
	mul.f32 	%f817, %f816, 0f437FFD71;
	cvt.rzi.u32.f32	%r297, %f817;
	cvt.u16.u32	%rs16, %r295;
	cvt.u16.u32	%rs17, %r297;
	cvt.u16.u32	%rs18, %r296;
	mov.u16 	%rs19, 255;
	st.v4.u8 	[%rd70], {%rs16, %rs18, %rs17, %rs19};
	ld.global.u32 	%r408, [imageEnabled];

BB0_123:
	cvt.u64.u32	%rd15, %r4;
	cvt.u64.u32	%rd16, %r5;
	and.b32  	%r298, %r408, 4;
	setp.eq.s32	%p134, %r298, 0;
	@%p134 bra 	BB0_127;

	ld.global.u32 	%r299, [additive];
	setp.eq.s32	%p135, %r299, 0;
	// inline asm
	{  cvt.rn.f16.f32 %rs20, %f392;}

	// inline asm
	@%p135 bra 	BB0_126;

	mov.u64 	%rd89, image_HDR;
	cvta.global.u64 	%rd78, %rd89;
	mov.u32 	%r303, 8;
	// inline asm
	call (%rd77), _rt_buffer_get_64, (%rd78, %r98, %r303, %rd15, %rd16, %rd24, %rd24);
	// inline asm
	ld.v4.u16 	{%rs27, %rs28, %rs29, %rs30}, [%rd77];
	// inline asm
	{  cvt.f32.f16 %f819, %rs27;}

	// inline asm
	// inline asm
	{  cvt.f32.f16 %f820, %rs28;}

	// inline asm
	// inline asm
	{  cvt.f32.f16 %f821, %rs29;}

	// inline asm
	// inline asm
	call (%rd83), _rt_buffer_get_64, (%rd78, %r98, %r303, %rd15, %rd16, %rd24, %rd24);
	// inline asm
	add.f32 	%f822, %f162, %f819;
	add.f32 	%f823, %f163, %f820;
	add.f32 	%f824, %f164, %f821;
	// inline asm
	{  cvt.rn.f16.f32 %rs26, %f824;}

	// inline asm
	// inline asm
	{  cvt.rn.f16.f32 %rs25, %f823;}

	// inline asm
	// inline asm
	{  cvt.rn.f16.f32 %rs24, %f822;}

	// inline asm
	st.v4.u16 	[%rd83], {%rs24, %rs25, %rs26, %rs20};
	bra.uni 	BB0_127;

BB0_126:
	mov.u64 	%rd96, image_HDR;
	cvta.global.u64 	%rd91, %rd96;
	mov.u32 	%r305, 8;
	// inline asm
	call (%rd90), _rt_buffer_get_64, (%rd91, %r98, %r305, %rd15, %rd16, %rd24, %rd24);
	// inline asm
	// inline asm
	{  cvt.rn.f16.f32 %rs33, %f164;}

	// inline asm
	// inline asm
	{  cvt.rn.f16.f32 %rs32, %f163;}

	// inline asm
	// inline asm
	{  cvt.rn.f16.f32 %rs31, %f162;}

	// inline asm
	st.v4.u16 	[%rd90], {%rs31, %rs32, %rs33, %rs20};

BB0_127:
	selp.f32	%f829, 0f3F800000, 0f3E800000, %p148;
	mul.f32 	%f204, %f829, %f971;
	mul.f32 	%f205, %f829, %f972;
	mul.f32 	%f206, %f829, %f973;
	mul.f32 	%f207, %f829, %f974;
	ld.global.f32 	%f830, [lightColor];
	mul.f32 	%f208, %f204, %f830;
	ld.global.f32 	%f831, [lightColor+4];
	mul.f32 	%f209, %f204, %f831;
	ld.global.f32 	%f832, [lightColor+8];
	mul.f32 	%f210, %f204, %f832;
	ld.global.u32 	%r306, [additive];
	setp.eq.s32	%p136, %r306, 0;
	// inline asm
	{  cvt.rn.f16.f32 %rs34, %f392;}

	// inline asm
	@%p136 bra 	BB0_129;

	mov.u64 	%rd109, image_RNM0;
	cvta.global.u64 	%rd98, %rd109;
	mov.u32 	%r310, 8;
	// inline asm
	call (%rd97), _rt_buffer_get_64, (%rd98, %r98, %r310, %rd15, %rd16, %rd24, %rd24);
	// inline asm
	ld.v4.u16 	{%rs41, %rs42, %rs43, %rs44}, [%rd97];
	// inline asm
	{  cvt.f32.f16 %f833, %rs41;}

	// inline asm
	// inline asm
	{  cvt.f32.f16 %f834, %rs42;}

	// inline asm
	// inline asm
	{  cvt.f32.f16 %f835, %rs43;}

	// inline asm
	// inline asm
	call (%rd103), _rt_buffer_get_64, (%rd98, %r98, %r310, %rd15, %rd16, %rd24, %rd24);
	// inline asm
	add.f32 	%f836, %f208, %f833;
	add.f32 	%f837, %f209, %f834;
	add.f32 	%f838, %f210, %f835;
	// inline asm
	{  cvt.rn.f16.f32 %rs40, %f838;}

	// inline asm
	// inline asm
	{  cvt.rn.f16.f32 %rs39, %f837;}

	// inline asm
	// inline asm
	{  cvt.rn.f16.f32 %rs38, %f836;}

	// inline asm
	st.v4.u16 	[%rd103], {%rs38, %rs39, %rs40, %rs34};
	bra.uni 	BB0_130;

BB0_129:
	mov.u64 	%rd116, image_RNM0;
	cvta.global.u64 	%rd111, %rd116;
	mov.u32 	%r312, 8;
	// inline asm
	call (%rd110), _rt_buffer_get_64, (%rd111, %r98, %r312, %rd15, %rd16, %rd24, %rd24);
	// inline asm
	// inline asm
	{  cvt.rn.f16.f32 %rs47, %f210;}

	// inline asm
	// inline asm
	{  cvt.rn.f16.f32 %rs46, %f209;}

	// inline asm
	// inline asm
	{  cvt.rn.f16.f32 %rs45, %f208;}

	// inline asm
	st.v4.u16 	[%rd110], {%rs45, %rs46, %rs47, %rs34};

BB0_130:
	mov.f32 	%f843, 0f34000000;
	max.f32 	%f844, %f204, %f843;
	div.rn.f32 	%f845, %f205, %f844;
	fma.rn.f32 	%f211, %f845, 0f3F000000, 0f3F000000;
	div.rn.f32 	%f846, %f206, %f844;
	fma.rn.f32 	%f212, %f846, 0f3F000000, 0f3F000000;
	div.rn.f32 	%f847, %f207, %f844;
	fma.rn.f32 	%f213, %f847, 0f3F000000, 0f3F000000;
	ld.global.u32 	%r313, [additive];
	setp.eq.s32	%p137, %r313, 0;
	// inline asm
	{  cvt.rn.f16.f32 %rs48, %f392;}

	// inline asm
	@%p137 bra 	BB0_132;

	mov.u64 	%rd129, image_RNM1;
	cvta.global.u64 	%rd118, %rd129;
	mov.u32 	%r317, 8;
	// inline asm
	call (%rd117), _rt_buffer_get_64, (%rd118, %r98, %r317, %rd15, %rd16, %rd24, %rd24);
	// inline asm
	ld.v4.u16 	{%rs55, %rs56, %rs57, %rs58}, [%rd117];
	// inline asm
	{  cvt.f32.f16 %f848, %rs55;}

	// inline asm
	// inline asm
	{  cvt.f32.f16 %f849, %rs56;}

	// inline asm
	// inline asm
	{  cvt.f32.f16 %f850, %rs57;}

	// inline asm
	// inline asm
	call (%rd123), _rt_buffer_get_64, (%rd118, %r98, %r317, %rd15, %rd16, %rd24, %rd24);
	// inline asm
	add.f32 	%f851, %f211, %f848;
	add.f32 	%f852, %f211, %f849;
	add.f32 	%f853, %f211, %f850;
	// inline asm
	{  cvt.rn.f16.f32 %rs54, %f853;}

	// inline asm
	// inline asm
	{  cvt.rn.f16.f32 %rs53, %f852;}

	// inline asm
	// inline asm
	{  cvt.rn.f16.f32 %rs52, %f851;}

	// inline asm
	st.v4.u16 	[%rd123], {%rs52, %rs53, %rs54, %rs48};
	bra.uni 	BB0_133;

BB0_132:
	mov.u64 	%rd136, image_RNM1;
	cvta.global.u64 	%rd131, %rd136;
	mov.u32 	%r319, 8;
	// inline asm
	call (%rd130), _rt_buffer_get_64, (%rd131, %r98, %r319, %rd15, %rd16, %rd24, %rd24);
	// inline asm
	// inline asm
	{  cvt.rn.f16.f32 %rs59, %f211;}

	// inline asm
	st.v4.u16 	[%rd130], {%rs59, %rs59, %rs59, %rs48};

BB0_133:
	ld.global.u32 	%r320, [additive];
	setp.eq.s32	%p138, %r320, 0;
	// inline asm
	{  cvt.rn.f16.f32 %rs60, %f392;}

	// inline asm
	@%p138 bra 	BB0_135;

	mov.u64 	%rd149, image_RNM2;
	cvta.global.u64 	%rd138, %rd149;
	mov.u32 	%r324, 8;
	// inline asm
	call (%rd137), _rt_buffer_get_64, (%rd138, %r98, %r324, %rd15, %rd16, %rd24, %rd24);
	// inline asm
	ld.v4.u16 	{%rs67, %rs68, %rs69, %rs70}, [%rd137];
	// inline asm
	{  cvt.f32.f16 %f856, %rs67;}

	// inline asm
	// inline asm
	{  cvt.f32.f16 %f857, %rs68;}

	// inline asm
	// inline asm
	{  cvt.f32.f16 %f858, %rs69;}

	// inline asm
	// inline asm
	call (%rd143), _rt_buffer_get_64, (%rd138, %r98, %r324, %rd15, %rd16, %rd24, %rd24);
	// inline asm
	add.f32 	%f859, %f212, %f856;
	add.f32 	%f860, %f212, %f857;
	add.f32 	%f861, %f212, %f858;
	// inline asm
	{  cvt.rn.f16.f32 %rs66, %f861;}

	// inline asm
	// inline asm
	{  cvt.rn.f16.f32 %rs65, %f860;}

	// inline asm
	// inline asm
	{  cvt.rn.f16.f32 %rs64, %f859;}

	// inline asm
	st.v4.u16 	[%rd143], {%rs64, %rs65, %rs66, %rs60};
	bra.uni 	BB0_136;

BB0_135:
	mov.u64 	%rd156, image_RNM2;
	cvta.global.u64 	%rd151, %rd156;
	mov.u32 	%r326, 8;
	// inline asm
	call (%rd150), _rt_buffer_get_64, (%rd151, %r98, %r326, %rd15, %rd16, %rd24, %rd24);
	// inline asm
	// inline asm
	{  cvt.rn.f16.f32 %rs71, %f212;}

	// inline asm
	st.v4.u16 	[%rd150], {%rs71, %rs71, %rs71, %rs60};

BB0_136:
	ld.global.u32 	%r327, [additive];
	setp.eq.s32	%p139, %r327, 0;
	// inline asm
	{  cvt.rn.f16.f32 %rs72, %f392;}

	// inline asm
	@%p139 bra 	BB0_138;

	mov.u64 	%rd169, image_RNM3;
	cvta.global.u64 	%rd158, %rd169;
	mov.u32 	%r331, 8;
	// inline asm
	call (%rd157), _rt_buffer_get_64, (%rd158, %r98, %r331, %rd15, %rd16, %rd24, %rd24);
	// inline asm
	ld.v4.u16 	{%rs79, %rs80, %rs81, %rs82}, [%rd157];
	// inline asm
	{  cvt.f32.f16 %f864, %rs79;}

	// inline asm
	// inline asm
	{  cvt.f32.f16 %f865, %rs80;}

	// inline asm
	// inline asm
	{  cvt.f32.f16 %f866, %rs81;}

	// inline asm
	// inline asm
	call (%rd163), _rt_buffer_get_64, (%rd158, %r98, %r331, %rd15, %rd16, %rd24, %rd24);
	// inline asm
	add.f32 	%f867, %f213, %f864;
	add.f32 	%f868, %f213, %f865;
	add.f32 	%f869, %f213, %f866;
	// inline asm
	{  cvt.rn.f16.f32 %rs78, %f869;}

	// inline asm
	// inline asm
	{  cvt.rn.f16.f32 %rs77, %f868;}

	// inline asm
	// inline asm
	{  cvt.rn.f16.f32 %rs76, %f867;}

	// inline asm
	st.v4.u16 	[%rd163], {%rs76, %rs77, %rs78, %rs72};
	bra.uni 	BB0_159;

BB0_138:
	mov.u64 	%rd176, image_RNM3;
	cvta.global.u64 	%rd171, %rd176;
	mov.u32 	%r333, 8;
	// inline asm
	call (%rd170), _rt_buffer_get_64, (%rd171, %r98, %r333, %rd15, %rd16, %rd24, %rd24);
	// inline asm
	// inline asm
	{  cvt.rn.f16.f32 %rs83, %f213;}

	// inline asm
	st.v4.u16 	[%rd170], {%rs83, %rs83, %rs83, %rs72};

BB0_159:
	ret;
}