//
// 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 1 .b8 rnd_seeds2[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[44];
	.reg .b64 	%SP;
	.reg .b64 	%SPL;
	.reg .pred 	%p<150>;
	.reg .b16 	%rs<159>;
	.reg .f32 	%f<1321>;
	.reg .b32 	%r<422>;
	.reg .b64 	%rd<330>;


	mov.u64 	%rd329, __local_depot0;
	cvta.local.u64 	%SP, %rd329;
	ld.global.u32 	%r1, [lightMeshBufferSize];
	abs.s32 	%r2, %r1;
	ld.global.u32 	%r3, [samples];
	ld.global.v2.u32 	{%r105, %r106}, [pixelID];
	cvt.u64.u32	%rd24, %r105;
	cvt.u64.u32	%rd25, %r106;
	mov.u64 	%rd28, uvnormal;
	cvta.global.u64 	%rd23, %rd28;
	mov.u32 	%r103, 2;
	mov.u32 	%r104, 4;
	mov.u64 	%rd27, 0;
	// inline asm
	call (%rd22), _rt_buffer_get_64, (%rd23, %r103, %r104, %rd24, %rd25, %rd27, %rd27);
	// inline asm
	ld.u32 	%r4, [%rd22];
	shr.u32 	%r109, %r4, 16;
	cvt.u16.u32	%rs1, %r109;
	and.b16  	%rs7, %rs1, 255;
	cvt.u16.u32	%rs8, %r4;
	or.b16  	%rs9, %rs8, %rs7;
	setp.eq.s16	%p7, %rs9, 0;
	mov.f32 	%f1166, 0f00000000;
	mov.f32 	%f1167, %f1166;
	mov.f32 	%f1168, %f1166;
	@%p7 bra 	BB0_2;

	ld.u8 	%rs10, [%rd22+1];
	and.b16  	%rs12, %rs8, 255;
	cvt.rn.f32.u16	%f358, %rs12;
	div.rn.f32 	%f359, %f358, 0f437F0000;
	fma.rn.f32 	%f360, %f359, 0f40000000, 0fBF800000;
	cvt.rn.f32.u16	%f361, %rs10;
	div.rn.f32 	%f362, %f361, 0f437F0000;
	fma.rn.f32 	%f363, %f362, 0f40000000, 0fBF800000;
	cvt.rn.f32.u16	%f364, %rs7;
	div.rn.f32 	%f365, %f364, 0f437F0000;
	fma.rn.f32 	%f366, %f365, 0f40000000, 0fBF800000;
	mul.f32 	%f367, %f363, %f363;
	fma.rn.f32 	%f368, %f360, %f360, %f367;
	fma.rn.f32 	%f369, %f366, %f366, %f368;
	sqrt.rn.f32 	%f370, %f369;
	rcp.rn.f32 	%f371, %f370;
	mul.f32 	%f1166, %f360, %f371;
	mul.f32 	%f1167, %f363, %f371;
	mul.f32 	%f1168, %f366, %f371;

BB0_2:
	ld.global.v2.u32 	{%r110, %r111}, [pixelID];
	ld.global.v2.u32 	{%r113, %r114}, [tileInfo];
	add.s32 	%r5, %r110, %r113;
	add.s32 	%r6, %r111, %r114;
	setp.eq.f32	%p8, %f1167, 0f00000000;
	setp.eq.f32	%p9, %f1166, 0f00000000;
	and.pred  	%p10, %p9, %p8;
	setp.eq.f32	%p11, %f1168, 0f00000000;
	and.pred  	%p12, %p10, %p11;
	@%p12 bra 	BB0_142;
	bra.uni 	BB0_3;

BB0_142:
	ld.global.u32 	%r421, [imageEnabled];
	and.b32  	%r343, %r421, 1;
	setp.eq.b32	%p141, %r343, 1;
	@!%p141 bra 	BB0_144;
	bra.uni 	BB0_143;

BB0_143:
	cvt.u64.u32	%rd206, %r5;
	cvt.u64.u32	%rd207, %r6;
	mov.u64 	%rd210, image;
	cvta.global.u64 	%rd205, %rd210;
	mov.u64 	%rd209, 0;
	// inline asm
	call (%rd204), _rt_buffer_get_64, (%rd205, %r103, %r104, %rd206, %rd207, %rd209, %rd209);
	// inline asm
	mov.u16 	%rs91, 0;
	st.v4.u8 	[%rd204], {%rs91, %rs91, %rs91, %rs91};
	ld.global.u32 	%r421, [imageEnabled];

BB0_144:
	and.b32  	%r346, %r421, 8;
	setp.eq.s32	%p142, %r346, 0;
	@%p142 bra 	BB0_146;

	cvt.u64.u32	%rd214, %r6;
	cvt.u64.u32	%rd213, %r5;
	mov.u64 	%rd217, image_Mask;
	cvta.global.u64 	%rd212, %rd217;
	mov.u64 	%rd216, 0;
	// inline asm
	call (%rd211), _rt_buffer_get_64, (%rd212, %r103, %r103, %rd213, %rd214, %rd216, %rd216);
	// inline asm
	mov.f32 	%f1121, 0f00000000;
	cvt.rzi.u32.f32	%r349, %f1121;
	cvt.u16.u32	%rs92, %r349;
	mov.u16 	%rs93, 0;
	st.v2.u8 	[%rd211], {%rs92, %rs93};
	ld.global.u32 	%r421, [imageEnabled];

BB0_146:
	cvt.u64.u32	%rd20, %r5;
	cvt.u64.u32	%rd21, %r6;
	and.b32  	%r350, %r421, 4;
	setp.eq.s32	%p143, %r350, 0;
	@%p143 bra 	BB0_150;

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

	mov.u64 	%rd230, image_HDR;
	cvta.global.u64 	%rd219, %rd230;
	mov.u32 	%r355, 8;
	mov.u64 	%rd229, 0;
	// inline asm
	call (%rd218), _rt_buffer_get_64, (%rd219, %r103, %r355, %rd20, %rd21, %rd229, %rd229);
	// inline asm
	ld.v4.u16 	{%rs100, %rs101, %rs102, %rs103}, [%rd218];
	// inline asm
	{  cvt.f32.f16 %f1122, %rs100;}

	// inline asm
	// inline asm
	{  cvt.f32.f16 %f1123, %rs101;}

	// inline asm
	// inline asm
	{  cvt.f32.f16 %f1124, %rs102;}

	// inline asm
	// inline asm
	call (%rd224), _rt_buffer_get_64, (%rd219, %r103, %r355, %rd20, %rd21, %rd229, %rd229);
	// inline asm
	add.f32 	%f1125, %f1122, 0f00000000;
	add.f32 	%f1126, %f1123, 0f00000000;
	add.f32 	%f1127, %f1124, 0f00000000;
	// inline asm
	{  cvt.rn.f16.f32 %rs99, %f1127;}

	// inline asm
	// inline asm
	{  cvt.rn.f16.f32 %rs98, %f1126;}

	// inline asm
	// inline asm
	{  cvt.rn.f16.f32 %rs97, %f1125;}

	// inline asm
	mov.u16 	%rs104, 0;
	st.v4.u16 	[%rd224], {%rs97, %rs98, %rs99, %rs104};
	bra.uni 	BB0_150;

BB0_3:
	ld.global.v2.u32 	{%r123, %r124}, [pixelID];
	cvt.u64.u32	%rd31, %r123;
	cvt.u64.u32	%rd32, %r124;
	mov.u64 	%rd41, uvpos;
	cvta.global.u64 	%rd30, %rd41;
	mov.u32 	%r119, 12;
	// inline asm
	call (%rd29), _rt_buffer_get_64, (%rd30, %r103, %r119, %rd31, %rd32, %rd27, %rd27);
	// inline asm
	ld.f32 	%f9, [%rd29+8];
	ld.f32 	%f8, [%rd29+4];
	ld.f32 	%f7, [%rd29];
	mul.f32 	%f388, %f7, 0f3456BF95;
	mul.f32 	%f389, %f8, 0f3456BF95;
	mul.f32 	%f390, %f9, 0f3456BF95;
	abs.f32 	%f10, %f1166;
	div.rn.f32 	%f391, %f388, %f10;
	abs.f32 	%f392, %f1167;
	div.rn.f32 	%f393, %f389, %f392;
	abs.f32 	%f11, %f1168;
	div.rn.f32 	%f394, %f390, %f11;
	abs.f32 	%f395, %f391;
	abs.f32 	%f396, %f393;
	abs.f32 	%f397, %f394;
	mov.f32 	%f398, 0f38D1B717;
	max.f32 	%f399, %f395, %f398;
	max.f32 	%f400, %f396, %f398;
	max.f32 	%f401, %f397, %f398;
	fma.rn.f32 	%f12, %f1166, %f399, %f7;
	fma.rn.f32 	%f13, %f1167, %f400, %f8;
	fma.rn.f32 	%f14, %f1168, %f401, %f9;
	ld.global.v2.u32 	{%r127, %r128}, [pixelID];
	cvt.u64.u32	%rd37, %r127;
	cvt.u64.u32	%rd38, %r128;
	mov.u64 	%rd42, rnd_seeds;
	cvta.global.u64 	%rd36, %rd42;
	// inline asm
	call (%rd35), _rt_buffer_get_64, (%rd36, %r103, %r104, %rd37, %rd38, %rd27, %rd27);
	// inline asm
	ld.u32 	%r397, [%rd35];
	mov.pred 	%p13, 0;
	mov.f32 	%f19, 0f00000000;
	mov.u32 	%r9, 0;
	setp.lt.s32	%p14, %r2, 1;
	@%p14 bra 	BB0_4;

	ld.global.f32 	%f15, [lightPointSize];
	mul.f32 	%f16, %f12, 0f3456BF95;
	mul.f32 	%f17, %f13, 0f3456BF95;
	mul.f32 	%f18, %f14, 0f3456BF95;
	mov.f32 	%f19, 0f00000000;
	mov.u32 	%r390, 0;
	abs.f32 	%f541, %f17;
	abs.f32 	%f542, %f16;
	max.f32 	%f543, %f542, %f541;
	abs.f32 	%f544, %f18;
	max.f32 	%f545, %f543, %f544;
	mov.u32 	%r9, %r390;
	mov.f32 	%f20, %f19;
	mov.f32 	%f21, %f19;
	mov.f32 	%f22, %f19;
	mov.f32 	%f23, %f19;
	mov.f32 	%f24, %f19;
	mov.f32 	%f25, %f19;
	mov.f32 	%f26, %f19;
	mov.f32 	%f27, %f19;
	mov.f32 	%f28, %f19;
	mov.f32 	%f29, %f19;
	mov.f32 	%f30, %f19;
	mov.f32 	%f31, %f19;
	mov.f32 	%f32, %f19;
	mov.f32 	%f33, %f19;
	mov.f32 	%f34, %f19;

BB0_6:
	mul.lo.s32 	%r10, %r390, 3;
	cvt.s64.s32	%rd45, %r10;
	mov.u64 	%rd49, lightMeshBuffer;
	cvta.global.u64 	%rd44, %rd49;
	mov.u32 	%r133, 1;
	mov.u64 	%rd48, 0;
	// inline asm
	call (%rd43), _rt_buffer_get_64, (%rd44, %r133, %r119, %rd45, %rd48, %rd48, %rd48);
	// inline asm
	ld.f32 	%f418, [%rd43];
	sub.f32 	%f419, %f418, %f7;
	ld.f32 	%f420, [%rd43+4];
	sub.f32 	%f421, %f420, %f8;
	ld.f32 	%f422, [%rd43+8];
	sub.f32 	%f423, %f422, %f9;
	mul.f32 	%f424, %f421, %f421;
	fma.rn.f32 	%f425, %f419, %f419, %f424;
	fma.rn.f32 	%f426, %f423, %f423, %f425;
	sqrt.rn.f32 	%f35, %f426;
	rcp.rn.f32 	%f427, %f35;
	mul.f32 	%f36, %f419, %f427;
	mul.f32 	%f37, %f421, %f427;
	mul.f32 	%f38, %f423, %f427;
	mul.f32 	%f428, %f1167, %f37;
	fma.rn.f32 	%f429, %f1166, %f36, %f428;
	fma.rn.f32 	%f39, %f1168, %f38, %f429;
	setp.leu.f32	%p15, %f39, 0f00000000;
	@%p15 bra 	BB0_23;

	mov.u64 	%rd321, 0;
	setp.eq.s32	%p16, %r3, 0;
	mul.f32 	%f430, %f35, %f35;
	div.rn.f32 	%f40, %f15, %f430;
	ld.global.v2.u32 	{%r137, %r138}, [pixelID];
	cvt.u64.u32	%rd56, %r138;
	cvt.u64.u32	%rd57, %r137;
	and.b64  	%rd52, %rd57, 255;
	and.b64  	%rd53, %rd56, 255;
	mov.u64 	%rd58, rnd_seeds2;
	cvta.global.u64 	%rd51, %rd58;
	mov.u32 	%r136, 3;
	// inline asm
	call (%rd50), _rt_buffer_get_64, (%rd51, %r103, %r136, %rd52, %rd53, %rd321, %rd321);
	// inline asm
	@%p16 bra 	BB0_9;

	ld.u8 	%rs14, [%rd50];
	cvt.rn.f32.u16	%f431, %rs14;
	div.rn.f32 	%f432, %f431, 0fC37F0000;
	fma.rn.f32 	%f433, %f432, 0f3F333333, 0f3F800000;
	setp.gt.f32	%p18, %f40, %f433;
	mov.pred 	%p149, -1;
	@%p18 bra 	BB0_25;

