//
// Generated by NVIDIA NVVM Compiler
//
// Compiler Build ID: CL-31968024
// Cuda compilation tools, release 12.0, V12.0.76
// Based on NVVM 7.0.1
//

.version 8.0
.target sm_52
.address_size 64

	// .globl	__raygen__oxMain
.const .align 16 .b8 params[1184];

.visible .entry __raygen__oxMain()
{
	.reg .pred 	%p<34>;
	.reg .b16 	%rs<20>;
	.reg .f32 	%f<418>;
	.reg .b32 	%r<672>;
	.reg .b64 	%rd<42>;


	// begin inline asm
	call (%r7), _optix_get_launch_index_x, ();
	// end inline asm
	// begin inline asm
	call (%r8), _optix_get_launch_index_y, ();
	// end inline asm
	ld.const.u64 	%rd6, [params+416];
	cvta.to.global.u64 	%rd7, %rd6;
	ld.const.u32 	%r13, [params+408];
	mad.lo.s32 	%r14, %r13, %r8, %r7;
	mul.wide.u32 	%rd8, %r14, 4;
	add.s64 	%rd1, %rd7, %rd8;
	ld.global.v2.u8 	{%rs7, %rs19}, [%rd1];
	or.b16  	%rs9, %rs7, %rs19;
	and.b16  	%rs10, %rs9, 255;
	setp.eq.s16 	%p1, %rs10, 0;
	@%p1 bra 	$L__BB0_2;

	ld.global.u8 	%rs18, [%rd1+2];
	bra.uni 	$L__BB0_3;

$L__BB0_2:
	ld.global.u8 	%rs18, [%rd1+2];
	setp.eq.s16 	%p2, %rs18, 0;
	mov.f32 	%f396, 0f00000000;
	mov.u16 	%rs19, 0;
	mov.f32 	%f397, %f396;
	mov.f32 	%f398, %f396;
	@%p2 bra 	$L__BB0_4;

$L__BB0_3:
	cvt.rn.f32.u16 	%f124, %rs7;
	div.rn.f32 	%f125, %f124, 0f437F0000;
	fma.rn.f32 	%f126, %f125, 0f40000000, 0fBF800000;
	and.b16  	%rs13, %rs19, 255;
	cvt.rn.f32.u16 	%f127, %rs13;
	div.rn.f32 	%f128, %f127, 0f437F0000;
	fma.rn.f32 	%f129, %f128, 0f40000000, 0fBF800000;
	cvt.rn.f32.u16 	%f130, %rs18;
	div.rn.f32 	%f131, %f130, 0f437F0000;
	fma.rn.f32 	%f132, %f131, 0f40000000, 0fBF800000;
	mul.f32 	%f133, %f129, %f129;
	fma.rn.f32 	%f134, %f126, %f126, %f133;
	fma.rn.f32 	%f135, %f132, %f132, %f134;
	sqrt.rn.f32 	%f136, %f135;
	rcp.rn.f32 	%f137, %f136;
	mul.f32 	%f398, %f137, %f132;
	mul.f32 	%f397, %f137, %f129;
	mul.f32 	%f396, %f126, %f137;

$L__BB0_4:
	ld.const.v2.u32 	{%r15, %r16}, [params];
	add.s32 	%r3, %r15, %r7;
	add.s32 	%r4, %r16, %r8;
	setp.eq.f32 	%p3, %f396, 0f00000000;
	setp.eq.f32 	%p4, %f397, 0f00000000;
	and.pred  	%p5, %p3, %p4;
	setp.eq.f32 	%p6, %f398, 0f00000000;
	and.pred  	%p7, %p6, %p5;
	@%p7 bra 	$L__BB0_26;
	bra.uni 	$L__BB0_5;

$L__BB0_26:
	ld.const.u64 	%rd37, [params+224];
	cvta.to.global.u64 	%rd38, %rd37;
	ld.const.u32 	%r669, [params+216];
	mad.lo.s32 	%r670, %r669, %r4, %r3;
	mul.wide.u32 	%rd39, %r670, 16;
	add.s64 	%rd40, %rd38, %rd39;
	mov.f32 	%f391, 0f00000000;
	st.global.v4.f32 	[%rd40], {%f391, %f391, %f391, %f391};
	bra.uni 	$L__BB0_27;

$L__BB0_5:
	ld.const.u64 	%rd9, [params+432];
	cvta.to.global.u64 	%rd10, %rd9;
	ld.const.u32 	%r19, [params+424];
	mad.lo.s32 	%r20, %r19, %r8, %r7;
	mul.wide.u32 	%rd11, %r20, 16;
	add.s64 	%rd12, %rd10, %rd11;
	ld.global.v4.f32 	{%f138, %f139, %f140, %f141}, [%rd12];
	abs.f32 	%f146, %f138;
	setp.le.f32 	%p8, %f146, 0f7F800000;
	selp.f32 	%f10, %f138, 0f00000000, %p8;
	abs.f32 	%f147, %f139;
	setp.le.f32 	%p9, %f147, 0f7F800000;
	selp.f32 	%f11, %f139, 0f00000000, %p9;
	abs.f32 	%f148, %f140;
	setp.le.f32 	%p10, %f148, 0f7F800000;
	selp.f32 	%f12, %f140, 0f00000000, %p10;
	add.f32 	%f13, %f141, 0f38D1B717;
	mov.f32 	%f149, 0f38D1B717;
	mul.f32 	%f150, %f10, 0f3456BF95;
	mul.f32 	%f151, %f11, 0f3456BF95;
	mul.f32 	%f152, %f12, 0f3456BF95;
	abs.f32 	%f14, %f396;
	div.rn.f32 	%f153, %f150, %f14;
	abs.f32 	%f154, %f397;
	div.rn.f32 	%f155, %f151, %f154;
	abs.f32 	%f15, %f398;
	div.rn.f32 	%f156, %f152, %f15;
	abs.f32 	%f157, %f153;
	abs.f32 	%f158, %f155;
	abs.f32 	%f159, %f156;
	max.f32 	%f160, %f157, %f149;
	max.f32 	%f161, %f158, %f149;
	max.f32 	%f162, %f159, %f149;
	mul.f32 	%f16, %f396, %f160;
	mul.f32 	%f17, %f397, %f161;
	mul.f32 	%f18, %f398, %f162;
	add.f32 	%f19, %f10, %f16;
	add.f32 	%f20, %f11, %f17;
	add.f32 	%f21, %f12, %f18;
	ld.const.u64 	%rd13, [params+448];
	cvta.to.global.u64 	%rd14, %rd13;
	ld.const.u32 	%r21, [params+440];
	mad.lo.s32 	%r22, %r21, %r8, %r7;
	mul.wide.u32 	%rd15, %r22, 16;
	add.s64 	%rd16, %rd14, %rd15;
	ld.global.v4.f32 	{%f163, %f164, %f165, %f166}, [%rd16];
	cvt.rzi.u32.f32 	%r23, %f166;
	ld.const.u64 	%rd17, [params+712];
	cvta.to.global.u64 	%rd18, %rd17;
	cvt.u64.u32 	%rd19, %r23;
	add.s64 	%rd2, %rd18, %rd19;
	ld.global.u8 	%r671, [%rd2];
	sub.f32 	%f25, %f163, %f10;
	sub.f32 	%f26, %f164, %f11;
	sub.f32 	%f27, %f165, %f12;
	mul.f32 	%f168, %f26, %f26;
	fma.rn.f32 	%f169, %f25, %f25, %f168;
	fma.rn.f32 	%f170, %f27, %f27, %f169;
	sqrt.rn.f32 	%f28, %f170;
	setp.ne.s32 	%p11, %r671, 255;
	setp.gt.f32 	%p12, %f28, 0f3727C5AC;
	and.pred  	%p13, %p11, %p12;
	@%p13 bra 	$L__BB0_7;
	bra.uni 	$L__BB0_6;

$L__BB0_7:
	rcp.rn.f32 	%f180, %f28;
	mul.f32 	%f174, %f25, %f180;
	mul.f32 	%f175, %f26, %f180;
	mul.f32 	%f176, %f27, %f180;
	ld.const.u64 	%rd41, [params+96];
	mov.f32 	%f179, 0f00000000;
	mov.u32 	%r60, 2;
	mov.u32 	%r62, 1;
	mov.u32 	%r63, 1065353216;
	mov.u32 	%r94, 0;
	// begin inline asm
	call(%r24,%r25,%r26,%r27,%r28,%r29,%r30,%r31,%r32,%r33,%r34,%r35,%r36,%r37,%r38,%r39,%r40,%r41,%r42,%r43,%r44,%r45,%r46,%r47,%r48,%r49,%r50,%r51,%r52,%r53,%r54,%r55),_optix_trace_typed_32,(%r94,%rd41,%f19,%f20,%f21,%f174,%f175,%f176,%f149,%f28,%f179,%r62,%r94,%r62,%r60,%r62,%r62,%r63,%r94,%r94,%r94,%r94,%r94,%r94,%r94,%r94,%r94,%r94,%r94,%r94,%r94,%r94,%r94,%r94,%r94,%r94,%r94,%r94,%r94,%r94,%r94,%r94,%r94,%r94,%r94,%r94,%r94,%r94,%r94);
	// end inline asm
	mov.b32 	%f181, %r24;
	setp.neu.f32 	%p14, %f181, 0f00000000;
	@%p14 bra 	$L__BB0_9;

	mov.u16 	%rs15, 255;
	st.global.u8 	[%rd2], %rs15;
	mov.u32 	%r671, 255;
	bra.uni 	$L__BB0_9;

$L__BB0_6:
	ld.const.u64 	%rd41, [params+96];

$L__BB0_9:
	setp.eq.s32 	%p15, %r671, 0;
	mov.u32 	%r166, 0;
	selp.f32 	%f404, %f163, %f10, %p15;
	add.f32 	%f182, %f16, %f404;
	selp.f32 	%f405, %f164, %f11, %p15;
	add.f32 	%f183, %f17, %f405;
	selp.f32 	%f406, %f165, %f12, %p15;
	add.f32 	%f184, %f18, %f406;
	neg.f32 	%f192, %f397;
	setp.gt.f32 	%p16, %f14, %f15;
	selp.f32 	%f193, %f192, 0f00000000, %p16;
	mov.f32 	%f190, 0f00000000;
	neg.f32 	%f194, %f398;
	selp.f32 	%f195, %f396, %f194, %p16;
	selp.f32 	%f196, 0f00000000, %f397, %p16;
	mul.f32 	%f197, %f195, %f195;
	fma.rn.f32 	%f198, %f193, %f193, %f197;
	fma.rn.f32 	%f199, %f196, %f196, %f198;
	sqrt.rn.f32 	%f200, %f199;
	rcp.rn.f32 	%f201, %f200;
	mul.f32 	%f32, %f193, %f201;
	mul.f32 	%f33, %f195, %f201;
	mul.f32 	%f34, %f196, %f201;
	mul.f32 	%f202, %f398, %f33;
	mul.f32 	%f203, %f397, %f34;
	sub.f32 	%f35, %f202, %f203;
	mul.f32 	%f204, %f396, %f34;
	mul.f32 	%f205, %f398, %f32;
	sub.f32 	%f36, %f204, %f205;
	mul.f32 	%f206, %f397, %f32;
	mul.f32 	%f207, %f396, %f33;
	sub.f32 	%f37, %f206, %f207;
	mov.f32 	%f208, 0f3F8147AE;
	sqrt.rn.f32 	%f209, %f208;
	rcp.rn.f32 	%f38, %f209;
	mul.f32 	%f40, %f38, 0f00000000;
	mul.f32 	%f210, %f38, 0f3DCCCCCD;
	mul.f32 	%f211, %f182, 0f3456BF95;
	mul.f32 	%f212, %f183, 0f3456BF95;
	mul.f32 	%f213, %f184, 0f3456BF95;
	abs.f32 	%f214, %f211;
	abs.f32 	%f215, %f212;
	abs.f32 	%f216, %f213;
	max.f32 	%f217, %f214, %f215;
	max.f32 	%f218, %f217, %f216;
	mov.f32 	%f219, 0f38D1B717;
	max.f32 	%f41, %f218, %f219;
	mul.f32 	%f220, %f35, %f38;
	mul.f32 	%f221, %f36, %f38;
	mul.f32 	%f222, %f37, %f38;
	mul.f32 	%f45, %f32, %f40;
	mul.f32 	%f46, %f33, %f40;
	mul.f32 	%f47, %f34, %f40;
	sub.f32 	%f223, %f45, %f220;
	sub.f32 	%f224, %f46, %f221;
	sub.f32 	%f225, %f47, %f222;
	mul.f32 	%f48, %f396, %f210;
	mul.f32 	%f49, %f397, %f210;
	mul.f32 	%f50, %f398, %f210;
	add.f32 	%f185, %f223, %f48;
	add.f32 	%f186, %f224, %f49;
	add.f32 	%f187, %f225, %f50;
	mul.f32 	%f189, %f13, 0f3FB504F3;
	mov.u32 	%r129, 1;
	mov.u32 	%r132, 2;
	mov.u32 	%r134, 3;
	// begin inline asm
	call(%r96,%r97,%r98,%r99,%r100,%r101,%r102,%r103,%r104,%r105,%r106,%r107,%r108,%r109,%r110,%r111,%r112,%r113,%r114,%r115,%r116,%r117,%r118,%r119,%r120,%r121,%r122,%r123,%r124,%r125,%r126,%r127),_optix_trace_typed_32,(%r166,%rd41,%f182,%f183,%f184,%f185,%f186,%f187,%f190,%f189,%f190,%r129,%r166,%r166,%r132,%r166,%r134,%r166,%r166,%r166,%r166,%r166,%r166,%r166,%r166,%r166,%r166,%r166,%r166,%r166,%r166,%r166,%r166,%r166,%r166,%r166,%r166,%r166,%r166,%r166,%r166,%r166,%r166,%r166,%r166,%r166,%r166,%r166,%r166);
	// end inline asm
	mov.b32 	%f55, %r96;
	mov.b32 	%f56, %r97;
	mov.b32 	%f57, %r98;
	add.f32 	%f226, %f55, %f56;
	add.f32 	%f227, %f226, %f57;
	setp.eq.f32 	%p17, %f227, 0f00000000;
	mov.f32 	%f403, 0f47C34F80;
	@%p17 bra 	$L__BB0_11;

	mul.f32 	%f228, %f56, %f56;
	fma.rn.f32 	%f229, %f55, %f55, %f228;
	fma.rn.f32 	%f230, %f57, %f57, %f229;
	sqrt.rn.f32 	%f231, %f230;
	rcp.rn.f32 	%f232, %f231;
	mul.f32 	%f233, %f232, %f55;
	mul.f32 	%f234, %f232, %f56;
	mul.f32 	%f235, %f232, %f57;
	fma.rn.f32 	%f236, %f41, %f233, %f182;
	fma.rn.f32 	%f237, %f41, %f234, %f183;
	fma.rn.f32 	%f238, %f41, %f235, %f184;
	setp.lt.f32 	%p18, %f231, 0f47C34F80;
	fma.rn.f32 	%f239, %f231, %f185, %f236;
	fma.rn.f32 	%f240, %f231, %f186, %f237;
	fma.rn.f32 	%f241, %f231, %f187, %f238;
	min.f32 	%f403, %f231, 0f47C34F80;
	selp.f32 	%f406, %f241, %f406, %p18;
	selp.f32 	%f405, %f240, %f405, %p18;
	selp.f32 	%f404, %f239, %f404, %p18;

$L__BB0_11:
	fma.rn.f32 	%f251, %f35, %f38, %f45;
	fma.rn.f32 	%f252, %f36, %f38, %f46;
	fma.rn.f32 	%f253, %f37, %f38, %f47;
	add.f32 	%f245, %f251, %f48;
	add.f32 	%f246, %f252, %f49;
	add.f32 	%f247, %f253, %f50;
	// begin inline asm
	call(%r167,%r168,%r169,%r170,%r171,%r172,%r173,%r174,%r175,%r176,%r177,%r178,%r179,%r180,%r181,%r182,%r183,%r184,%r185,%r186,%r187,%r188,%r189,%r190,%r191,%r192,%r193,%r194,%r195,%r196,%r197,%r198),_optix_trace_typed_32,(%r166,%rd41,%f182,%f183,%f184,%f245,%f246,%f247,%f190,%f189,%f190,%r129,%r166,%r166,%r132,%r166,%r134,%r166,%r166,%r166,%r166,%r166,%r166,%r166,%r166,%r166,%r166,%r166,%r166,%r166,%r166,%r166,%r166,%r166,%r166,%r166,%r166,%r166,%r166,%r166,%r166,%r166,%r166,%r166,%r166,%r166,%r166,%r166,%r166);
	// end inline asm
	mov.b32 	%f69, %r167;
	mov.b32 	%f70, %r168;
	mov.b32 	%f71, %r169;
	add.f32 	%f254, %f69, %f70;
	add.f32 	%f255, %f254, %f71;
	setp.eq.f32 	%p19, %f255, 0f00000000;
	@%p19 bra 	$L__BB0_13;

	mul.f32 	%f256, %f70, %f70;
	fma.rn.f32 	%f257, %f69, %f69, %f256;
	fma.rn.f32 	%f258, %f71, %f71, %f257;
	sqrt.rn.f32 	%f259, %f258;
	rcp.rn.f32 	%f260, %f259;
	mul.f32 	%f261, %f260, %f69;
	mul.f32 	%f262, %f260, %f70;
	mul.f32 	%f263, %f260, %f71;
	fma.rn.f32 	%f264, %f41, %f261, %f182;
	fma.rn.f32 	%f265, %f41, %f262, %f183;
	fma.rn.f32 	%f266, %f41, %f263, %f184;
	setp.lt.f32 	%p20, %f259, %f403;
	fma.rn.f32 	%f267, %f259, %f245, %f264;
	fma.rn.f32 	%f268, %f259, %f246, %f265;
	fma.rn.f32 	%f269, %f259, %f247, %f266;
	selp.f32 	%f403, %f259, %f403, %p20;
	selp.f32 	%f406, %f269, %f406, %p20;
	selp.f32 	%f405, %f268, %f405, %p20;
	selp.f32 	%f404, %f267, %f404, %p20;

$L__BB0_13:
	neg.f32 	%f394, %f38;
	mul.f32 	%f393, %f38, 0f00000000;
	mul.f32 	%f80, %f35, %f393;
	fma.rn.f32 	%f279, %f32, %f394, %f80;
	mul.f32 	%f81, %f36, %f393;
	fma.rn.f32 	%f280, %f33, %f394, %f81;
	mul.f32 	%f82, %f37, %f393;
	fma.rn.f32 	%f281, %f34, %f394, %f82;
	add.f32 	%f273, %f279, %f48;
	add.f32 	%f274, %f280, %f49;
	add.f32 	%f275, %f281, %f50;
	mov.f32 	%f278, 0f00000000;
	mov.u32 	%r271, 1;
	mov.u32 	%r274, 2;
	mov.u32 	%r276, 3;
	mov.u32 	%r308, 0;
	// begin inline asm
	call(%r238,%r239,%r240,%r241,%r242,%r243,%r244,%r245,%r246,%r247,%r248,%r249,%r250,%r251,%r252,%r253,%r254,%r255,%r256,%r257,%r258,%r259,%r260,%r261,%r262,%r263,%r264,%r265,%r266,%r267,%r268,%r269),_optix_trace_typed_32,(%r308,%rd41,%f182,%f183,%f184,%f273,%f274,%f275,%f278,%f189,%f278,%r271,%r308,%r308,%r274,%r308,%r276,%r308,%r308,%r308,%r308,%r308,%r308,%r308,%r308,%r308,%r308,%r308,%r308,%r308,%r308,%r308,%r308,%r308,%r308,%r308,%r308,%r308,%r308,%r308,%r308,%r308,%r308,%r308,%r308,%r308,%r308,%r308,%r308);
	// end inline asm
	mov.b32 	%f86, %r238;
	mov.b32 	%f87, %r239;
	mov.b32 	%f88, %r240;
	add.f32 	%f282, %f86, %f87;
	add.f32 	%f283, %f282, %f88;
	setp.eq.f32 	%p21, %f283, 0f00000000;
	@%p21 bra 	$L__BB0_15;

	mul.f32 	%f284, %f87, %f87;
	fma.rn.f32 	%f285, %f86, %f86, %f284;
	fma.rn.f32 	%f286, %f88, %f88, %f285;
	sqrt.rn.f32 	%f287, %f286;
	rcp.rn.f32 	%f288, %f287;
	mul.f32 	%f289, %f288, %f86;
	mul.f32 	%f290, %f288, %f87;
	mul.f32 	%f291, %f288, %f88;
	fma.rn.f32 	%f292, %f41, %f289, %f182;
	fma.rn.f32 	%f293, %f41, %f290, %f183;
	fma.rn.f32 	%f294, %f41, %f291, %f184;
	setp.lt.f32 	%p22, %f287, %f403;
	fma.rn.f32 	%f295, %f287, %f273, %f292;
	fma.rn.f32 	%f296, %f287, %f274, %f293;
	fma.rn.f32 	%f297, %f287, %f275, %f294;
	selp.f32 	%f403, %f287, %f403, %p22;
	selp.f32 	%f406, %f297, %f406, %p22;
	selp.f32 	%f405, %f296, %f405, %p22;
	selp.f32 	%f404, %f295, %f404, %p22;

$L__BB0_15:
	fma.rn.f32 	%f307, %f32, %f38, %f80;
	fma.rn.f32 	%f308, %f33, %f38, %f81;
	fma.rn.f32 	%f309, %f34, %f38, %f82;
	add.f32 	%f301, %f307, %f48;
	add.f32 	%f302, %f308, %f49;
	add.f32 	%f303, %f309, %f50;
	// begin inline asm
	call(%r309,%r310,%r311,%r312,%r313,%r314,%r315,%r316,%r317,%r318,%r319,%r320,%r321,%r322,%r323,%r324,%r325,%r326,%r327,%r328,%r329,%r330,%r331,%r332,%r333,%r334,%r335,%r336,%r337,%r338,%r339,%r340),_optix_trace_typed_32,(%r308,%rd41,%f182,%f183,%f184,%f301,%f302,%f303,%f278,%f189,%f278,%r271,%r308,%r308,%r274,%r308,%r276,%r308,%r308,%r308,%r308,%r308,%r308,%r308,%r308,%r308,%r308,%r308,%r308,%r308,%r308,%r308,%r308,%r308,%r308,%r308,%r308,%r308,%r308,%r308,%r308,%r308,%r308,%r308,%r308,%r308,%r308,%r308,%r308);
	// end inline asm
	mov.b32 	%f100, %r309;
	mov.b32 	%f101, %r310;
	mov.b32 	%f102, %r311;
	add.f32 	%f310, %f100, %f101;
	add.f32 	%f311, %f310, %f102;
	setp.eq.f32 	%p23, %f311, 0f00000000;
	@%p23 bra 	$L__BB0_17;

	mul.f32 	%f312, %f101, %f101;
	fma.rn.f32 	%f313, %f100, %f100, %f312;
	fma.rn.f32 	%f314, %f102, %f102, %f313;
	sqrt.rn.f32 	%f315, %f314;
	rcp.rn.f32 	%f316, %f315;
	mul.f32 	%f317, %f316, %f100;
	mul.f32 	%f318, %f316, %f101;
	mul.f32 	%f319, %f316, %f102;
	fma.rn.f32 	%f320, %f41, %f317, %f182;
	fma.rn.f32 	%f321, %f41, %f318, %f183;
	fma.rn.f32 	%f322, %f41, %f319, %f184;
	setp.lt.f32 	%p24, %f315, %f403;
	fma.rn.f32 	%f323, %f315, %f301, %f320;
	fma.rn.f32 	%f324, %f315, %f302, %f321;
	fma.rn.f32 	%f325, %f315, %f303, %f322;
	selp.f32 	%f406, %f325, %f406, %p24;
	selp.f32 	%f405, %f324, %f405, %p24;
	selp.f32 	%f404, %f323, %f404, %p24;

$L__BB0_17:
	abs.f32 	%f326, %f404;
	setp.le.f32 	%p25, %f326, 0f7F800000;
	selp.f32 	%f414, %f404, 0f00000000, %p25;
	abs.f32 	%f327, %f405;
	setp.le.f32 	%p26, %f327, 0f7F800000;
	selp.f32 	%f415, %f405, 0f00000000, %p26;
	abs.f32 	%f328, %f406;
	setp.le.f32 	%p27, %f328, 0f7F800000;
	selp.f32 	%f416, %f406, 0f00000000, %p27;
	ld.const.f32 	%f112, [params+720];
	setp.eq.f32 	%p28, %f112, 0f00000000;
	@%p28 bra 	$L__BB0_19;

	fma.rn.f32 	%f414, %f396, %f112, %f414;
	fma.rn.f32 	%f415, %f397, %f112, %f415;
	fma.rn.f32 	%f416, %f398, %f112, %f416;

$L__BB0_19:
	ld.const.u64 	%rd25, [params+224];
	cvta.to.global.u64 	%rd26, %rd25;
	ld.const.u32 	%r380, [params+216];
	mad.lo.s32 	%r381, %r380, %r4, %r3;
	mul.wide.u32 	%rd27, %r381, 16;
	add.s64 	%rd28, %rd26, %rd27;
	mov.f32 	%f329, 0f3F800000;
	st.global.v4.f32 	[%rd28], {%f414, %f415, %f416, %f329};
	ld.const.u8 	%rs16, [params+104];
	and.b16  	%rs17, %rs16, 32;
	setp.eq.s16 	%p29, %rs17, 0;
	mov.f32 	%f417, 0f437F0000;
	@%p29 bra 	$L__BB0_27;

	add.f32 	%f395, %f141, 0f38D1B717;
	mul.f32 	%f337, %f395, 0f41000000;
	mov.f32 	%f336, 0f38D1B717;
	mov.f32 	%f338, 0f00000000;
	mov.u32 	%r415, 1;
	mov.u32 	%r418, 2;
	mov.u32 	%r420, 3;
	mov.u32 	%r452, 0;
	// begin inline asm
	call(%r382,%r383,%r384,%r385,%r386,%r387,%r388,%r389,%r390,%r391,%r392,%r393,%r394,%r395,%r396,%r397,%r398,%r399,%r400,%r401,%r402,%r403,%r404,%r405,%r406,%r407,%r408,%r409,%r410,%r411,%r412,%r413),_optix_trace_typed_32,(%r452,%rd41,%f182,%f183,%f184,%f185,%f186,%f187,%f336,%f337,%f338,%r415,%r452,%r452,%r418,%r452,%r420,%r452,%r452,%r452,%r452,%r452,%r452,%r452,%r452,%r452,%r452,%r452,%r452,%r452,%r452,%r452,%r452,%r452,%r452,%r452,%r452,%r452,%r452,%r452,%r452,%r452,%r452,%r452,%r452,%r452,%r452,%r452,%r452);
	// end inline asm
	mov.b32 	%f340, %r382;
	mov.b32 	%f341, %r383;
	mov.b32 	%f342, %r384;
	add.f32 	%f343, %f340, %f341;
	add.f32 	%f344, %f343, %f342;
	setp.neu.f32 	%p30, %f344, 0f00000000;
	@%p30 bra 	$L__BB0_25;

	// begin inline asm
	call(%r453,%r454,%r455,%r456,%r457,%r458,%r459,%r460,%r461,%r462,%r463,%r464,%r465,%r466,%r467,%r468,%r469,%r470,%r471,%r472,%r473,%r474,%r475,%r476,%r477,%r478,%r479,%r480,%r481,%r482,%r483,%r484),_optix_trace_typed_32,(%r452,%rd41,%f182,%f183,%f184,%f245,%f246,%f247,%f336,%f337,%f338,%r415,%r452,%r452,%r418,%r452,%r420,%r452,%r452,%r452,%r452,%r452,%r452,%r452,%r452,%r452,%r452,%r452,%r452,%r452,%r452,%r452,%r452,%r452,%r452,%r452,%r452,%r452,%r452,%r452,%r452,%r452,%r452,%r452,%r452,%r452,%r452,%r452,%r452);
	// end inline asm
	mov.b32 	%f355, %r453;
	mov.b32 	%f356, %r454;
	mov.b32 	%f357, %r455;
	add.f32 	%f358, %f355, %f356;
	add.f32 	%f359, %f358, %f357;
	setp.neu.f32 	%p31, %f359, 0f00000000;
	@%p31 bra 	$L__BB0_25;

	mov.f32 	%f366, 0f38D1B717;
	mov.f32 	%f368, 0f00000000;
	mov.u32 	%r557, 1;
	mov.u32 	%r560, 2;
	mov.u32 	%r562, 3;
	mov.u32 	%r594, 0;
	// begin inline asm
	call(%r524,%r525,%r526,%r527,%r528,%r529,%r530,%r531,%r532,%r533,%r534,%r535,%r536,%r537,%r538,%r539,%r540,%r541,%r542,%r543,%r544,%r545,%r546,%r547,%r548,%r549,%r550,%r551,%r552,%r553,%r554,%r555),_optix_trace_typed_32,(%r594,%rd41,%f182,%f183,%f184,%f273,%f274,%f275,%f366,%f337,%f368,%r557,%r594,%r594,%r560,%r594,%r562,%r594,%r594,%r594,%r594,%r594,%r594,%r594,%r594,%r594,%r594,%r594,%r594,%r594,%r594,%r594,%r594,%r594,%r594,%r594,%r594,%r594,%r594,%r594,%r594,%r594,%r594,%r594,%r594,%r594,%r594,%r594,%r594);
	// end inline asm
	mov.b32 	%f370, %r524;
	mov.b32 	%f371, %r525;
	mov.b32 	%f372, %r526;
	add.f32 	%f373, %f370, %f371;
	add.f32 	%f374, %f373, %f372;
	setp.neu.f32 	%p32, %f374, 0f00000000;
	@%p32 bra 	$L__BB0_25;

	// begin inline asm
	call(%r595,%r596,%r597,%r598,%r599,%r600,%r601,%r602,%r603,%r604,%r605,%r606,%r607,%r608,%r609,%r610,%r611,%r612,%r613,%r614,%r615,%r616,%r617,%r618,%r619,%r620,%r621,%r622,%r623,%r624,%r625,%r626),_optix_trace_typed_32,(%r594,%rd41,%f182,%f183,%f184,%f301,%f302,%f303,%f366,%f337,%f368,%r557,%r594,%r594,%r560,%r594,%r562,%r594,%r594,%r594,%r594,%r594,%r594,%r594,%r594,%r594,%r594,%r594,%r594,%r594,%r594,%r594,%r594,%r594,%r594,%r594,%r594,%r594,%r594,%r594,%r594,%r594,%r594,%r594,%r594,%r594,%r594,%r594,%r594);
	// end inline asm
	mov.b32 	%f385, %r595;
	mov.b32 	%f386, %r596;
	mov.b32 	%f387, %r597;
	add.f32 	%f388, %f385, %f386;
	add.f32 	%f389, %f388, %f387;
	setp.neu.f32 	%p33, %f389, 0f00000000;
	@%p33 bra 	$L__BB0_25;

	mov.f32 	%f417, 0f00000000;

$L__BB0_25:
	ld.const.u64 	%rd33, [params+192];
	cvta.to.global.u64 	%rd34, %rd33;
	ld.const.u32 	%r666, [params+184];
	mad.lo.s32 	%r667, %r666, %r4, %r3;
	cvt.u64.u32 	%rd35, %r667;
	cvt.rzi.u32.f32 	%r668, %f417;
	add.s64 	%rd36, %rd34, %rd35;
	st.global.u8 	[%rd36], %r668;

$L__BB0_27:
	ret;

}