//
// 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 4 .b8 payload[4];
.global .align 4 .b8 ray[36];
.global .align 4 .f32 t;
.global .align 1 .b8 rnd_seeds[1];
.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_typeinfo7payloadE[8] = {82, 97, 121, 0, 4, 0, 0, 0};
.global .align 4 .b8 _ZN21rti_internal_typeinfo3rayE[8] = {82, 97, 121, 0, 36, 0, 0, 0};
.global .align 4 .b8 _ZN21rti_internal_typeinfo1tE[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 8 .b8 _ZN21rti_internal_typename7payloadE[8] = {82, 97, 121, 68, 97, 116, 97, 0};
.global .align 4 .b8 _ZN21rti_internal_typename3rayE[4] = {82, 97, 121, 0};
.global .align 8 .b8 _ZN21rti_internal_typename1tE[6] = {102, 108, 111, 97, 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_typeenum7payloadE = 4919;
.global .align 4 .u32 _ZN21rti_internal_typeenum3rayE = 4919;
.global .align 4 .u32 _ZN21rti_internal_typeenum1tE = 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 16 .b8 _ZN21rti_internal_semantic7payloadE[10] = {114, 116, 80, 97, 121, 108, 111, 97, 100, 0};
.global .align 16 .b8 _ZN21rti_internal_semantic3rayE[13] = {114, 116, 67, 117, 114, 114, 101, 110, 116, 82, 97, 121, 0};
.global .align 16 .b8 _ZN21rti_internal_semantic1tE[23] = {114, 116, 73, 110, 116, 101, 114, 115, 101, 99, 116, 105, 111, 110, 68, 105, 115, 116, 97, 110, 99, 101, 0};
.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_annotation7payloadE[1];
.global .align 1 .b8 _ZN23rti_internal_annotation3rayE[1];
.global .align 1 .b8 _ZN23rti_internal_annotation1tE[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[32];
	.reg .b64 	%SP;
	.reg .b64 	%SPL;
	.reg .pred 	%p<31>;
	.reg .f32 	%f<192>;
	.reg .b32 	%r<212>;
	.reg .b64 	%rd<37>;


	mov.u64 	%rd36, __local_depot0;
	cvta.local.u64 	%SP, %rd36;
	ld.global.f32 	%f62, [normal];
	ld.global.f32 	%f63, [normal+4];
	ld.global.f32 	%f64, [normal+8];
	mov.u32 	%r80, 7937;
	mov.f32 	%f191, 0f00000000;
	// inline asm
	call (%f58, %f59, %f60, %f61), _rt_transform_tuple, (%r80, %f62, %f63, %f64, %f191);
	// inline asm
	mul.f32 	%f67, %f59, %f59;
	fma.rn.f32 	%f68, %f58, %f58, %f67;
	fma.rn.f32 	%f69, %f60, %f60, %f68;
	sqrt.rn.f32 	%f70, %f69;
	rcp.rn.f32 	%f71, %f70;
	mul.f32 	%f1, %f58, %f71;
	mul.f32 	%f2, %f59, %f71;
	mul.f32 	%f3, %f60, %f71;
	ld.global.f32 	%f72, [ray+12];
	ld.global.f32 	%f73, [t];
	ld.global.f32 	%f74, [ray+16];
	ld.global.f32 	%f75, [ray+20];
	ld.global.f32 	%f76, [ray];
	fma.rn.f32 	%f4, %f73, %f72, %f76;
	ld.global.f32 	%f77, [ray+4];
	fma.rn.f32 	%f5, %f73, %f74, %f77;
	ld.global.f32 	%f78, [ray+8];
	fma.rn.f32 	%f6, %f73, %f75, %f78;
	abs.f32 	%f79, %f1;
	abs.f32 	%f80, %f3;
	setp.gt.f32	%p1, %f79, %f80;
	neg.f32 	%f81, %f2;
	neg.f32 	%f82, %f3;
	selp.f32	%f83, %f81, 0f00000000, %p1;
	selp.f32	%f84, %f1, %f82, %p1;
	selp.f32	%f85, 0f00000000, %f2, %p1;
	mul.f32 	%f86, %f84, %f84;
	fma.rn.f32 	%f87, %f83, %f83, %f86;
	fma.rn.f32 	%f88, %f85, %f85, %f87;
	sqrt.rn.f32 	%f89, %f88;
	rcp.rn.f32 	%f90, %f89;
	mul.f32 	%f7, %f83, %f90;
	mul.f32 	%f8, %f84, %f90;
	mul.f32 	%f9, %f85, %f90;
	mul.f32 	%f91, %f3, %f8;
	mul.f32 	%f92, %f2, %f9;
	sub.f32 	%f10, %f91, %f92;
	mul.f32 	%f93, %f1, %f9;
	mul.f32 	%f94, %f3, %f7;
	sub.f32 	%f11, %f93, %f94;
	mul.f32 	%f95, %f2, %f7;
	mul.f32 	%f96, %f1, %f8;
	sub.f32 	%f12, %f95, %f96;
	ld.global.v2.u32 	{%r84, %r85}, [pixelID];
	cvt.u64.u32	%rd15, %r84;
	cvt.u64.u32	%rd16, %r85;
	mov.u64 	%rd19, rnd_seeds;
	cvta.global.u64 	%rd14, %rd19;
	mov.u32 	%r81, 2;
	mov.u32 	%r82, 4;
	mov.u64 	%rd18, 0;
	// inline asm
	call (%rd13), _rt_buffer_get_64, (%rd14, %r81, %r82, %rd15, %rd16, %rd18, %rd18);
	// inline asm
	ld.u32 	%r191, [%rd13];
	mov.u32 	%r83, 0;
	st.global.u32 	[payload], %r83;
	mov.u32 	%r188, %r83;

BB0_1:
	cvt.rn.f32.s32	%f14, %r188;
	mov.u32 	%r190, %r83;

BB0_2:
	mad.lo.s32 	%r89, %r191, 1664525, 1013904223;
	and.b32  	%r90, %r89, 16777215;
	cvt.rn.f32.u32	%f97, %r90;
	fma.rn.f32 	%f98, %f97, 0f33800000, %f14;
	mul.f32 	%f99, %f98, 0f3DCCCCCD;
	mad.lo.s32 	%r191, %r89, 1664525, 1013904223;
	and.b32  	%r91, %r191, 16777215;
	cvt.rn.f32.u32	%f100, %r91;
	cvt.rn.f32.s32	%f101, %r190;
	fma.rn.f32 	%f102, %f100, 0f33800000, %f101;
	mul.f32 	%f103, %f102, 0f3DCCCCCD;
	sqrt.rn.f32 	%f16, %f99;
	mul.f32 	%f185, %f103, 0f40C90FDB;
	abs.f32 	%f18, %f185;
	setp.neu.f32	%p2, %f18, 0f7F800000;
	mov.f32 	%f179, %f185;
	@%p2 bra 	BB0_4;

	mov.f32 	%f104, 0f00000000;
	mul.rn.f32 	%f179, %f185, %f104;

BB0_4:
	mul.f32 	%f105, %f179, 0f3F22F983;
	cvt.rni.s32.f32	%r201, %f105;
	cvt.rn.f32.s32	%f106, %r201;
	neg.f32 	%f107, %f106;
	mov.f32 	%f108, 0f3FC90FDA;
	fma.rn.f32 	%f109, %f107, %f108, %f179;
	mov.f32 	%f110, 0f33A22168;
	fma.rn.f32 	%f111, %f107, %f110, %f109;
	mov.f32 	%f112, 0f27C234C5;
	fma.rn.f32 	%f180, %f107, %f112, %f111;
	abs.f32 	%f113, %f179;
	setp.leu.f32	%p3, %f113, 0f47CE4780;
	@%p3 bra 	BB0_15;

	mov.b32 	 %r8, %f179;
	shr.u32 	%r9, %r8, 23;
	shl.b32 	%r94, %r8, 8;
	or.b32  	%r10, %r94, -2147483648;
	add.u64 	%rd21, %SP, 4;
	cvta.to.local.u64 	%rd33, %rd21;
	mov.u32 	%r193, 0;
	mov.u64 	%rd32, __cudart_i2opi_f;
	mov.u32 	%r192, -6;

BB0_6:
	.pragma "nounroll";
	ld.const.u32 	%r97, [%rd32];
	// inline asm
	{
	mad.lo.cc.u32   %r95, %r97, %r10, %r193;
	madc.hi.u32     %r193, %r97, %r10,  0;
	}
	// inline asm
	st.local.u32 	[%rd33], %r95;
	add.s64 	%rd33, %rd33, 4;
	add.s64 	%rd32, %rd32, 4;
	add.s32 	%r192, %r192, 1;
	setp.ne.s32	%p4, %r192, 0;
	@%p4 bra 	BB0_6;

	and.b32  	%r100, %r9, 255;
	add.s32 	%r101, %r100, -128;
	shr.u32 	%r102, %r101, 5;
	and.b32  	%r15, %r8, -2147483648;
	cvta.to.local.u64 	%rd23, %rd21;
	st.local.u32 	[%rd23+24], %r193;
	mov.u32 	%r103, 6;
	sub.s32 	%r104, %r103, %r102;
	mul.wide.s32 	%rd24, %r104, 4;
	add.s64 	%rd6, %rd23, %rd24;
	ld.local.u32 	%r194, [%rd6];
	ld.local.u32 	%r195, [%rd6+-4];
	and.b32  	%r18, %r9, 31;
	setp.eq.s32	%p5, %r18, 0;
	@%p5 bra 	BB0_9;

	mov.u32 	%r105, 32;
	sub.s32 	%r106, %r105, %r18;
	shr.u32 	%r107, %r195, %r106;
	shl.b32 	%r108, %r194, %r18;
	add.s32 	%r194, %r107, %r108;
	ld.local.u32 	%r109, [%rd6+-8];
	shr.u32 	%r110, %r109, %r106;
	shl.b32 	%r111, %r195, %r18;
	add.s32 	%r195, %r110, %r111;

BB0_9:
	shr.u32 	%r112, %r195, 30;
	shl.b32 	%r113, %r194, 2;
	add.s32 	%r196, %r112, %r113;
	shl.b32 	%r24, %r195, 2;
	shr.u32 	%r114, %r196, 31;
	shr.u32 	%r115, %r194, 30;
	add.s32 	%r25, %r114, %r115;
	setp.eq.s32	%p6, %r114, 0;
	@%p6 bra 	BB0_10;
	bra.uni 	BB0_11;

BB0_10:
	mov.u32 	%r197, %r15;
	mov.u32 	%r198, %r24;
	bra.uni 	BB0_12;

BB0_11:
	not.b32 	%r116, %r196;
	neg.s32 	%r198, %r24;
	setp.eq.s32	%p7, %r24, 0;
	selp.u32	%r117, 1, 0, %p7;
	add.s32 	%r196, %r117, %r116;
	xor.b32  	%r197, %r15, -2147483648;

BB0_12:
	clz.b32 	%r200, %r196;
	setp.eq.s32	%p8, %r200, 0;
	shl.b32 	%r118, %r196, %r200;
	mov.u32 	%r119, 32;
	sub.s32 	%r120, %r119, %r200;
	shr.u32 	%r121, %r198, %r120;
	add.s32 	%r122, %r121, %r118;
	selp.b32	%r33, %r196, %r122, %p8;
	mov.u32 	%r123, -921707870;
	mul.hi.u32 	%r199, %r33, %r123;
	setp.eq.s32	%p9, %r15, 0;
	neg.s32 	%r124, %r25;
	selp.b32	%r201, %r25, %r124, %p9;
	setp.lt.s32	%p10, %r199, 1;
	@%p10 bra 	BB0_14;

	mul.lo.s32 	%r125, %r33, -921707870;
	shr.u32 	%r126, %r125, 31;
	shl.b32 	%r127, %r199, 1;
	add.s32 	%r199, %r126, %r127;
	add.s32 	%r200, %r200, 1;

BB0_14:
	mov.u32 	%r128, 126;
	sub.s32 	%r129, %r128, %r200;
	shl.b32 	%r130, %r129, 23;
	add.s32 	%r131, %r199, 1;
	shr.u32 	%r132, %r131, 7;
	add.s32 	%r133, %r132, 1;
	shr.u32 	%r134, %r133, 1;
	add.s32 	%r135, %r134, %r130;
	or.b32  	%r136, %r135, %r197;
	mov.b32 	 %f180, %r136;

BB0_15:
	mul.rn.f32 	%f24, %f180, %f180;
	add.s32 	%r41, %r201, 1;
	and.b32  	%r42, %r41, 1;
	setp.eq.s32	%p11, %r42, 0;
	@%p11 bra 	BB0_17;
	bra.uni 	BB0_16;

BB0_17:
	mov.f32 	%f116, 0f3C08839E;
	mov.f32 	%f117, 0fB94CA1F9;
	fma.rn.f32 	%f181, %f117, %f24, %f116;
	bra.uni 	BB0_18;

BB0_16:
	mov.f32 	%f114, 0fBAB6061A;
	mov.f32 	%f115, 0f37CCF5CE;
	fma.rn.f32 	%f181, %f115, %f24, %f114;

BB0_18:
	@%p11 bra 	BB0_20;
	bra.uni 	BB0_19;

BB0_20:
	mov.f32 	%f121, 0fBE2AAAA3;
	fma.rn.f32 	%f122, %f181, %f24, %f121;
	mov.f32 	%f123, 0f00000000;
	fma.rn.f32 	%f182, %f122, %f24, %f123;
	bra.uni 	BB0_21;

BB0_19:
	mov.f32 	%f118, 0f3D2AAAA5;
	fma.rn.f32 	%f119, %f181, %f24, %f118;
	mov.f32 	%f120, 0fBF000000;
	fma.rn.f32 	%f182, %f119, %f24, %f120;

BB0_21:
	fma.rn.f32 	%f183, %f182, %f180, %f180;
	@%p11 bra 	BB0_23;

	mov.f32 	%f124, 0f3F800000;
	fma.rn.f32 	%f183, %f182, %f24, %f124;

BB0_23:
	and.b32  	%r137, %r41, 2;
	setp.eq.s32	%p14, %r137, 0;
	@%p14 bra 	BB0_25;

	mov.f32 	%f125, 0f00000000;
	mov.f32 	%f126, 0fBF800000;
	fma.rn.f32 	%f183, %f183, %f126, %f125;

BB0_25:
	@%p2 bra 	BB0_27;

	mov.f32 	%f127, 0f00000000;
	mul.rn.f32 	%f185, %f185, %f127;

BB0_27:
	mul.f32 	%f128, %f185, 0f3F22F983;
	cvt.rni.s32.f32	%r211, %f128;
	cvt.rn.f32.s32	%f129, %r211;
	neg.f32 	%f130, %f129;
	fma.rn.f32 	%f132, %f130, %f108, %f185;
	fma.rn.f32 	%f134, %f130, %f110, %f132;
	fma.rn.f32 	%f186, %f130, %f112, %f134;
	abs.f32 	%f136, %f185;
	setp.leu.f32	%p16, %f136, 0f47CE4780;
	@%p16 bra 	BB0_38;

	mov.b32 	 %r44, %f185;
	shr.u32 	%r45, %r44, 23;
	shl.b32 	%r140, %r44, 8;
	or.b32  	%r46, %r140, -2147483648;
	add.u64 	%rd26, %SP, 4;
	cvta.to.local.u64 	%rd35, %rd26;
	mov.u32 	%r203, 0;
	mov.u64 	%rd34, __cudart_i2opi_f;
	mov.u32 	%r202, -6;

BB0_29:
	.pragma "nounroll";
	ld.const.u32 	%r143, [%rd34];
	// inline asm
	{
	mad.lo.cc.u32   %r141, %r143, %r46, %r203;
	madc.hi.u32     %r203, %r143, %r46,  0;
	}
	// inline asm
	st.local.u32 	[%rd35], %r141;
	add.s64 	%rd35, %rd35, 4;
	add.s64 	%rd34, %rd34, 4;
	add.s32 	%r202, %r202, 1;
	setp.ne.s32	%p17, %r202, 0;
	@%p17 bra 	BB0_29;

	and.b32  	%r146, %r45, 255;
	add.s32 	%r147, %r146, -128;
	shr.u32 	%r148, %r147, 5;
	and.b32  	%r51, %r44, -2147483648;
	cvta.to.local.u64 	%rd28, %rd26;
	st.local.u32 	[%rd28+24], %r203;
	mov.u32 	%r149, 6;
	sub.s32 	%r150, %r149, %r148;
	mul.wide.s32 	%rd29, %r150, 4;
	add.s64 	%rd12, %rd28, %rd29;
	ld.local.u32 	%r204, [%rd12];
	ld.local.u32 	%r205, [%rd12+-4];
	and.b32  	%r54, %r45, 31;
	setp.eq.s32	%p18, %r54, 0;
	@%p18 bra 	BB0_32;

	mov.u32 	%r151, 32;
	sub.s32 	%r152, %r151, %r54;
	shr.u32 	%r153, %r205, %r152;
	shl.b32 	%r154, %r204, %r54;
	add.s32 	%r204, %r153, %r154;
	ld.local.u32 	%r155, [%rd12+-8];
	shr.u32 	%r156, %r155, %r152;
	shl.b32 	%r157, %r205, %r54;
	add.s32 	%r205, %r156, %r157;

BB0_32:
	shr.u32 	%r158, %r205, 30;
	shl.b32 	%r159, %r204, 2;
	add.s32 	%r206, %r158, %r159;
	shl.b32 	%r60, %r205, 2;
	shr.u32 	%r160, %r206, 31;
	shr.u32 	%r161, %r204, 30;
	add.s32 	%r61, %r160, %r161;
	setp.eq.s32	%p19, %r160, 0;
	@%p19 bra 	BB0_33;
	bra.uni 	BB0_34;

BB0_33:
	mov.u32 	%r207, %r51;
	mov.u32 	%r208, %r60;
	bra.uni 	BB0_35;

BB0_34:
	not.b32 	%r162, %r206;
	neg.s32 	%r208, %r60;
	setp.eq.s32	%p20, %r60, 0;
	selp.u32	%r163, 1, 0, %p20;
	add.s32 	%r206, %r163, %r162;
	xor.b32  	%r207, %r51, -2147483648;

BB0_35:
	clz.b32 	%r210, %r206;
	setp.eq.s32	%p21, %r210, 0;
	shl.b32 	%r164, %r206, %r210;
	mov.u32 	%r165, 32;
	sub.s32 	%r166, %r165, %r210;
	shr.u32 	%r167, %r208, %r166;
	add.s32 	%r168, %r167, %r164;
	selp.b32	%r69, %r206, %r168, %p21;
	mov.u32 	%r169, -921707870;
	mul.hi.u32 	%r209, %r69, %r169;
	setp.eq.s32	%p22, %r51, 0;
	neg.s32 	%r170, %r61;
	selp.b32	%r211, %r61, %r170, %p22;
	setp.lt.s32	%p23, %r209, 1;
	@%p23 bra 	BB0_37;

	mul.lo.s32 	%r171, %r69, -921707870;
	shr.u32 	%r172, %r171, 31;
	shl.b32 	%r173, %r209, 1;
	add.s32 	%r209, %r172, %r173;
	add.s32 	%r210, %r210, 1;

BB0_37:
	mov.u32 	%r174, 126;
	sub.s32 	%r175, %r174, %r210;
	shl.b32 	%r176, %r175, 23;
	add.s32 	%r177, %r209, 1;
	shr.u32 	%r178, %r177, 7;
	add.s32 	%r179, %r178, 1;
	shr.u32 	%r180, %r179, 1;
	add.s32 	%r181, %r180, %r176;
	or.b32  	%r182, %r181, %r207;
	mov.b32 	 %f186, %r182;

BB0_38:
	mul.rn.f32 	%f41, %f186, %f186;
	and.b32  	%r77, %r211, 1;
	setp.eq.s32	%p24, %r77, 0;
	@%p24 bra 	BB0_40;
	bra.uni 	BB0_39;

BB0_40:
	mov.f32 	%f139, 0f3C08839E;
	mov.f32 	%f140, 0fB94CA1F9;
	fma.rn.f32 	%f187, %f140, %f41, %f139;
	bra.uni 	BB0_41;

BB0_39:
	mov.f32 	%f137, 0fBAB6061A;
	mov.f32 	%f138, 0f37CCF5CE;
	fma.rn.f32 	%f187, %f138, %f41, %f137;

BB0_41:
	@%p24 bra 	BB0_43;
	bra.uni 	BB0_42;

BB0_43:
	mov.f32 	%f144, 0fBE2AAAA3;
	fma.rn.f32 	%f145, %f187, %f41, %f144;
	mov.f32 	%f146, 0f00000000;
	fma.rn.f32 	%f188, %f145, %f41, %f146;
	bra.uni 	BB0_44;

BB0_42:
	mov.f32 	%f141, 0f3D2AAAA5;
	fma.rn.f32 	%f142, %f187, %f41, %f141;
	mov.f32 	%f143, 0fBF000000;
	fma.rn.f32 	%f188, %f142, %f41, %f143;

BB0_44:
	fma.rn.f32 	%f189, %f188, %f186, %f186;
	@%p24 bra 	BB0_46;

	mov.f32 	%f147, 0f3F800000;
	fma.rn.f32 	%f189, %f188, %f41, %f147;

BB0_46:
	and.b32  	%r183, %r211, 2;
	setp.eq.s32	%p27, %r183, 0;
	@%p27 bra 	BB0_48;

	mov.f32 	%f148, 0f00000000;
	mov.f32 	%f149, 0fBF800000;
	fma.rn.f32 	%f189, %f189, %f149, %f148;

BB0_48:
	mul.f32 	%f150, %f16, %f183;
	mul.f32 	%f151, %f150, %f150;
	mov.f32 	%f152, 0f3F800000;
	sub.f32 	%f153, %f152, %f151;
	mul.f32 	%f154, %f16, %f189;
	mul.f32 	%f155, %f154, %f154;
	sub.f32 	%f156, %f153, %f155;
	mov.f32 	%f157, 0f00000000;
	max.f32 	%f158, %f157, %f156;
	sqrt.rn.f32 	%f159, %f158;
	mul.f32 	%f160, %f7, %f154;
	mul.f32 	%f161, %f8, %f154;
	mul.f32 	%f162, %f9, %f154;
	fma.rn.f32 	%f163, %f10, %f150, %f160;
	fma.rn.f32 	%f164, %f11, %f150, %f161;
	fma.rn.f32 	%f165, %f12, %f150, %f162;
	fma.rn.f32 	%f53, %f1, %f159, %f163;
	fma.rn.f32 	%f54, %f2, %f159, %f164;
	fma.rn.f32 	%f55, %f3, %f159, %f165;
	setp.leu.f32	%p28, %f54, 0f00000000;
	@%p28 bra 	BB0_50;

	add.u64 	%rd30, %SP, 0;
	cvta.to.local.u64 	%rd31, %rd30;
	mov.u32 	%r187, 1065353216;
	st.local.u32 	[%rd31], %r187;
	ld.global.u32 	%r184, [root];
	mov.u32 	%r185, 1;
	mov.f32 	%f172, 0f3A83126F;
	mov.f32 	%f173, 0f6C4ECB8F;
	// inline asm
	call _rt_trace_64, (%r184, %f4, %f5, %f6, %f53, %f54, %f55, %r185, %f172, %f173, %rd30, %r82);
	// inline asm
	ld.global.f32 	%f174, [payload];
	ld.local.f32 	%f175, [%rd31];
	add.f32 	%f191, %f175, %f174;
	st.global.f32 	[payload], %f191;

BB0_50:
	add.s32 	%r190, %r190, 1;
	setp.lt.s32	%p29, %r190, 10;
	@%p29 bra 	BB0_2;

	add.s32 	%r188, %r188, 1;
	setp.lt.s32	%p30, %r188, 10;
	@%p30 bra 	BB0_1;

	div.rn.f32 	%f176, %f191, 0f42C80000;
	st.global.f32 	[payload], %f176;
	ret;
}