BB0_9:
	ld.global.f32 	%f436, [lightInvCutoff];
	mul.f32 	%f41, %f35, %f436;
	mov.f32 	%f440, 0f40800000;
	abs.f32 	%f43, %f41;
	setp.lt.f32	%p19, %f43, 0f00800000;
	mul.f32 	%f442, %f43, 0f4B800000;
	selp.f32	%f443, 0fC3170000, 0fC2FE0000, %p19;
	selp.f32	%f444, %f442, %f43, %p19;
	mov.b32 	 %r141, %f444;
	and.b32  	%r142, %r141, 8388607;
	or.b32  	%r143, %r142, 1065353216;
	mov.b32 	 %f445, %r143;
	shr.u32 	%r144, %r141, 23;
	cvt.rn.f32.u32	%f446, %r144;
	add.f32 	%f447, %f443, %f446;
	setp.gt.f32	%p20, %f445, 0f3FB504F3;
	mul.f32 	%f448, %f445, 0f3F000000;
	add.f32 	%f449, %f447, 0f3F800000;
	selp.f32	%f450, %f448, %f445, %p20;
	selp.f32	%f451, %f449, %f447, %p20;
	add.f32 	%f452, %f450, 0fBF800000;
	add.f32 	%f435, %f450, 0f3F800000;
	// inline asm
	rcp.approx.ftz.f32 %f434,%f435;
	// inline asm
	add.f32 	%f453, %f452, %f452;
	mul.f32 	%f454, %f434, %f453;
	mul.f32 	%f455, %f454, %f454;
	mov.f32 	%f456, 0f3C4CAF63;
	mov.f32 	%f457, 0f3B18F0FE;
	fma.rn.f32 	%f458, %f457, %f455, %f456;
	mov.f32 	%f459, 0f3DAAAABD;
	fma.rn.f32 	%f460, %f458, %f455, %f459;
	mul.rn.f32 	%f461, %f460, %f455;
	mul.rn.f32 	%f462, %f461, %f454;
	sub.f32 	%f463, %f452, %f454;
	neg.f32 	%f464, %f454;
	add.f32 	%f465, %f463, %f463;
	fma.rn.f32 	%f466, %f464, %f452, %f465;
	mul.rn.f32 	%f467, %f434, %f466;
	add.f32 	%f468, %f462, %f454;
	sub.f32 	%f469, %f454, %f468;
	add.f32 	%f470, %f462, %f469;
	add.f32 	%f471, %f467, %f470;
	add.f32 	%f472, %f468, %f471;
	sub.f32 	%f473, %f468, %f472;
	add.f32 	%f474, %f471, %f473;
	mov.f32 	%f475, 0f3F317200;
	mul.rn.f32 	%f476, %f451, %f475;
	mov.f32 	%f477, 0f35BFBE8E;
	mul.rn.f32 	%f478, %f451, %f477;
	add.f32 	%f479, %f476, %f472;
	sub.f32 	%f480, %f476, %f479;
	add.f32 	%f481, %f472, %f480;
	add.f32 	%f482, %f474, %f481;
	add.f32 	%f483, %f478, %f482;
	add.f32 	%f484, %f479, %f483;
	sub.f32 	%f485, %f479, %f484;
	add.f32 	%f486, %f483, %f485;
	mul.rn.f32 	%f487, %f440, %f484;
	neg.f32 	%f488, %f487;
	fma.rn.f32 	%f489, %f440, %f484, %f488;
	fma.rn.f32 	%f490, %f440, %f486, %f489;
	mov.f32 	%f491, 0f00000000;
	fma.rn.f32 	%f492, %f491, %f484, %f490;
	add.rn.f32 	%f493, %f487, %f492;
	neg.f32 	%f494, %f493;
	add.rn.f32 	%f495, %f487, %f494;
	add.rn.f32 	%f496, %f495, %f492;
	mov.b32 	 %r145, %f493;
	setp.eq.s32	%p21, %r145, 1118925336;
	add.s32 	%r146, %r145, -1;
	mov.b32 	 %f497, %r146;
	add.f32 	%f498, %f496, 0f37000000;
	selp.f32	%f499, %f497, %f493, %p21;
	selp.f32	%f44, %f498, %f496, %p21;
	mul.f32 	%f500, %f499, 0f3FB8AA3B;
	cvt.rzi.f32.f32	%f501, %f500;
	mov.f32 	%f502, 0fBF317200;
	fma.rn.f32 	%f503, %f501, %f502, %f499;
	mov.f32 	%f504, 0fB5BFBE8E;
	fma.rn.f32 	%f505, %f501, %f504, %f503;
	mul.f32 	%f506, %f505, 0f3FB8AA3B;
	ex2.approx.ftz.f32 	%f507, %f506;
	add.f32 	%f508, %f501, 0f00000000;
	ex2.approx.f32 	%f509, %f508;
	mul.f32 	%f510, %f507, %f509;
	setp.lt.f32	%p22, %f499, 0fC2D20000;
	selp.f32	%f511, 0f00000000, %f510, %p22;
	setp.gt.f32	%p23, %f499, 0f42D20000;
	selp.f32	%f1185, 0f7F800000, %f511, %p23;
	setp.eq.f32	%p24, %f1185, 0f7F800000;
	@%p24 bra 	BB0_11;

	fma.rn.f32 	%f1185, %f1185, %f44, %f1185;

BB0_11:
	mov.f32 	%f1162, 0f40000000;
	cvt.rzi.f32.f32	%f1161, %f1162;
	add.f32 	%f1160, %f1161, %f1161;
	mov.f32 	%f1159, 0f40800000;
	sub.f32 	%f1158, %f1159, %f1160;
	abs.f32 	%f1157, %f1158;
	setp.lt.f32	%p25, %f41, 0f00000000;
	setp.eq.f32	%p26, %f1157, 0f3F800000;
	and.pred  	%p1, %p25, %p26;
	mov.b32 	 %r147, %f1185;
	xor.b32  	%r148, %r147, -2147483648;
	mov.b32 	 %f512, %r148;
	selp.f32	%f1187, %f512, %f1185, %p1;
	setp.eq.f32	%p27, %f41, 0f00000000;
	@%p27 bra 	BB0_14;
	bra.uni 	BB0_12;

BB0_14:
	add.f32 	%f515, %f41, %f41;
	selp.f32	%f1187, %f515, 0f00000000, %p26;
	bra.uni 	BB0_15;

BB0_12:
	setp.geu.f32	%p28, %f41, 0f00000000;
	@%p28 bra 	BB0_15;

	mov.f32 	%f1165, 0f40800000;
	cvt.rzi.f32.f32	%f514, %f1165;
	setp.neu.f32	%p29, %f514, 0f40800000;
	selp.f32	%f1187, 0f7FFFFFFF, %f1187, %p29;

BB0_15:
	add.f32 	%f516, %f43, 0f40800000;
	mov.b32 	 %r149, %f516;
	setp.lt.s32	%p31, %r149, 2139095040;
	@%p31 bra 	BB0_20;

	setp.gtu.f32	%p32, %f43, 0f7F800000;
	@%p32 bra 	BB0_19;
	bra.uni 	BB0_17;

BB0_19:
	add.f32 	%f1187, %f41, 0f40800000;
	bra.uni 	BB0_20;

BB0_17:
	setp.neu.f32	%p33, %f43, 0f7F800000;
	@%p33 bra 	BB0_20;

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

BB0_20:
	mov.u32 	%r387, 1;
	mov.u64 	%rd320, lightMeshBuffer;
	cvta.global.u64 	%rd319, %rd320;
	mul.lo.s32 	%r386, %r390, 3;
	mov.u64 	%rd318, 0;
	mov.f32 	%f517, 0f3F800000;
	sub.f32 	%f518, %f517, %f1187;
	setp.eq.f32	%p34, %f41, 0f3F800000;
	selp.f32	%f519, 0f00000000, %f518, %p34;
	cvt.sat.f32.f32	%f520, %f519;
	mul.f32 	%f521, %f40, %f520;
	add.s32 	%r152, %r386, 1;
	cvt.s64.s32	%rd61, %r152;
	// inline asm
	call (%rd59), _rt_buffer_get_64, (%rd319, %r387, %r119, %rd61, %rd318, %rd318, %rd318);
	// inline asm
	ld.f32 	%f522, [%rd59];
	mul.f32 	%f523, %f36, %f522;
	ld.f32 	%f524, [%rd59+4];
	mul.f32 	%f525, %f37, %f524;
	neg.f32 	%f526, %f525;
	sub.f32 	%f527, %f526, %f523;
	ld.f32 	%f528, [%rd59+8];
	mul.f32 	%f529, %f38, %f528;
	sub.f32 	%f530, %f527, %f529;
	cvt.sat.f32.f32	%f531, %f530;
	mul.f32 	%f55, %f521, %f531;
	setp.leu.f32	%p35, %f55, 0f3727C5AC;
	@%p35 bra 	BB0_22;

	mov.u32 	%r389, 1;
	mov.u64 	%rd324, lightMeshBuffer;
	cvta.global.u64 	%rd323, %rd324;
	mul.lo.s32 	%r388, %r390, 3;
	mov.u64 	%rd322, 0;
	cvt.sat.f32.f32	%f540, %f39;
	add.u64 	%rd66, %SP, 12;
	cvta.to.local.u64 	%rd73, %rd66;
	max.f32 	%f538, %f545, %f398;
	sub.f32 	%f539, %f35, %f538;
	mov.u32 	%r158, 1065353216;
	st.local.u32 	[%rd73], %r158;
	ld.global.u32 	%r153, [root];
	// inline asm
	call _rt_trace_64, (%r153, %f12, %f13, %f14, %f36, %f37, %f38, %r389, %f538, %f539, %rd66, %r104);
	// inline asm
	add.s32 	%r159, %r388, 2;
	cvt.s64.s32	%rd69, %r159;
	// inline asm
	call (%rd67), _rt_buffer_get_64, (%rd323, %r389, %r119, %rd69, %rd322, %rd322, %rd322);
	// inline asm
	ld.f32 	%f547, [%rd67];
	ld.local.f32 	%f548, [%rd73];
	mul.f32 	%f549, %f548, %f547;
	ld.f32 	%f550, [%rd67+4];
	mul.f32 	%f551, %f548, %f550;
	ld.f32 	%f552, [%rd67+8];
	mul.f32 	%f553, %f548, %f552;
	mul.f32 	%f554, %f55, 0f3EA2F983;
	mul.f32 	%f555, %f554, %f549;
	mul.f32 	%f556, %f554, %f551;
	mul.f32 	%f557, %f554, %f553;
	fma.rn.f32 	%f22, %f540, %f555, %f22;
	fma.rn.f32 	%f21, %f540, %f556, %f21;
	fma.rn.f32 	%f20, %f540, %f557, %f20;
	mul.f32 	%f558, %f39, 0f40800000;
	cvt.sat.f32.f32	%f559, %f558;
	mul.f32 	%f560, %f555, %f559;
	mul.f32 	%f561, %f556, %f559;
	mul.f32 	%f562, %f559, %f557;
	fma.rn.f32 	%f28, %f36, %f560, %f28;
	fma.rn.f32 	%f27, %f36, %f561, %f27;
	fma.rn.f32 	%f26, %f36, %f562, %f26;
	fma.rn.f32 	%f31, %f37, %f560, %f31;
	fma.rn.f32 	%f30, %f37, %f561, %f30;
	fma.rn.f32 	%f29, %f37, %f562, %f29;
	fma.rn.f32 	%f34, %f38, %f560, %f34;
	fma.rn.f32 	%f33, %f38, %f561, %f33;
	fma.rn.f32 	%f32, %f38, %f562, %f32;
	add.f32 	%f25, %f25, %f560;
	add.f32 	%f24, %f24, %f561;
	add.f32 	%f23, %f23, %f562;
	add.f32 	%f19, %f19, %f548;

BB0_22:
	add.s32 	%r9, %r9, 1;

BB0_23:
	add.s32 	%r390, %r390, 1;
	setp.lt.s32	%p37, %r390, %r2;
	@%p37 bra 	BB0_6;
	bra.uni 	BB0_24;

BB0_4:
	mov.f32 	%f20, %f19;
	mov.f32 	%f21, %f19;
	mov.f32 	%f22, %f19;
	mov.f32 	%f23, %f19;
	mov.f32 	%f24, %f19;
	mov.f32 	%f25, %f19;
	mov.f32 	%f26, %f19;
	mov.f32 	%f27, %f19;
	mov.f32 	%f28, %f19;
	mov.f32 	%f29, %f19;
	mov.f32 	%f30, %f19;
	mov.f32 	%f31, %f19;
	mov.f32 	%f32, %f19;
	mov.f32 	%f33, %f19;
	mov.f32 	%f34, %f19;

BB0_24:
	mov.pred 	%p149, %p13;

BB0_25:
	cvt.rn.f32.s32	%f563, %r9;
	mov.f32 	%f564, 0f3F800000;
	max.f32 	%f565, %f563, %f564;
	rcp.rn.f32 	%f566, %f565;
	mul.f32 	%f1305, %f22, %f566;
	mul.f32 	%f1306, %f21, %f566;
	mul.f32 	%f1307, %f20, %f566;
	div.rn.f32 	%f1308, %f19, %f565;
	mul.f32 	%f1302, %f25, %f566;
	mul.f32 	%f1303, %f24, %f566;
	mul.f32 	%f1304, %f23, %f566;
	mul.f32 	%f1299, %f28, %f566;
	mul.f32 	%f1300, %f27, %f566;
	mul.f32 	%f1301, %f26, %f566;
	mul.f32 	%f1296, %f31, %f566;
	mul.f32 	%f1297, %f30, %f566;
	mul.f32 	%f1298, %f29, %f566;
	mul.f32 	%f1293, %f34, %f566;
	mul.f32 	%f1294, %f33, %f566;
	mul.f32 	%f1295, %f32, %f566;
	@!%p149 bra 	BB0_78;
	bra.uni 	BB0_26;

BB0_26:
	abs.f32 	%f1164, %f1168;
	abs.f32 	%f1163, %f1166;
	setp.gt.f32	%p38, %f1163, %f1164;
	neg.f32 	%f582, %f1167;
	selp.f32	%f583, %f582, 0f00000000, %p38;
	neg.f32 	%f584, %f1168;
	selp.f32	%f585, %f1166, %f584, %p38;
	selp.f32	%f586, 0f00000000, %f1167, %p38;
	mul.f32 	%f587, %f585, %f585;
	fma.rn.f32 	%f588, %f583, %f583, %f587;
	fma.rn.f32 	%f589, %f586, %f586, %f588;
	sqrt.rn.f32 	%f590, %f589;
	rcp.rn.f32 	%f591, %f590;
	mul.f32 	%f136, %f583, %f591;
	mul.f32 	%f137, %f585, %f591;
	mul.f32 	%f138, %f586, %f591;
	mov.f32 	%f581, 0f00000000;
	setp.lt.s32	%p39, %r3, 1;
	mov.f32 	%f1251, %f581;
	mov.f32 	%f1252, %f581;
	mov.f32 	%f1253, %f581;
	mov.f32 	%f1254, %f581;
	mov.f32 	%f1255, %f581;
	mov.f32 	%f1256, %f581;
	mov.f32 	%f1257, %f581;
	mov.f32 	%f1258, %f581;
	mov.f32 	%f1259, %f581;
	mov.f32 	%f1260, %f581;
	mov.f32 	%f1261, %f581;
	mov.f32 	%f1262, %f581;
	mov.f32 	%f1263, %f581;
	mov.f32 	%f1264, %f581;
	mov.f32 	%f1265, %f581;
	@%p39 bra 	BB0_77;

	cvt.rn.f32.s32	%f607, %r3;
	rcp.rn.f32 	%f139, %f607;
	add.u64 	%rd75, %SP, 16;
	cvta.to.local.u64 	%rd3, %rd75;
	mul.f32 	%f140, %f12, 0f3456BF95;
	mul.f32 	%f141, %f13, 0f3456BF95;
	mul.f32 	%f142, %f14, 0f3456BF95;
	add.u64 	%rd76, %SP, 0;
	cvta.to.local.u64 	%rd4, %rd76;
	mul.f32 	%f608, %f1166, %f137;
	mul.f32 	%f609, %f1167, %f136;
	sub.f32 	%f143, %f609, %f608;
	mul.f32 	%f610, %f1168, %f136;
	mul.f32 	%f611, %f1166, %f138;
	sub.f32 	%f144, %f611, %f610;
	mul.f32 	%f612, %f1167, %f138;
	mul.f32 	%f613, %f1168, %f137;
	sub.f32 	%f145, %f613, %f612;
	mov.f32 	%f1251, 0f00000000;
	mov.u32 	%r160, 0;
	abs.f32 	%f614, %f141;
	abs.f32 	%f615, %f140;
	max.f32 	%f616, %f615, %f614;
	abs.f32 	%f617, %f142;
	max.f32 	%f618, %f616, %f617;
	mov.u32 	%r394, %r160;
	mov.f32 	%f1252, %f1251;
	mov.f32 	%f1253, %f1251;
	mov.f32 	%f1254, %f1251;
	mov.f32 	%f1255, %f1251;
	mov.f32 	%f1256, %f1251;
	mov.f32 	%f1257, %f1251;
	mov.f32 	%f1258, %f1251;
	mov.f32 	%f1259, %f1251;
	mov.f32 	%f1260, %f1251;
	mov.f32 	%f1261, %f1251;
	mov.f32 	%f1262, %f1251;
	mov.f32 	%f1263, %f1251;
	mov.f32 	%f1264, %f1251;
	mov.f32 	%f1265, %f1251;

BB0_28:
	cvt.rn.f32.s32	%f161, %r394;
	max.f32 	%f162, %f618, %f398;
	mov.u32 	%r396, %r160;

BB0_29:
	mad.lo.s32 	%r162, %r397, 1664525, 1013904223;
	and.b32  	%r163, %r162, 16777215;
	cvt.rn.f32.u32	%f620, %r163;
	fma.rn.f32 	%f621, %f620, 0f33800000, %f161;
	mul.f32 	%f178, %f139, %f621;
	mad.lo.s32 	%r397, %r162, 1664525, 1013904223;
	and.b32  	%r164, %r397, 16777215;
	cvt.rn.f32.u32	%f622, %r164;
	cvt.rn.f32.s32	%f623, %r396;
	fma.rn.f32 	%f624, %f622, 0f33800000, %f623;
	mul.f32 	%f625, %f139, %f624;
	mul.f32 	%f626, %f178, %f178;
	sub.f32 	%f628, %f564, %f626;
	mov.f32 	%f629, 0f00000000;
	max.f32 	%f630, %f629, %f628;
	sqrt.rn.f32 	%f179, %f630;
	mul.f32 	%f1272, %f625, 0f40C90FDB;
	abs.f32 	%f181, %f1272;
	setp.neu.f32	%p40, %f181, 0f7F800000;
	mov.f32 	%f1266, %f1272;
	@%p40 bra 	BB0_31;

	mul.rn.f32 	%f1266, %f1272, %f629;

BB0_31:
	mul.f32 	%f632, %f1266, 0f3F22F983;
	cvt.rni.s32.f32	%r407, %f632;
	cvt.rn.f32.s32	%f633, %r407;
	neg.f32 	%f634, %f633;
	mov.f32 	%f635, 0f3FC90FDA;
	fma.rn.f32 	%f636, %f634, %f635, %f1266;
	mov.f32 	%f637, 0f33A22168;
	fma.rn.f32 	%f638, %f634, %f637, %f636;
	mov.f32 	%f639, 0f27C234C5;
	fma.rn.f32 	%f1267, %f634, %f639, %f638;
	abs.f32 	%f640, %f1266;
	setp.leu.f32	%p41, %f640, 0f47CE4780;
	@%p41 bra 	BB0_42;

	mov.b32 	 %r21, %f1266;
	shr.u32 	%r22, %r21, 23;
	shl.b32 	%r167, %r21, 8;
	or.b32  	%r23, %r167, -2147483648;
	mov.u32 	%r398, 0;
	mov.u64 	%rd326, 0;
	mov.u64 	%rd325, %rd3;
	mov.u32 	%r399, %r398;

BB0_33:
	.pragma "nounroll";
	shl.b64 	%rd78, %rd326, 2;
	mov.u64 	%rd79, __cudart_i2opi_f;
	add.s64 	%rd80, %rd79, %rd78;
	ld.const.u32 	%r170, [%rd80];
	// inline asm
	{
	mad.lo.cc.u32   %r168, %r170, %r23, %r399;
	madc.hi.u32     %r399, %r170, %r23,  0;
	}
	// inline asm
	st.local.u32 	[%rd325], %r168;
	add.s32 	%r398, %r398, 1;
	cvt.s64.s32	%rd326, %r398;
	mul.wide.s32 	%rd81, %r398, 4;
	add.s64 	%rd325, %rd3, %rd81;
	setp.ne.s32	%p42, %r398, 6;
	@%p42 bra 	BB0_33;

	and.b32  	%r173, %r22, 255;
	add.s32 	%r174, %r173, -128;
	shr.u32 	%r175, %r174, 5;
	and.b32  	%r28, %r21, -2147483648;
	st.local.u32 	[%rd3+24], %r399;
	mov.u32 	%r176, 6;
	sub.s32 	%r177, %r176, %r175;
	mul.wide.s32 	%rd82, %r177, 4;
	add.s64 	%rd10, %rd3, %rd82;
	ld.local.u32 	%r400, [%rd10];
	ld.local.u32 	%r401, [%rd10+-4];
	and.b32  	%r31, %r22, 31;
	setp.eq.s32	%p43, %r31, 0;
	@%p43 bra 	BB0_36;

	mov.u32 	%r178, 32;
	sub.s32 	%r179, %r178, %r31;
	shr.u32 	%r180, %r401, %r179;
	shl.b32 	%r181, %r400, %r31;
	add.s32 	%r400, %r180, %r181;
	ld.local.u32 	%r182, [%rd10+-8];
	shr.u32 	%r183, %r182, %r179;
	shl.b32 	%r184, %r401, %r31;
	add.s32 	%r401, %r183, %r184;

BB0_36:
	shr.u32 	%r185, %r401, 30;
	shl.b32 	%r186, %r400, 2;
	add.s32 	%r402, %r185, %r186;
	shl.b32 	%r37, %r401, 2;
	shr.u32 	%r187, %r402, 31;
	shr.u32 	%r188, %r400, 30;
	add.s32 	%r38, %r187, %r188;
	setp.eq.s32	%p44, %r187, 0;
	@%p44 bra 	BB0_37;
	bra.uni 	BB0_38;

BB0_37:
	mov.u32 	%r403, %r28;
	mov.u32 	%r404, %r37;
	bra.uni 	BB0_39;

BB0_38:
	not.b32 	%r189, %r402;
	neg.s32 	%r404, %r37;
	setp.eq.s32	%p45, %r37, 0;
	selp.u32	%r190, 1, 0, %p45;
	add.s32 	%r402, %r190, %r189;
	xor.b32  	%r403, %r28, -2147483648;

BB0_39:
	clz.b32 	%r406, %r402;
	setp.eq.s32	%p46, %r406, 0;
	shl.b32 	%r191, %r402, %r406;
	mov.u32 	%r192, 32;
	sub.s32 	%r193, %r192, %r406;
	shr.u32 	%r194, %r404, %r193;
	add.s32 	%r195, %r194, %r191;
	selp.b32	%r46, %r402, %r195, %p46;
	mov.u32 	%r196, -921707870;
	mul.hi.u32 	%r405, %r46, %r196;
	setp.eq.s32	%p47, %r28, 0;
	neg.s32 	%r197, %r38;
	selp.b32	%r407, %r38, %r197, %p47;
	setp.lt.s32	%p48, %r405, 1;
	@%p48 bra 	BB0_41;

	mul.lo.s32 	%r198, %r46, -921707870;
	shr.u32 	%r199, %r198, 31;
	shl.b32 	%r200, %r405, 1;
	add.s32 	%r405, %r199, %r200;
	add.s32 	%r406, %r406, 1;

BB0_41:
	mov.u32 	%r201, 126;
	sub.s32 	%r202, %r201, %r406;
	shl.b32 	%r203, %r202, 23;
	add.s32 	%r204, %r405, 1;
	shr.u32 	%r205, %r204, 7;
	add.s32 	%r206, %r205, 1;
	shr.u32 	%r207, %r206, 1;
	add.s32 	%r208, %r207, %r203;
	or.b32  	%r209, %r208, %r403;
	mov.b32 	 %f1267, %r209;

BB0_42:
	mul.rn.f32 	%f187, %f1267, %f1267;
	add.s32 	%r54, %r407, 1;
	and.b32  	%r55, %r54, 1;
	setp.eq.s32	%p49, %r55, 0;
	@%p49 bra 	BB0_44;
	bra.uni 	BB0_43;

BB0_44:
	mov.f32 	%f643, 0f3C08839E;
	mov.f32 	%f644, 0fB94CA1F9;
	fma.rn.f32 	%f1268, %f644, %f187, %f643;
	bra.uni 	BB0_45;

BB0_43:
	mov.f32 	%f641, 0fBAB6061A;
	mov.f32 	%f642, 0f37CCF5CE;
	fma.rn.f32 	%f1268, %f642, %f187, %f641;

BB0_45:
	@%p49 bra 	BB0_47;
	bra.uni 	BB0_46;

BB0_47:
	mov.f32 	%f648, 0fBE2AAAA3;
	fma.rn.f32 	%f649, %f1268, %f187, %f648;
	fma.rn.f32 	%f1269, %f649, %f187, %f629;
	bra.uni 	BB0_48;

BB0_46:
	mov.f32 	%f645, 0f3D2AAAA5;
	fma.rn.f32 	%f646, %f1268, %f187, %f645;
	mov.f32 	%f647, 0fBF000000;
	fma.rn.f32 	%f1269, %f646, %f187, %f647;

BB0_48:
	fma.rn.f32 	%f1270, %f1269, %f1267, %f1267;
	@%p49 bra 	BB0_50;

	fma.rn.f32 	%f1270, %f1269, %f187, %f564;

BB0_50:
	and.b32  	%r210, %r54, 2;
	setp.eq.s32	%p52, %r210, 0;
	@%p52 bra 	BB0_52;

	mov.f32 	%f653, 0fBF800000;
	fma.rn.f32 	%f1270, %f1270, %f653, %f629;

BB0_52:
	@%p40 bra 	BB0_54;

	mul.rn.f32 	%f1272, %f1272, %f629;

BB0_54:
	mul.f32 	%f655, %f1272, 0f3F22F983;
	cvt.rni.s32.f32	%r417, %f655;
	cvt.rn.f32.s32	%f656, %r417;
	neg.f32 	%f657, %f656;
	fma.rn.f32 	%f659, %f657, %f635, %f1272;
	fma.rn.f32 	%f661, %f657, %f637, %f659;
	fma.rn.f32 	%f1273, %f657, %f639, %f661;
	abs.f32 	%f663, %f1272;
	setp.leu.f32	%p54, %f663, 0f47CE4780;
	@%p54 bra 	BB0_65;

	mov.b32 	 %r57, %f1272;
	shr.u32 	%r58, %r57, 23;
	shl.b32 	%r213, %r57, 8;
	or.b32  	%r59, %r213, -2147483648;
	mov.u32 	%r408, 0;
	mov.u64 	%rd327, %rd3;
	mov.u64 	%rd328, %rd27;
	mov.u32 	%r409, %r408;

BB0_56:
	.pragma "nounroll";
	shl.b64 	%rd84, %rd328, 2;
	mov.u64 	%rd85, __cudart_i2opi_f;
	add.s64 	%rd86, %rd85, %rd84;
	ld.const.u32 	%r216, [%rd86];
	// inline asm
	{
	mad.lo.cc.u32   %r214, %r216, %r59, %r409;
	madc.hi.u32     %r409, %r216, %r59,  0;
	}
	// inline asm
	st.local.u32 	[%rd327], %r214;
	add.s32 	%r408, %r408, 1;
	cvt.s64.s32	%rd328, %r408;
	mul.wide.s32 	%rd87, %r408, 4;
	add.s64 	%rd327, %rd3, %rd87;
	setp.ne.s32	%p55, %r408, 6;
	@%p55 bra 	BB0_56;

	and.b32  	%r219, %r58, 255;
	add.s32 	%r220, %r219, -128;
	shr.u32 	%r221, %r220, 5;
	and.b32  	%r64, %r57, -2147483648;
	st.local.u32 	[%rd3+24], %r409;
	mov.u32 	%r222, 6;
	sub.s32 	%r223, %r222, %r221;
	mul.wide.s32 	%rd88, %r223, 4;
	add.s64 	%rd16, %rd3, %rd88;
	ld.local.u32 	%r410, [%rd16];
	ld.local.u32 	%r411, [%rd16+-4];
	and.b32  	%r67, %r58, 31;
	setp.eq.s32	%p56, %r67, 0;
	@%p56 bra 	BB0_59;

	mov.u32 	%r224, 32;
	sub.s32 	%r225, %r224, %r67;
	shr.u32 	%r226, %r411, %r225;
	shl.b32 	%r227, %r410, %r67;
	add.s32 	%r410, %r226, %r227;
	ld.local.u32 	%r228, [%rd16+-8];
	shr.u32 	%r229, %r228, %r225;
	shl.b32 	%r230, %r411, %r67;
	add.s32 	%r411, %r229, %r230;

BB0_59:
	shr.u32 	%r231, %r411, 30;
	shl.b32 	%r232, %r410, 2;
	add.s32 	%r412, %r231, %r232;
	shl.b32 	%r73, %r411, 2;
	shr.u32 	%r233, %r412, 31;
	shr.u32 	%r234, %r410, 30;
	add.s32 	%r74, %r233, %r234;
	setp.eq.s32	%p57, %r233, 0;
	@%p57 bra 	BB0_60;
	bra.uni 	BB0_61;

BB0_60:
	mov.u32 	%r413, %r64;
	mov.u32 	%r414, %r73;
	bra.uni 	BB0_62;

BB0_61:
	not.b32 	%r235, %r412;
	neg.s32 	%r414, %r73;
	setp.eq.s32	%p58, %r73, 0;
	selp.u32	%r236, 1, 0, %p58;
	add.s32 	%r412, %r236, %r235;
	xor.b32  	%r413, %r64, -2147483648;

BB0_62:
	clz.b32 	%r416, %r412;
	setp.eq.s32	%p59, %r416, 0;
	shl.b32 	%r237, %r412, %r416;
	mov.u32 	%r238, 32;
	sub.s32 	%r239, %r238, %r416;
	shr.u32 	%r240, %r414, %r239;
	add.s32 	%r241, %r240, %r237;
	selp.b32	%r82, %r412, %r241, %p59;
	mov.u32 	%r242, -921707870;
	mul.hi.u32 	%r415, %r82, %r242;
	setp.eq.s32	%p60, %r64, 0;
	neg.s32 	%r243, %r74;
	selp.b32	%r417, %r74, %r243, %p60;
	setp.lt.s32	%p61, %r415, 1;
	@%p61 bra 	BB0_64;

	mul.lo.s32 	%r244, %r82, -921707870;
	shr.u32 	%r245, %r244, 31;
	shl.b32 	%r246, %r415, 1;
	add.s32 	%r415, %r245, %r246;
	add.s32 	%r416, %r416, 1;

BB0_64:
	mov.u32 	%r247, 126;
	sub.s32 	%r248, %r247, %r416;
	shl.b32 	%r249, %r248, 23;
	add.s32 	%r250, %r415, 1;
	shr.u32 	%r251, %r250, 7;
	add.s32 	%r252, %r251, 1;
	shr.u32 	%r253, %r252, 1;
	add.s32 	%r254, %r253, %r249;
	or.b32  	%r255, %r254, %r413;
	mov.b32 	 %f1273, %r255;

BB0_65:
	mul.rn.f32 	%f204, %f1273, %f1273;
	and.b32  	%r90, %r417, 1;
	setp.eq.s32	%p62, %r90, 0;
	@%p62 bra 	BB0_67;
	bra.uni 	BB0_66;

BB0_67:
	mov.f32 	%f666, 0f3C08839E;
	mov.f32 	%f667, 0fB94CA1F9;
	fma.rn.f32 	%f1274, %f667, %f204, %f666;
	bra.uni 	BB0_68;

BB0_66:
	mov.f32 	%f664, 0fBAB6061A;
	mov.f32 	%f665, 0f37CCF5CE;
	fma.rn.f32 	%f1274, %f665, %f204, %f664;

BB0_68:
	@%p62 bra 	BB0_70;
	bra.uni 	BB0_69;

BB0_70:
	mov.f32 	%f671, 0fBE2AAAA3;
	fma.rn.f32 	%f672, %f1274, %f204, %f671;
	fma.rn.f32 	%f1275, %f672, %f204, %f629;
	bra.uni 	BB0_71;

BB0_69:
	mov.f32 	%f668, 0f3D2AAAA5;
	fma.rn.f32 	%f669, %f1274, %f204, %f668;
	mov.f32 	%f670, 0fBF000000;
	fma.rn.f32 	%f1275, %f669, %f204, %f670;

BB0_71:
	fma.rn.f32 	%f1276, %f1275, %f1273, %f1273;
	@%p62 bra 	BB0_73;

	fma.rn.f32 	%f1276, %f1275, %f204, %f564;

BB0_73:
	and.b32  	%r256, %r417, 2;
	setp.eq.s32	%p65, %r256, 0;
	@%p65 bra 	BB0_75;

	mov.f32 	%f676, 0fBF800000;
	fma.rn.f32 	%f1276, %f1276, %f676, %f629;

BB0_75:
	mul.f32 	%f685, %f179, %f1270;
	mul.f32 	%f686, %f179, %f1276;
	mul.f32 	%f687, %f136, %f686;
	mul.f32 	%f688, %f137, %f686;
	mul.f32 	%f689, %f138, %f686;
	fma.rn.f32 	%f690, %f145, %f685, %f687;
	fma.rn.f32 	%f691, %f144, %f685, %f688;
	fma.rn.f32 	%f692, %f143, %f685, %f689;
	fma.rn.f32 	%f680, %f1166, %f178, %f690;
	fma.rn.f32 	%f681, %f1167, %f178, %f691;
	fma.rn.f32 	%f682, %f1168, %f178, %f692;
	mov.u32 	%r258, 0;
	st.local.u32 	[%rd4+8], %r258;
	st.local.u32 	[%rd4+4], %r258;
	st.local.u32 	[%rd4], %r258;
	ld.global.u32 	%r257, [root];
	mov.f32 	%f684, 0f6C4ECB8F;
	// inline asm
	call _rt_trace_64, (%r257, %f12, %f13, %f14, %f680, %f681, %f682, %r258, %f162, %f684, %rd76, %r119);
	// inline asm
	ld.local.f32 	%f693, [%rd4];
	max.f32 	%f695, %f693, %f629;
	ld.local.f32 	%f696, [%rd4+4];
	max.f32 	%f697, %f696, %f629;
	ld.local.f32 	%f698, [%rd4+8];
	max.f32 	%f699, %f698, %f629;
	fma.rn.f32 	%f1259, %f680, %f695, %f1259;
	fma.rn.f32 	%f1258, %f680, %f697, %f1258;
	fma.rn.f32 	%f1257, %f680, %f699, %f1257;
	fma.rn.f32 	%f1262, %f681, %f695, %f1262;
	fma.rn.f32 	%f1261, %f681, %f697, %f1261;
	fma.rn.f32 	%f1260, %f681, %f699, %f1260;
	fma.rn.f32 	%f1265, %f682, %f695, %f1265;
	fma.rn.f32 	%f1264, %f682, %f697, %f1264;
	fma.rn.f32 	%f1263, %f682, %f699, %f1263;
	add.f32 	%f1256, %f1256, %f695;
	add.f32 	%f1255, %f1255, %f697;
	add.f32 	%f1254, %f1254, %f699;
	mul.f32 	%f700, %f1167, %f681;
	fma.rn.f32 	%f701, %f1166, %f680, %f700;
	fma.rn.f32 	%f702, %f1168, %f682, %f701;
	cvt.sat.f32.f32	%f703, %f702;
	fma.rn.f32 	%f1253, %f695, %f703, %f1253;
	fma.rn.f32 	%f1252, %f697, %f703, %f1252;
	fma.rn.f32 	%f1251, %f699, %f703, %f1251;
	add.s32 	%r396, %r396, 1;
	setp.lt.s32	%p66, %r396, %r3;
	@%p66 bra 	BB0_29;

	add.s32 	%r394, %r394, 1;
	setp.lt.s32	%p67, %r394, %r3;
	@%p67 bra 	BB0_28;

BB0_77:
	mul.lo.s32 	%r260, %r3, %r3;
	cvt.rn.f32.s32	%f704, %r260;
	rcp.rn.f32 	%f705, %f704;
	mul.f32 	%f706, %f1253, %f705;
	mul.f32 	%f707, %f1252, %f705;
	mul.f32 	%f708, %f1251, %f705;
	div.rn.f32 	%f1308, %f581, %f704;
	mul.f32 	%f1302, %f1256, %f705;
	mul.f32 	%f1303, %f1255, %f705;
	mul.f32 	%f1304, %f1254, %f705;
	mul.f32 	%f1299, %f1259, %f705;
	mul.f32 	%f1300, %f1258, %f705;
	mul.f32 	%f1301, %f1257, %f705;
	mul.f32 	%f1296, %f1262, %f705;
	mul.f32 	%f1297, %f1261, %f705;
	mul.f32 	%f1298, %f1260, %f705;
	mul.f32 	%f1293, %f1265, %f705;
	mul.f32 	%f1294, %f1264, %f705;
	mul.f32 	%f1295, %f1263, %f705;
	fma.rn.f32 	%f1305, %f1253, %f705, %f706;
	fma.rn.f32 	%f1306, %f1252, %f705, %f707;
	fma.rn.f32 	%f1307, %f1251, %f705, %f708;

BB0_78:
	setp.gt.s32	%p68, %r1, -1;
	selp.f32	%f278, %f1308, %f1305, %p68;
	ld.global.u32 	%r419, [imageEnabled];
	and.b32  	%r261, %r419, 8;
	setp.eq.s32	%p69, %r261, 0;
	@%p69 bra 	BB0_91;

	cvt.u64.u32	%rd92, %r5;
	cvt.u64.u32	%rd93, %r6;
	mov.u64 	%rd96, image_Mask;
	cvta.global.u64 	%rd91, %rd96;
	// inline asm
	call (%rd90), _rt_buffer_get_64, (%rd91, %r103, %r103, %rd92, %rd93, %rd27, %rd27);
	// inline asm
	mov.f32 	%f712, 0f3E68BA2E;
	cvt.rzi.f32.f32	%f713, %f712;
	fma.rn.f32 	%f714, %f713, 0fC0000000, 0f3EE8BA2E;
	abs.f32 	%f279, %f714;
	abs.f32 	%f280, %f278;
	setp.lt.f32	%p70, %f280, 0f00800000;
	mul.f32 	%f715, %f280, 0f4B800000;
	selp.f32	%f716, 0fC3170000, 0fC2FE0000, %p70;
	selp.f32	%f717, %f715, %f280, %p70;
	mov.b32 	 %r264, %f717;
	and.b32  	%r265, %r264, 8388607;
	or.b32  	%r266, %r265, 1065353216;
	mov.b32 	 %f718, %r266;
	shr.u32 	%r267, %r264, 23;
	cvt.rn.f32.u32	%f719, %r267;
	add.f32 	%f720, %f716, %f719;
	setp.gt.f32	%p71, %f718, 0f3FB504F3;
	mul.f32 	%f721, %f718, 0f3F000000;
	add.f32 	%f722, %f720, 0f3F800000;
	selp.f32	%f723, %f721, %f718, %p71;
	selp.f32	%f724, %f722, %f720, %p71;
	add.f32 	%f725, %f723, 0fBF800000;
	add.f32 	%f711, %f723, 0f3F800000;
	// inline asm
	rcp.approx.ftz.f32 %f710,%f711;
	// inline asm
	add.f32 	%f726, %f725, %f725;
	mul.f32 	%f727, %f710, %f726;
	mul.f32 	%f728, %f727, %f727;
	mov.f32 	%f729, 0f3C4CAF63;
	mov.f32 	%f730, 0f3B18F0FE;
	fma.rn.f32 	%f731, %f730, %f728, %f729;
	mov.f32 	%f732, 0f3DAAAABD;
	fma.rn.f32 	%f733, %f731, %f728, %f732;
	mul.rn.f32 	%f734, %f733, %f728;
	mul.rn.f32 	%f735, %f734, %f727;
	sub.f32 	%f736, %f725, %f727;
	neg.f32 	%f737, %f727;
	add.f32 	%f738, %f736, %f736;
	fma.rn.f32 	%f739, %f737, %f725, %f738;
	mul.rn.f32 	%f740, %f710, %f739;
	add.f32 	%f741, %f735, %f727;
	sub.f32 	%f742, %f727, %f741;
	add.f32 	%f743, %f735, %f742;
	add.f32 	%f744, %f740, %f743;
	add.f32 	%f745, %f741, %f744;
	sub.f32 	%f746, %f741, %f745;
	add.f32 	%f747, %f744, %f746;
	mov.f32 	%f748, 0f3F317200;
	mul.rn.f32 	%f749, %f724, %f748;
	mov.f32 	%f750, 0f35BFBE8E;
	mul.rn.f32 	%f751, %f724, %f750;
	add.f32 	%f752, %f749, %f745;
	sub.f32 	%f753, %f749, %f752;
	add.f32 	%f754, %f745, %f753;
	add.f32 	%f755, %f747, %f754;
	add.f32 	%f756, %f751, %f755;
	add.f32 	%f757, %f752, %f756;
	sub.f32 	%f758, %f752, %f757;
	add.f32 	%f759, %f756, %f758;
	mov.f32 	%f760, 0f3EE8BA2E;
	mul.rn.f32 	%f761, %f760, %f757;
	neg.f32 	%f762, %f761;
	fma.rn.f32 	%f763, %f760, %f757, %f762;
	fma.rn.f32 	%f764, %f760, %f759, %f763;
	mov.f32 	%f765, 0f00000000;
	fma.rn.f32 	%f766, %f765, %f757, %f764;
	add.rn.f32 	%f767, %f761, %f766;
	neg.f32 	%f768, %f767;
	add.rn.f32 	%f769, %f761, %f768;
	add.rn.f32 	%f770, %f769, %f766;
	mov.b32 	 %r268, %f767;
	setp.eq.s32	%p72, %r268, 1118925336;
	add.s32 	%r269, %r268, -1;
	mov.b32 	 %f771, %r269;
	add.f32 	%f772, %f770, 0f37000000;
	selp.f32	%f773, %f771, %f767, %p72;
	selp.f32	%f281, %f772, %f770, %p72;
	mul.f32 	%f774, %f773, 0f3FB8AA3B;
	cvt.rzi.f32.f32	%f775, %f774;
	mov.f32 	%f776, 0fBF317200;
	fma.rn.f32 	%f777, %f775, %f776, %f773;
	mov.f32 	%f778, 0fB5BFBE8E;
	fma.rn.f32 	%f779, %f775, %f778, %f777;
	mul.f32 	%f780, %f779, 0f3FB8AA3B;
	ex2.approx.ftz.f32 	%f781, %f780;
	add.f32 	%f782, %f775, 0f00000000;
	ex2.approx.f32 	%f783, %f782;
	mul.f32 	%f784, %f781, %f783;
	setp.lt.f32	%p73, %f773, 0fC2D20000;
	selp.f32	%f785, 0f00000000, %f784, %p73;
	setp.gt.f32	%p74, %f773, 0f42D20000;
	selp.f32	%f1309, 0f7F800000, %f785, %p74;
	setp.eq.f32	%p75, %f1309, 0f7F800000;
	@%p75 bra 	BB0_81;

	fma.rn.f32 	%f1309, %f1309, %f281, %f1309;

BB0_81:
	setp.lt.f32	%p76, %f278, 0f00000000;
	setp.eq.f32	%p77, %f279, 0f3F800000;
	and.pred  	%p3, %p76, %p77;
	mov.b32 	 %r270, %f1309;
	xor.b32  	%r271, %r270, -2147483648;
	mov.b32 	 %f786, %r271;
	selp.f32	%f1311, %f786, %f1309, %p3;
	setp.eq.f32	%p78, %f278, 0f00000000;
	@%p78 bra 	BB0_84;
	bra.uni 	BB0_82;

BB0_84:
	add.f32 	%f789, %f278, %f278;
	selp.f32	%f1311, %f789, 0f00000000, %p77;
	bra.uni 	BB0_85;

BB0_149:
	mov.u64 	%rd237, image_HDR;
	cvta.global.u64 	%rd232, %rd237;
	mov.u32 	%r357, 8;
	mov.u64 	%rd236, 0;
	// inline asm
	call (%rd231), _rt_buffer_get_64, (%rd232, %r103, %r357, %rd20, %rd21, %rd236, %rd236);
	// inline asm
	mov.f32 	%f1128, 0f00000000;
	// inline asm
	{  cvt.rn.f16.f32 %rs105, %f1128;}

	// inline asm
	mov.u16 	%rs106, 0;
	st.v4.u16 	[%rd231], {%rs105, %rs105, %rs105, %rs106};

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

	mov.u64 	%rd250, image_RNM0;
	cvta.global.u64 	%rd239, %rd250;
	mov.u32 	%r362, 8;
	mov.u64 	%rd249, 0;
	// inline asm
	call (%rd238), _rt_buffer_get_64, (%rd239, %r103, %r362, %rd20, %rd21, %rd249, %rd249);
	// inline asm
	ld.v4.u16 	{%rs113, %rs114, %rs115, %rs116}, [%rd238];
	// inline asm
	{  cvt.f32.f16 %f1129, %rs113;}

	// inline asm
	// inline asm
	{  cvt.f32.f16 %f1130, %rs114;}

	// inline asm
	// inline asm
	{  cvt.f32.f16 %f1131, %rs115;}

	// inline asm
	// inline asm
	call (%rd244), _rt_buffer_get_64, (%rd239, %r103, %r362, %rd20, %rd21, %rd249, %rd249);
	// inline asm
	add.f32 	%f1132, %f1129, 0f00000000;
	add.f32 	%f1133, %f1130, 0f00000000;
	add.f32 	%f1134, %f1131, 0f00000000;
	// inline asm
	{  cvt.rn.f16.f32 %rs112, %f1134;}

	// inline asm
	// inline asm
	{  cvt.rn.f16.f32 %rs111, %f1133;}

	// inline asm
	// inline asm
	{  cvt.rn.f16.f32 %rs110, %f1132;}

	// inline asm
	mov.u16 	%rs117, 0;
	st.v4.u16 	[%rd244], {%rs110, %rs111, %rs112, %rs117};
	bra.uni 	BB0_153;

BB0_152:
	mov.u64 	%rd257, image_RNM0;
	cvta.global.u64 	%rd252, %rd257;
	mov.u32 	%r364, 8;
	mov.u64 	%rd256, 0;
	// inline asm
	call (%rd251), _rt_buffer_get_64, (%rd252, %r103, %r364, %rd20, %rd21, %rd256, %rd256);
	// inline asm
	mov.f32 	%f1135, 0f00000000;
	// inline asm
	{  cvt.rn.f16.f32 %rs118, %f1135;}

	// inline asm
	mov.u16 	%rs119, 0;
	st.v4.u16 	[%rd251], {%rs118, %rs118, %rs118, %rs119};

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

	mov.u64 	%rd270, image_RNM1;
	cvta.global.u64 	%rd259, %rd270;
	mov.u32 	%r369, 8;
	mov.u64 	%rd269, 0;
	// inline asm
	call (%rd258), _rt_buffer_get_64, (%rd259, %r103, %r369, %rd20, %rd21, %rd269, %rd269);
	// inline asm
	ld.v4.u16 	{%rs126, %rs127, %rs128, %rs129}, [%rd258];
	// inline asm
	{  cvt.f32.f16 %f1136, %rs126;}

	// inline asm
	// inline asm
	{  cvt.f32.f16 %f1137, %rs127;}

	// inline asm
	// inline asm
	{  cvt.f32.f16 %f1138, %rs128;}

	// inline asm
	// inline asm
	call (%rd264), _rt_buffer_get_64, (%rd259, %r103, %r369, %rd20, %rd21, %rd269, %rd269);
	// inline asm
	add.f32 	%f1139, %f1136, 0f00000000;
	add.f32 	%f1140, %f1137, 0f00000000;
	add.f32 	%f1141, %f1138, 0f00000000;
	// inline asm
	{  cvt.rn.f16.f32 %rs125, %f1141;}

	// inline asm
	// inline asm
	{  cvt.rn.f16.f32 %rs124, %f1140;}

	// inline asm
	// inline asm
	{  cvt.rn.f16.f32 %rs123, %f1139;}

	// inline asm
	mov.u16 	%rs130, 0;
	st.v4.u16 	[%rd264], {%rs123, %rs124, %rs125, %rs130};
	bra.uni 	BB0_156;

BB0_155:
	mov.u64 	%rd277, image_RNM1;
	cvta.global.u64 	%rd272, %rd277;
	mov.u32 	%r371, 8;
	mov.u64 	%rd276, 0;
	// inline asm
	call (%rd271), _rt_buffer_get_64, (%rd272, %r103, %r371, %rd20, %rd21, %rd276, %rd276);
	// inline asm
	mov.f32 	%f1142, 0f00000000;
	// inline asm
	{  cvt.rn.f16.f32 %rs131, %f1142;}

	// inline asm
	mov.u16 	%rs132, 0;
	st.v4.u16 	[%rd271], {%rs131, %rs131, %rs131, %rs132};

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

	mov.u64 	%rd290, image_RNM2;
	cvta.global.u64 	%rd279, %rd290;
	mov.u32 	%r376, 8;
	mov.u64 	%rd289, 0;
	// inline asm
	call (%rd278), _rt_buffer_get_64, (%rd279, %r103, %r376, %rd20, %rd21, %rd289, %rd289);
	// inline asm
	ld.v4.u16 	{%rs139, %rs140, %rs141, %rs142}, [%rd278];
	// inline asm
	{  cvt.f32.f16 %f1143, %rs139;}

	// inline asm
	// inline asm
	{  cvt.f32.f16 %f1144, %rs140;}

	// inline asm
	// inline asm
	{  cvt.f32.f16 %f1145, %rs141;}

	// inline asm
	// inline asm
	call (%rd284), _rt_buffer_get_64, (%rd279, %r103, %r376, %rd20, %rd21, %rd289, %rd289);
	// inline asm
	add.f32 	%f1146, %f1143, 0f00000000;
	add.f32 	%f1147, %f1144, 0f00000000;
	add.f32 	%f1148, %f1145, 0f00000000;
	// inline asm
	{  cvt.rn.f16.f32 %rs138, %f1148;}

	// inline asm
	// inline asm
	{  cvt.rn.f16.f32 %rs137, %f1147;}

	// inline asm
	// inline asm
	{  cvt.rn.f16.f32 %rs136, %f1146;}

	// inline asm
	mov.u16 	%rs143, 0;
	st.v4.u16 	[%rd284], {%rs136, %rs137, %rs138, %rs143};
	bra.uni 	BB0_159;

BB0_158:
	mov.u64 	%rd297, image_RNM2;
	cvta.global.u64 	%rd292, %rd297;
	mov.u32 	%r378, 8;
	mov.u64 	%rd296, 0;
	// inline asm
	call (%rd291), _rt_buffer_get_64, (%rd292, %r103, %r378, %rd20, %rd21, %rd296, %rd296);
	// inline asm
	mov.f32 	%f1149, 0f00000000;
	// inline asm
	{  cvt.rn.f16.f32 %rs144, %f1149;}

	// inline asm
	mov.u16 	%rs145, 0;
	st.v4.u16 	[%rd291], {%rs144, %rs144, %rs144, %rs145};

BB0_159:
	ld.global.u32 	%r379, [additive];
	setp.eq.s32	%p148, %r379, 0;
	@%p148 bra 	BB0_161;

	mov.u64 	%rd310, image_RNM3;
	cvta.global.u64 	%rd299, %rd310;
	mov.u32 	%r383, 8;
	mov.u64 	%rd309, 0;
	// inline asm
	call (%rd298), _rt_buffer_get_64, (%rd299, %r103, %r383, %rd20, %rd21, %rd309, %rd309);
	// inline asm
	ld.v4.u16 	{%rs152, %rs153, %rs154, %rs155}, [%rd298];
	// inline asm
	{  cvt.f32.f16 %f1150, %rs152;}

	// inline asm
	// inline asm
	{  cvt.f32.f16 %f1151, %rs153;}

	// inline asm
	// inline asm
	{  cvt.f32.f16 %f1152, %rs154;}

	// inline asm
	// inline asm
	call (%rd304), _rt_buffer_get_64, (%rd299, %r103, %r383, %rd20, %rd21, %rd309, %rd309);
	// inline asm
	add.f32 	%f1153, %f1150, 0f00000000;
	add.f32 	%f1154, %f1151, 0f00000000;
	add.f32 	%f1155, %f1152, 0f00000000;
	// inline asm
	{  cvt.rn.f16.f32 %rs151, %f1155;}

	// inline asm
	// inline asm
	{  cvt.rn.f16.f32 %rs150, %f1154;}

	// inline asm
	// inline asm
	{  cvt.rn.f16.f32 %rs149, %f1153;}

	// inline asm
	mov.u16 	%rs156, 0;
	st.v4.u16 	[%rd304], {%rs149, %rs150, %rs151, %rs156};
	bra.uni 	BB0_162;

BB0_161:
	mov.u64 	%rd317, image_RNM3;
	cvta.global.u64 	%rd312, %rd317;
	mov.u32 	%r385, 8;
	mov.u64 	%rd316, 0;
	// inline asm
	call (%rd311), _rt_buffer_get_64, (%rd312, %r103, %r385, %rd20, %rd21, %rd316, %rd316);
	// inline asm
	mov.f32 	%f1156, 0f00000000;
	// inline asm
	{  cvt.rn.f16.f32 %rs157, %f1156;}

	// inline asm
	mov.u16 	%rs158, 0;
	st.v4.u16 	[%rd311], {%rs157, %rs157, %rs157, %rs158};
	bra.uni 	BB0_162;

BB0_82:
	setp.geu.f32	%p79, %f278, 0f00000000;
	@%p79 bra 	BB0_85;

	cvt.rzi.f32.f32	%f788, %f760;
	setp.neu.f32	%p80, %f788, 0f3EE8BA2E;
	selp.f32	%f1311, 0f7FFFFFFF, %f1311, %p80;

BB0_85:
	add.f32 	%f790, %f280, 0f3EE8BA2E;
	mov.b32 	 %r272, %f790;
	setp.lt.s32	%p82, %r272, 2139095040;
	@%p82 bra 	BB0_90;

	setp.gtu.f32	%p83, %f280, 0f7F800000;
	@%p83 bra 	BB0_89;
	bra.uni 	BB0_87;

BB0_89:
	add.f32 	%f1311, %f278, 0f3EE8BA2E;
	bra.uni 	BB0_90;

BB0_87:
	setp.neu.f32	%p84, %f280, 0f7F800000;
	@%p84 bra 	BB0_90;

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

BB0_90:
	mul.f32 	%f791, %f1311, 0f437F0000;
	setp.eq.f32	%p85, %f278, 0f3F800000;
	selp.f32	%f792, 0f437F0000, %f791, %p85;
	cvt.rzi.u32.f32	%r273, %f792;
	cvt.u16.u32	%rs15, %r273;
	mov.u16 	%rs16, 255;
	st.v2.u8 	[%rd90], {%rs15, %rs16};
	ld.global.u32 	%r419, [imageEnabled];

BB0_91:
	ld.global.f32 	%f793, [lightColor];
	mul.f32 	%f292, %f1305, %f793;
	ld.global.f32 	%f794, [lightColor+4];
	mul.f32 	%f293, %f1306, %f794;
	ld.global.f32 	%f795, [lightColor+8];
	mul.f32 	%f294, %f1307, %f795;
	and.b32  	%r274, %r419, 1;
	setp.eq.b32	%p86, %r274, 1;
	@!%p86 bra 	BB0_126;
	bra.uni 	BB0_92;

BB0_92:
	mov.f32 	%f798, 0f3E666666;
	cvt.rzi.f32.f32	%f799, %f798;
	fma.rn.f32 	%f800, %f799, 0fC0000000, 0f3EE66666;
	abs.f32 	%f295, %f800;
	abs.f32 	%f296, %f292;
	setp.lt.f32	%p87, %f296, 0f00800000;
	mul.f32 	%f801, %f296, 0f4B800000;
	selp.f32	%f802, 0fC3170000, 0fC2FE0000, %p87;
	selp.f32	%f803, %f801, %f296, %p87;
	mov.b32 	 %r275, %f803;
	and.b32  	%r276, %r275, 8388607;
	or.b32  	%r277, %r276, 1065353216;
	mov.b32 	 %f804, %r277;
	shr.u32 	%r278, %r275, 23;
	cvt.rn.f32.u32	%f805, %r278;
	add.f32 	%f806, %f802, %f805;
	setp.gt.f32	%p88, %f804, 0f3FB504F3;
	mul.f32 	%f807, %f804, 0f3F000000;
	add.f32 	%f808, %f806, 0f3F800000;
	selp.f32	%f809, %f807, %f804, %p88;
	selp.f32	%f810, %f808, %f806, %p88;
	add.f32 	%f811, %f809, 0fBF800000;
	add.f32 	%f797, %f809, 0f3F800000;
	// inline asm
	rcp.approx.ftz.f32 %f796,%f797;
	// inline asm
	add.f32 	%f812, %f811, %f811;
	mul.f32 	%f813, %f796, %f812;
	mul.f32 	%f814, %f813, %f813;
	mov.f32 	%f815, 0f3C4CAF63;
	mov.f32 	%f816, 0f3B18F0FE;
	fma.rn.f32 	%f817, %f816, %f814, %f815;
	mov.f32 	%f818, 0f3DAAAABD;
	fma.rn.f32 	%f819, %f817, %f814, %f818;
	mul.rn.f32 	%f820, %f819, %f814;
	mul.rn.f32 	%f821, %f820, %f813;
	sub.f32 	%f822, %f811, %f813;
	neg.f32 	%f823, %f813;
	add.f32 	%f824, %f822, %f822;
	fma.rn.f32 	%f825, %f823, %f811, %f824;
	mul.rn.f32 	%f826, %f796, %f825;
	add.f32 	%f827, %f821, %f813;
	sub.f32 	%f828, %f813, %f827;
	add.f32 	%f829, %f821, %f828;
	add.f32 	%f830, %f826, %f829;
	add.f32 	%f831, %f827, %f830;
	sub.f32 	%f832, %f827, %f831;
	add.f32 	%f833, %f830, %f832;
	mov.f32 	%f834, 0f3F317200;
	mul.rn.f32 	%f835, %f810, %f834;
	mov.f32 	%f836, 0f35BFBE8E;
	mul.rn.f32 	%f837, %f810, %f836;
	add.f32 	%f838, %f835, %f831;
	sub.f32 	%f839, %f835, %f838;
	add.f32 	%f840, %f831, %f839;
	add.f32 	%f841, %f833, %f840;
	add.f32 	%f842, %f837, %f841;
	add.f32 	%f843, %f838, %f842;
	sub.f32 	%f844, %f838, %f843;
	add.f32 	%f845, %f842, %f844;
	mov.f32 	%f846, 0f3EE66666;
	mul.rn.f32 	%f847, %f846, %f843;
	neg.f32 	%f848, %f847;
	fma.rn.f32 	%f849, %f846, %f843, %f848;
	fma.rn.f32 	%f850, %f846, %f845, %f849;
	mov.f32 	%f851, 0f00000000;
	fma.rn.f32 	%f852, %f851, %f843, %f850;
	add.rn.f32 	%f853, %f847, %f852;
	neg.f32 	%f854, %f853;
	add.rn.f32 	%f855, %f847, %f854;
	add.rn.f32 	%f856, %f855, %f852;
	mov.b32 	 %r279, %f853;
	setp.eq.s32	%p89, %r279, 1118925336;
	add.s32 	%r280, %r279, -1;
	mov.b32 	 %f857, %r280;
	add.f32 	%f858, %f856, 0f37000000;
	selp.f32	%f859, %f857, %f853, %p89;
	selp.f32	%f297, %f858, %f856, %p89;
	mul.f32 	%f860, %f859, 0f3FB8AA3B;
	cvt.rzi.f32.f32	%f861, %f860;
	mov.f32 	%f862, 0fBF317200;
	fma.rn.f32 	%f863, %f861, %f862, %f859;
	mov.f32 	%f864, 0fB5BFBE8E;
	fma.rn.f32 	%f865, %f861, %f864, %f863;
	mul.f32 	%f866, %f865, 0f3FB8AA3B;
	ex2.approx.ftz.f32 	%f867, %f866;
	add.f32 	%f868, %f861, 0f00000000;
	ex2.approx.f32 	%f869, %f868;
	mul.f32 	%f870, %f867, %f869;
	setp.lt.f32	%p90, %f859, 0fC2D20000;
	selp.f32	%f871, 0f00000000, %f870, %p90;
	setp.gt.f32	%p91, %f859, 0f42D20000;
	selp.f32	%f1312, 0f7F800000, %f871, %p91;
	setp.eq.f32	%p92, %f1312, 0f7F800000;
	@%p92 bra 	BB0_94;

	fma.rn.f32 	%f1312, %f1312, %f297, %f1312;

BB0_94:
	setp.lt.f32	%p93, %f292, 0f00000000;
	setp.eq.f32	%p94, %f295, 0f3F800000;
	and.pred  	%p4, %p93, %p94;
	mov.b32 	 %r281, %f1312;
	xor.b32  	%r282, %r281, -2147483648;
	mov.b32 	 %f872, %r282;
	selp.f32	%f1314, %f872, %f1312, %p4;
	setp.eq.f32	%p95, %f292, 0f00000000;
	@%p95 bra 	BB0_97;
	bra.uni 	BB0_95;

BB0_97:
	add.f32 	%f875, %f292, %f292;
	selp.f32	%f1314, %f875, 0f00000000, %p94;
	bra.uni 	BB0_98;

BB0_95:
	setp.geu.f32	%p96, %f292, 0f00000000;
	@%p96 bra 	BB0_98;

	cvt.rzi.f32.f32	%f874, %f846;
	setp.neu.f32	%p97, %f874, 0f3EE66666;
	selp.f32	%f1314, 0f7FFFFFFF, %f1314, %p97;

BB0_98:
	add.f32 	%f876, %f296, 0f3EE66666;
	mov.b32 	 %r283, %f876;
	setp.lt.s32	%p99, %r283, 2139095040;
	@%p99 bra 	BB0_103;

	setp.gtu.f32	%p100, %f296, 0f7F800000;
	@%p100 bra 	BB0_102;
	bra.uni 	BB0_100;

BB0_102:
	add.f32 	%f1314, %f292, 0f3EE66666;
	bra.uni 	BB0_103;

BB0_100:
	setp.neu.f32	%p101, %f296, 0f7F800000;
	@%p101 bra 	BB0_103;

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

BB0_103:
	setp.eq.f32	%p102, %f292, 0f3F800000;
	selp.f32	%f308, 0f3F800000, %f1314, %p102;
	abs.f32 	%f309, %f293;
	setp.lt.f32	%p103, %f309, 0f00800000;
	mul.f32 	%f879, %f309, 0f4B800000;
	selp.f32	%f880, 0fC3170000, 0fC2FE0000, %p103;
	selp.f32	%f881, %f879, %f309, %p103;
	mov.b32 	 %r284, %f881;
	and.b32  	%r285, %r284, 8388607;
	or.b32  	%r286, %r285, 1065353216;
	mov.b32 	 %f882, %r286;
	shr.u32 	%r287, %r284, 23;
	cvt.rn.f32.u32	%f883, %r287;
	add.f32 	%f884, %f880, %f883;
	setp.gt.f32	%p104, %f882, 0f3FB504F3;
	mul.f32 	%f885, %f882, 0f3F000000;
	add.f32 	%f886, %f884, 0f3F800000;
	selp.f32	%f887, %f885, %f882, %p104;
	selp.f32	%f888, %f886, %f884, %p104;
	add.f32 	%f889, %f887, 0fBF800000;
	add.f32 	%f878, %f887, 0f3F800000;
	// inline asm
	rcp.approx.ftz.f32 %f877,%f878;
	// inline asm
	add.f32 	%f890, %f889, %f889;
	mul.f32 	%f891, %f877, %f890;
	mul.f32 	%f892, %f891, %f891;
	fma.rn.f32 	%f895, %f816, %f892, %f815;
	fma.rn.f32 	%f897, %f895, %f892, %f818;
	mul.rn.f32 	%f898, %f897, %f892;
	mul.rn.f32 	%f899, %f898, %f891;
	sub.f32 	%f900, %f889, %f891;
	neg.f32 	%f901, %f891;
	add.f32 	%f902, %f900, %f900;
	fma.rn.f32 	%f903, %f901, %f889, %f902;
	mul.rn.f32 	%f904, %f877, %f903;
	add.f32 	%f905, %f899, %f891;
	sub.f32 	%f906, %f891, %f905;
	add.f32 	%f907, %f899, %f906;
	add.f32 	%f908, %f904, %f907;
	add.f32 	%f909, %f905, %f908;
	sub.f32 	%f910, %f905, %f909;
	add.f32 	%f911, %f908, %f910;
	mul.rn.f32 	%f913, %f888, %f834;
	mul.rn.f32 	%f915, %f888, %f836;
	add.f32 	%f916, %f913, %f909;
	sub.f32 	%f917, %f913, %f916;
	add.f32 	%f918, %f909, %f917;
	add.f32 	%f919, %f911, %f918;
	add.f32 	%f920, %f915, %f919;
	add.f32 	%f921, %f916, %f920;
	sub.f32 	%f922, %f916, %f921;
	add.f32 	%f923, %f920, %f922;
	mul.rn.f32 	%f925, %f846, %f921;
	neg.f32 	%f926, %f925;
	fma.rn.f32 	%f927, %f846, %f921, %f926;
	fma.rn.f32 	%f928, %f846, %f923, %f927;
	fma.rn.f32 	%f930, %f851, %f921, %f928;
	add.rn.f32 	%f931, %f925, %f930;
	neg.f32 	%f932, %f931;
	add.rn.f32 	%f933, %f925, %f932;
	add.rn.f32 	%f934, %f933, %f930;
	mov.b32 	 %r288, %f931;
	setp.eq.s32	%p105, %r288, 1118925336;
	add.s32 	%r289, %r288, -1;
	mov.b32 	 %f935, %r289;
	add.f32 	%f936, %f934, 0f37000000;
	selp.f32	%f937, %f935, %f931, %p105;
	selp.f32	%f310, %f936, %f934, %p105;
	mul.f32 	%f938, %f937, 0f3FB8AA3B;
	cvt.rzi.f32.f32	%f939, %f938;
	fma.rn.f32 	%f941, %f939, %f862, %f937;
	fma.rn.f32 	%f943, %f939, %f864, %f941;
	mul.f32 	%f944, %f943, 0f3FB8AA3B;
	ex2.approx.ftz.f32 	%f945, %f944;
	add.f32 	%f946, %f939, 0f00000000;
	ex2.approx.f32 	%f947, %f946;
	mul.f32 	%f948, %f945, %f947;
	setp.lt.f32	%p106, %f937, 0fC2D20000;
	selp.f32	%f949, 0f00000000, %f948, %p106;
	setp.gt.f32	%p107, %f937, 0f42D20000;
	selp.f32	%f1315, 0f7F800000, %f949, %p107;
	setp.eq.f32	%p108, %f1315, 0f7F800000;
	@%p108 bra 	BB0_105;

	fma.rn.f32 	%f1315, %f1315, %f310, %f1315;

BB0_105:
	setp.lt.f32	%p109, %f293, 0f00000000;
	and.pred  	%p5, %p109, %p94;
	mov.b32 	 %r290, %f1315;
	xor.b32  	%r291, %r290, -2147483648;
	mov.b32 	 %f950, %r291;
	selp.f32	%f1317, %f950, %f1315, %p5;
	setp.eq.f32	%p111, %f293, 0f00000000;
	@%p111 bra 	BB0_108;
	bra.uni 	BB0_106;

BB0_108:
	add.f32 	%f953, %f293, %f293;
	selp.f32	%f1317, %f953, 0f00000000, %p94;
	bra.uni 	BB0_109;

BB0_106:
	setp.geu.f32	%p112, %f293, 0f00000000;
	@%p112 bra 	BB0_109;

	cvt.rzi.f32.f32	%f952, %f846;
	setp.neu.f32	%p113, %f952, 0f3EE66666;
	selp.f32	%f1317, 0f7FFFFFFF, %f1317, %p113;

BB0_109:
	add.f32 	%f954, %f309, 0f3EE66666;
	mov.b32 	 %r292, %f954;
	setp.lt.s32	%p115, %r292, 2139095040;
	@%p115 bra 	BB0_114;

	setp.gtu.f32	%p116, %f309, 0f7F800000;
	@%p116 bra 	BB0_113;
	bra.uni 	BB0_111;

BB0_113:
	add.f32 	%f1317, %f293, 0f3EE66666;
	bra.uni 	BB0_114;

BB0_111:
	setp.neu.f32	%p117, %f309, 0f7F800000;
	@%p117 bra 	BB0_114;

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

BB0_114:
	setp.eq.f32	%p118, %f293, 0f3F800000;
	selp.f32	%f321, 0f3F800000, %f1317, %p118;
	abs.f32 	%f322, %f294;
	setp.lt.f32	%p119, %f322, 0f00800000;
	mul.f32 	%f957, %f322, 0f4B800000;
	selp.f32	%f958, 0fC3170000, 0fC2FE0000, %p119;
	selp.f32	%f959, %f957, %f322, %p119;
	mov.b32 	 %r293, %f959;
	and.b32  	%r294, %r293, 8388607;
	or.b32  	%r295, %r294, 1065353216;
	mov.b32 	 %f960, %r295;
	shr.u32 	%r296, %r293, 23;
	cvt.rn.f32.u32	%f961, %r296;
	add.f32 	%f962, %f958, %f961;
	setp.gt.f32	%p120, %f960, 0f3FB504F3;
	mul.f32 	%f963, %f960, 0f3F000000;
	add.f32 	%f964, %f962, 0f3F800000;
	selp.f32	%f965, %f963, %f960, %p120;
	selp.f32	%f966, %f964, %f962, %p120;
	add.f32 	%f967, %f965, 0fBF800000;
	add.f32 	%f956, %f965, 0f3F800000;
	// inline asm
	rcp.approx.ftz.f32 %f955,%f956;
	// inline asm
	add.f32 	%f968, %f967, %f967;
	mul.f32 	%f969, %f955, %f968;
	mul.f32 	%f970, %f969, %f969;
	fma.rn.f32 	%f973, %f816, %f970, %f815;
	fma.rn.f32 	%f975, %f973, %f970, %f818;
	mul.rn.f32 	%f976, %f975, %f970;
	mul.rn.f32 	%f977, %f976, %f969;
	sub.f32 	%f978, %f967, %f969;
	neg.f32 	%f979, %f969;
	add.f32 	%f980, %f978, %f978;
	fma.rn.f32 	%f981, %f979, %f967, %f980;
	mul.rn.f32 	%f982, %f955, %f981;
	add.f32 	%f983, %f977, %f969;
	sub.f32 	%f984, %f969, %f983;
	add.f32 	%f985, %f977, %f984;
	add.f32 	%f986, %f982, %f985;
	add.f32 	%f987, %f983, %f986;
	sub.f32 	%f988, %f983, %f987;
	add.f32 	%f989, %f986, %f988;
	mul.rn.f32 	%f991, %f966, %f834;
	mul.rn.f32 	%f993, %f966, %f836;
	add.f32 	%f994, %f991, %f987;
	sub.f32 	%f995, %f991, %f994;
	add.f32 	%f996, %f987, %f995;
	add.f32 	%f997, %f989, %f996;
	add.f32 	%f998, %f993, %f997;
	add.f32 	%f999, %f994, %f998;
	sub.f32 	%f1000, %f994, %f999;
	add.f32 	%f1001, %f998, %f1000;
	mul.rn.f32 	%f1003, %f846, %f999;
	neg.f32 	%f1004, %f1003;
	fma.rn.f32 	%f1005, %f846, %f999, %f1004;
	fma.rn.f32 	%f1006, %f846, %f1001, %f1005;
	fma.rn.f32 	%f1008, %f851, %f999, %f1006;
	add.rn.f32 	%f1009, %f1003, %f1008;
	neg.f32 	%f1010, %f1009;
	add.rn.f32 	%f1011, %f1003, %f1010;
	add.rn.f32 	%f1012, %f1011, %f1008;
	mov.b32 	 %r297, %f1009;
	setp.eq.s32	%p121, %r297, 1118925336;
	add.s32 	%r298, %r297, -1;
	mov.b32 	 %f1013, %r298;
	add.f32 	%f1014, %f1012, 0f37000000;
	selp.f32	%f1015, %f1013, %f1009, %p121;
	selp.f32	%f323, %f1014, %f1012, %p121;
	mul.f32 	%f1016, %f1015, 0f3FB8AA3B;
	cvt.rzi.f32.f32	%f1017, %f1016;
	fma.rn.f32 	%f1019, %f1017, %f862, %f1015;
	fma.rn.f32 	%f1021, %f1017, %f864, %f1019;
	mul.f32 	%f1022, %f1021, 0f3FB8AA3B;
	ex2.approx.ftz.f32 	%f1023, %f1022;
	add.f32 	%f1024, %f1017, 0f00000000;
	ex2.approx.f32 	%f1025, %f1024;
	mul.f32 	%f1026, %f1023, %f1025;
	setp.lt.f32	%p122, %f1015, 0fC2D20000;
	selp.f32	%f1027, 0f00000000, %f1026, %p122;
	setp.gt.f32	%p123, %f1015, 0f42D20000;
	selp.f32	%f1318, 0f7F800000, %f1027, %p123;
	setp.eq.f32	%p124, %f1318, 0f7F800000;
	@%p124 bra 	BB0_116;

	fma.rn.f32 	%f1318, %f1318, %f323, %f1318;

BB0_116:
	setp.lt.f32	%p125, %f294, 0f00000000;
	and.pred  	%p6, %p125, %p94;
	mov.b32 	 %r299, %f1318;
	xor.b32  	%r300, %r299, -2147483648;
	mov.b32 	 %f1028, %r300;
	selp.f32	%f1320, %f1028, %f1318, %p6;
	setp.eq.f32	%p127, %f294, 0f00000000;
	@%p127 bra 	BB0_119;
	bra.uni 	BB0_117;

BB0_119:
	add.f32 	%f1031, %f294, %f294;
	selp.f32	%f1320, %f1031, 0f00000000, %p94;
	bra.uni 	BB0_120;

BB0_117:
	setp.geu.f32	%p128, %f294, 0f00000000;
	@%p128 bra 	BB0_120;

	cvt.rzi.f32.f32	%f1030, %f846;
	setp.neu.f32	%p129, %f1030, 0f3EE66666;
	selp.f32	%f1320, 0f7FFFFFFF, %f1320, %p129;

BB0_120:
	add.f32 	%f1032, %f322, 0f3EE66666;
	mov.b32 	 %r301, %f1032;
	setp.lt.s32	%p131, %r301, 2139095040;
	@%p131 bra 	BB0_125;

	setp.gtu.f32	%p132, %f322, 0f7F800000;
	@%p132 bra 	BB0_124;
	bra.uni 	BB0_122;

BB0_124:
	add.f32 	%f1320, %f294, 0f3EE66666;
	bra.uni 	BB0_125;

BB0_122:
	setp.neu.f32	%p133, %f322, 0f7F800000;
	@%p133 bra 	BB0_125;

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

BB0_125:
	setp.eq.f32	%p134, %f294, 0f3F800000;
	selp.f32	%f1033, 0f3F800000, %f1320, %p134;
	cvt.u64.u32	%rd100, %r6;
	cvt.u64.u32	%rd99, %r5;
	mov.u64 	%rd103, image;
	cvta.global.u64 	%rd98, %rd103;
	// inline asm
	call (%rd97), _rt_buffer_get_64, (%rd98, %r103, %r104, %rd99, %rd100, %rd27, %rd27);
	// inline asm
	cvt.sat.f32.f32	%f1034, %f1033;
	mul.f32 	%f1035, %f1034, 0f437FFD71;
	cvt.rzi.u32.f32	%r304, %f1035;
	cvt.sat.f32.f32	%f1036, %f321;
	mul.f32 	%f1037, %f1036, 0f437FFD71;
	cvt.rzi.u32.f32	%r305, %f1037;
	cvt.sat.f32.f32	%f1038, %f308;
	mul.f32 	%f1039, %f1038, 0f437FFD71;
	cvt.rzi.u32.f32	%r306, %f1039;
	cvt.u16.u32	%rs17, %r304;
	cvt.u16.u32	%rs18, %r306;
	cvt.u16.u32	%rs19, %r305;
	mov.u16 	%rs20, 255;
	st.v4.u8 	[%rd97], {%rs17, %rs19, %rs18, %rs20};
	ld.global.u32 	%r419, [imageEnabled];

BB0_126:
	cvt.u64.u32	%rd18, %r5;
	cvt.u64.u32	%rd19, %r6;
	and.b32  	%r307, %r419, 4;
	setp.eq.s32	%p135, %r307, 0;
	@%p135 bra 	BB0_130;

	ld.global.u32 	%r308, [additive];
	setp.eq.s32	%p136, %r308, 0;
	// inline asm
	{  cvt.rn.f16.f32 %rs21, %f564;}

	// inline asm
	@%p136 bra 	BB0_129;

	mov.u64 	%rd116, image_HDR;
	cvta.global.u64 	%rd105, %rd116;
	mov.u32 	%r312, 8;
	// inline asm
	call (%rd104), _rt_buffer_get_64, (%rd105, %r103, %r312, %rd18, %rd19, %rd27, %rd27);
	// inline asm
	ld.v4.u16 	{%rs28, %rs29, %rs30, %rs31}, [%rd104];
	// inline asm
	{  cvt.f32.f16 %f1041, %rs28;}

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

	// inline asm
	// inline asm
	{  cvt.f32.f16 %f1043, %rs30;}

	// inline asm
	// inline asm
	call (%rd110), _rt_buffer_get_64, (%rd105, %r103, %r312, %rd18, %rd19, %rd27, %rd27);
	// inline asm
	add.f32 	%f1044, %f292, %f1041;
	add.f32 	%f1045, %f293, %f1042;
	add.f32 	%f1046, %f294, %f1043;
	// inline asm
	{  cvt.rn.f16.f32 %rs27, %f1046;}

	// inline asm
	// inline asm
	{  cvt.rn.f16.f32 %rs26, %f1045;}

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

	// inline asm
	st.v4.u16 	[%rd110], {%rs25, %rs26, %rs27, %rs21};
	bra.uni 	BB0_130;

BB0_129:
	mov.u64 	%rd123, image_HDR;
	cvta.global.u64 	%rd118, %rd123;
	mov.u32 	%r314, 8;
	// inline asm
	call (%rd117), _rt_buffer_get_64, (%rd118, %r103, %r314, %rd18, %rd19, %rd27, %rd27);
	// inline asm
	// inline asm
	{  cvt.rn.f16.f32 %rs34, %f294;}

	// inline asm
	// inline asm
	{  cvt.rn.f16.f32 %rs33, %f293;}

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

	// inline asm
	st.v4.u16 	[%rd117], {%rs32, %rs33, %rs34, %rs21};

BB0_130:
	selp.f32	%f1051, 0f3F000000, 0f3E800000, %p149;
	mul.f32 	%f1052, %f1051, %f1302;
	mul.f32 	%f1053, %f1051, %f1303;
	mul.f32 	%f1054, %f1051, %f1304;
	mul.f32 	%f1055, %f1051, %f1299;
	mul.f32 	%f1056, %f1051, %f1300;
	mul.f32 	%f1057, %f1051, %f1301;
	mul.f32 	%f1058, %f1051, %f1296;
	mul.f32 	%f1059, %f1051, %f1297;
	mul.f32 	%f1060, %f1051, %f1298;
	mul.f32 	%f1061, %f1051, %f1293;
	mul.f32 	%f1062, %f1051, %f1294;
	mul.f32 	%f1063, %f1051, %f1295;
	ld.global.f32 	%f1064, [lightColor];
	mul.f32 	%f334, %f1052, %f1064;
	ld.global.f32 	%f1065, [lightColor+4];
	mul.f32 	%f335, %f1053, %f1065;
	ld.global.f32 	%f1066, [lightColor+8];
	mul.f32 	%f336, %f1054, %f1066;
	mul.f32 	%f337, %f1055, %f1064;
	mul.f32 	%f338, %f1056, %f1065;
	mul.f32 	%f339, %f1057, %f1066;
	mul.f32 	%f340, %f1058, %f1064;
	mul.f32 	%f341, %f1059, %f1065;
	mul.f32 	%f342, %f1060, %f1066;
	mul.f32 	%f343, %f1061, %f1064;
	mul.f32 	%f344, %f1062, %f1065;
	mul.f32 	%f345, %f1063, %f1066;
	ld.global.u32 	%r315, [additive];
	setp.eq.s32	%p137, %r315, 0;
	// inline asm
	{  cvt.rn.f16.f32 %rs35, %f564;}

	// inline asm
	@%p137 bra 	BB0_132;

	mov.u64 	%rd136, image_RNM0;
	cvta.global.u64 	%rd125, %rd136;
	mov.u32 	%r319, 8;
	// inline asm
	call (%rd124), _rt_buffer_get_64, (%rd125, %r103, %r319, %rd18, %rd19, %rd27, %rd27);
	// inline asm
	ld.v4.u16 	{%rs42, %rs43, %rs44, %rs45}, [%rd124];
	// inline asm
	{  cvt.f32.f16 %f1067, %rs42;}

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

	// inline asm
	// inline asm
	{  cvt.f32.f16 %f1069, %rs44;}

	// inline asm
	// inline asm
	call (%rd130), _rt_buffer_get_64, (%rd125, %r103, %r319, %rd18, %rd19, %rd27, %rd27);
	// inline asm
	add.f32 	%f1070, %f334, %f1067;
	add.f32 	%f1071, %f335, %f1068;
	add.f32 	%f1072, %f336, %f1069;
	// inline asm
	{  cvt.rn.f16.f32 %rs41, %f1072;}

	// inline asm
	// inline asm
	{  cvt.rn.f16.f32 %rs40, %f1071;}

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

	// inline asm
	st.v4.u16 	[%rd130], {%rs39, %rs40, %rs41, %rs35};
	bra.uni 	BB0_133;

BB0_132:
	mov.u64 	%rd143, image_RNM0;
	cvta.global.u64 	%rd138, %rd143;
	mov.u32 	%r321, 8;
	// inline asm
	call (%rd137), _rt_buffer_get_64, (%rd138, %r103, %r321, %rd18, %rd19, %rd27, %rd27);
	// inline asm
	// inline asm
	{  cvt.rn.f16.f32 %rs48, %f336;}

	// inline asm
	// inline asm
	{  cvt.rn.f16.f32 %rs47, %f335;}

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

	// inline asm
	st.v4.u16 	[%rd137], {%rs46, %rs47, %rs48, %rs35};

BB0_133:
	mov.f32 	%f1077, 0f34000000;
	max.f32 	%f1078, %f334, %f1077;
	div.rn.f32 	%f1079, %f337, %f1078;
	max.f32 	%f1080, %f335, %f1077;
	div.rn.f32 	%f1081, %f338, %f1080;
	max.f32 	%f1082, %f336, %f1077;
	div.rn.f32 	%f1083, %f339, %f1082;
	fma.rn.f32 	%f346, %f1079, 0f3F000000, 0f3F000000;
	fma.rn.f32 	%f347, %f1081, 0f3F000000, 0f3F000000;
	fma.rn.f32 	%f348, %f1083, 0f3F000000, 0f3F000000;
	div.rn.f32 	%f1084, %f340, %f1078;
	div.rn.f32 	%f1085, %f341, %f1080;
	div.rn.f32 	%f1086, %f342, %f1082;
	fma.rn.f32 	%f349, %f1084, 0f3F000000, 0f3F000000;
	fma.rn.f32 	%f350, %f1085, 0f3F000000, 0f3F000000;
	fma.rn.f32 	%f351, %f1086, 0f3F000000, 0f3F000000;
	div.rn.f32 	%f1087, %f343, %f1078;
	div.rn.f32 	%f1088, %f344, %f1080;
	div.rn.f32 	%f1089, %f345, %f1082;
	fma.rn.f32 	%f352, %f1087, 0f3F000000, 0f3F000000;
	fma.rn.f32 	%f353, %f1088, 0f3F000000, 0f3F000000;
	fma.rn.f32 	%f354, %f1089, 0f3F000000, 0f3F000000;
	ld.global.u32 	%r322, [additive];
	setp.eq.s32	%p138, %r322, 0;
	// inline asm
	{  cvt.rn.f16.f32 %rs49, %f564;}

	// inline asm
	@%p138 bra 	BB0_135;

	mov.u64 	%rd156, image_RNM1;
	cvta.global.u64 	%rd145, %rd156;
	mov.u32 	%r326, 8;
	// inline asm
	call (%rd144), _rt_buffer_get_64, (%rd145, %r103, %r326, %rd18, %rd19, %rd27, %rd27);
	// inline asm
	ld.v4.u16 	{%rs56, %rs57, %rs58, %rs59}, [%rd144];
	// inline asm
	{  cvt.f32.f16 %f1090, %rs56;}

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

	// inline asm
	// inline asm
	{  cvt.f32.f16 %f1092, %rs58;}

	// inline asm
	// inline asm
	call (%rd150), _rt_buffer_get_64, (%rd145, %r103, %r326, %rd18, %rd19, %rd27, %rd27);
	// inline asm
	add.f32 	%f1093, %f346, %f1090;
	add.f32 	%f1094, %f347, %f1091;
	add.f32 	%f1095, %f348, %f1092;
	// inline asm
	{  cvt.rn.f16.f32 %rs55, %f1095;}

	// inline asm
	// inline asm
	{  cvt.rn.f16.f32 %rs54, %f1094;}

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

	// inline asm
	st.v4.u16 	[%rd150], {%rs53, %rs54, %rs55, %rs49};
	bra.uni 	BB0_136;

BB0_135:
	mov.u64 	%rd163, image_RNM1;
	cvta.global.u64 	%rd158, %rd163;
	mov.u32 	%r328, 8;
	// inline asm
	call (%rd157), _rt_buffer_get_64, (%rd158, %r103, %r328, %rd18, %rd19, %rd27, %rd27);
	// inline asm
	// inline asm
	{  cvt.rn.f16.f32 %rs62, %f348;}

	// inline asm
	// inline asm
	{  cvt.rn.f16.f32 %rs61, %f347;}

	// inline asm
	// inline asm
	{  cvt.rn.f16.f32 %rs60, %f346;}

	// inline asm
	st.v4.u16 	[%rd157], {%rs60, %rs61, %rs62, %rs49};

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

	// inline asm
	@%p139 bra 	BB0_138;

	mov.u64 	%rd176, image_RNM2;
	cvta.global.u64 	%rd165, %rd176;
	mov.u32 	%r333, 8;
	// inline asm
	call (%rd164), _rt_buffer_get_64, (%rd165, %r103, %r333, %rd18, %rd19, %rd27, %rd27);
	// inline asm
	ld.v4.u16 	{%rs70, %rs71, %rs72, %rs73}, [%rd164];
	// inline asm
	{  cvt.f32.f16 %f1100, %rs70;}

	// inline asm
	// inline asm
	{  cvt.f32.f16 %f1101, %rs71;}

	// inline asm
	// inline asm
	{  cvt.f32.f16 %f1102, %rs72;}

	// inline asm
	// inline asm
	call (%rd170), _rt_buffer_get_64, (%rd165, %r103, %r333, %rd18, %rd19, %rd27, %rd27);
	// inline asm
	add.f32 	%f1103, %f349, %f1100;
	add.f32 	%f1104, %f350, %f1101;
	add.f32 	%f1105, %f351, %f1102;
	// inline asm
	{  cvt.rn.f16.f32 %rs69, %f1105;}

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

	// inline asm
	// inline asm
	{  cvt.rn.f16.f32 %rs67, %f1103;}

	// inline asm
	st.v4.u16 	[%rd170], {%rs67, %rs68, %rs69, %rs63};
	bra.uni 	BB0_139;

BB0_138:
	mov.u64 	%rd183, image_RNM2;
	cvta.global.u64 	%rd178, %rd183;
	mov.u32 	%r335, 8;
	// inline asm
	call (%rd177), _rt_buffer_get_64, (%rd178, %r103, %r335, %rd18, %rd19, %rd27, %rd27);
	// inline asm
	// inline asm
	{  cvt.rn.f16.f32 %rs76, %f351;}

	// inline asm
	// inline asm
	{  cvt.rn.f16.f32 %rs75, %f350;}

	// inline asm
	// inline asm
	{  cvt.rn.f16.f32 %rs74, %f349;}

	// inline asm
	st.v4.u16 	[%rd177], {%rs74, %rs75, %rs76, %rs63};

BB0_139:
	ld.global.u32 	%r336, [additive];
	setp.eq.s32	%p140, %r336, 0;
	// inline asm
	{  cvt.rn.f16.f32 %rs77, %f564;}

	// inline asm
	@%p140 bra 	BB0_141;

	mov.u64 	%rd196, image_RNM3;
	cvta.global.u64 	%rd185, %rd196;
	mov.u32 	%r340, 8;
	// inline asm
	call (%rd184), _rt_buffer_get_64, (%rd185, %r103, %r340, %rd18, %rd19, %rd27, %rd27);
	// inline asm
	ld.v4.u16 	{%rs84, %rs85, %rs86, %rs87}, [%rd184];
	// inline asm
	{  cvt.f32.f16 %f1110, %rs84;}

	// inline asm
	// inline asm
	{  cvt.f32.f16 %f1111, %rs85;}

	// inline asm
	// inline asm
	{  cvt.f32.f16 %f1112, %rs86;}

	// inline asm
	// inline asm
	call (%rd190), _rt_buffer_get_64, (%rd185, %r103, %r340, %rd18, %rd19, %rd27, %rd27);
	// inline asm
	add.f32 	%f1113, %f352, %f1110;
	add.f32 	%f1114, %f353, %f1111;
	add.f32 	%f1115, %f354, %f1112;
	// inline asm
	{  cvt.rn.f16.f32 %rs83, %f1115;}

	// inline asm
	// inline asm
	{  cvt.rn.f16.f32 %rs82, %f1114;}

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

	// inline asm
	st.v4.u16 	[%rd190], {%rs81, %rs82, %rs83, %rs77};
	bra.uni 	BB0_162;

BB0_141:
	mov.u64 	%rd203, image_RNM3;
	cvta.global.u64 	%rd198, %rd203;
	mov.u32 	%r342, 8;
	// inline asm
	call (%rd197), _rt_buffer_get_64, (%rd198, %r103, %r342, %rd18, %rd19, %rd27, %rd27);
	// inline asm
	// inline asm
	{  cvt.rn.f16.f32 %rs90, %f354;}

	// inline asm
	// inline asm
	{  cvt.rn.f16.f32 %rs89, %f353;}

	// inline asm
	// inline asm
	{  cvt.rn.f16.f32 %rs88, %f352;}

	// inline asm
	st.v4.u16 	[%rd197], {%rs88, %rs89, %rs90, %rs77};

BB0_162:
	ret;
}