//
// 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	__closesthit__oxMain
.const .align 16 .b8 params[1184];

.visible .entry __closesthit__oxMain()
{
	.reg .pred 	%p<21>;
	.reg .b16 	%rs<9>;
	.reg .f32 	%f<152>;
	.reg .b32 	%r<51>;
	.reg .b64 	%rd<3>;


	// begin inline asm
	call (%r6), _optix_read_primitive_idx, ();
	// end inline asm
	// begin inline asm
	call (%f24, %f25), _optix_get_triangle_barycentrics, ();
	// end inline asm
	mov.u32 	%r21, 0;
	// begin inline asm
	call (%rd1), _optix_call_direct_callable,(%r21);
	// end inline asm
	{ // callseq 0, 0
	.reg .b32 temp_param_reg;
	.param .b32 param0;
	st.param.b32 	[param0+0], %r6;
	.param .align 8 .b8 param1[8];
	st.param.v2.f32 	[param1+0], {%f24, %f25};
	.param .align 8 .b8 retval0[32];
	prototype_0 : .callprototype (.param .align 8 .b8 _[32]) _ (.param .b32 _, .param .align 8 .b8 _[8]);
	call (retval0), 
	%rd1, 
	(
	param0, 
	param1
	)
	, prototype_0;
	ld.param.v2.f32 	{%f1, %f2}, [retval0+0];
	ld.param.f32 	%f3, [retval0+8];
	ld.param.v4.b8 	{%rs1, %rs2, %rs3, %rs4}, [retval0+12];
	ld.param.v2.f32 	{%f7, %f5}, [retval0+16];
	ld.param.f32 	%f6, [retval0+24];
	ld.param.v4.b8 	{%rs5, %rs6, %rs7, %rs8}, [retval0+28];
	} // callseq 0
	mov.u32 	%r18, 1;
	div.rn.f32 	%f30, %f7, 0f41200000;
	mov.b32 	%r24, %f30;
	and.b32  	%r25, %r24, -2147483648;
	or.b32  	%r26, %r25, 1056964608;
	mov.b32 	%f31, %r26;
	add.rz.f32 	%f32, %f30, %f31;
	cvt.rzi.f32.f32 	%f33, %f32;
	cvt.rzi.s32.f32 	%r1, %f33;
	// begin inline asm
	call (%f26), _optix_get_world_ray_direction_x, ();
	// end inline asm
	// begin inline asm
	call (%f27), _optix_get_world_ray_direction_y, ();
	// end inline asm
	// begin inline asm
	call (%f28), _optix_get_world_ray_direction_z, ();
	// end inline asm
	mul.f32 	%f34, %f2, %f27;
	fma.rn.f32 	%f35, %f1, %f26, %f34;
	fma.rn.f32 	%f36, %f3, %f28, %f35;
	setp.gt.f32 	%p2, %f36, 0f00000000;
	// begin inline asm
	call (%f29), _optix_get_ray_tmax, ();
	// end inline asm
	// begin inline asm
	call (%r8), _optix_get_payload, (%r21);
	// end inline asm
	// begin inline asm
	call (%r10), _optix_get_payload, (%r18);
	// end inline asm
	mov.u32 	%r20, 2;
	// begin inline asm
	call (%r12), _optix_get_payload, (%r20);
	// end inline asm
	mov.u32 	%r22, 3;
	// begin inline asm
	call (%r14), _optix_get_payload, (%r22);
	// end inline asm
	// begin inline asm
	call _optix_set_payload, (%r21, %r21);
	// end inline asm
	// begin inline asm
	call _optix_set_payload, (%r18, %r21);
	// end inline asm
	// begin inline asm
	call _optix_set_payload, (%r20, %r21);
	// end inline asm
	mov.u32 	%r23, 2139095039;
	// begin inline asm
	call _optix_set_payload, (%r22, %r23);
	// end inline asm
	ld.const.u32 	%r27, [params+664];
	setp.ne.s32 	%p3, %r1, %r27;
	or.pred  	%p4, %p2, %p3;
	@%p4 bra 	$L__BB0_14;

	mov.b32 	%f9, %r14;
	ld.const.f32 	%f38, [params+660];
	mul.f32 	%f11, %f29, %f38;
	mov.f32 	%f39, 0f40000000;
	cvt.rzi.f32.f32 	%f40, %f39;
	add.f32 	%f41, %f40, %f40;
	mov.f32 	%f42, 0f40800000;
	sub.f32 	%f43, %f42, %f41;
	abs.f32 	%f12, %f43;
	abs.f32 	%f13, %f11;
	setp.lt.f32 	%p5, %f13, 0f00800000;
	mul.f32 	%f44, %f13, 0f4B800000;
	selp.f32 	%f45, %f44, %f13, %p5;
	selp.f32 	%f46, 0fC3170000, 0fC2FE0000, %p5;
	mov.b32 	%r28, %f45;
	and.b32  	%r29, %r28, 8388607;
	or.b32  	%r30, %r29, 1065353216;
	mov.b32 	%f47, %r30;
	shr.u32 	%r31, %r28, 23;
	cvt.rn.f32.u32 	%f48, %r31;
	add.f32 	%f49, %f46, %f48;
	setp.gt.f32 	%p6, %f47, 0f3FB504F3;
	mul.f32 	%f50, %f47, 0f3F000000;
	add.f32 	%f51, %f49, 0f3F800000;
	selp.f32 	%f52, %f51, %f49, %p6;
	selp.f32 	%f53, %f50, %f47, %p6;
	add.f32 	%f54, %f53, 0fBF800000;
	add.f32 	%f55, %f53, 0f3F800000;
	rcp.approx.ftz.f32 	%f56, %f55;
	add.f32 	%f57, %f54, %f54;
	mul.f32 	%f58, %f57, %f56;
	mul.f32 	%f59, %f58, %f58;
	mov.f32 	%f60, 0f3C4CAF63;
	mov.f32 	%f61, 0f3B18F0FE;
	fma.rn.f32 	%f62, %f61, %f59, %f60;
	mov.f32 	%f63, 0f3DAAAABD;
	fma.rn.f32 	%f64, %f62, %f59, %f63;
	mul.rn.f32 	%f65, %f64, %f59;
	mul.rn.f32 	%f66, %f65, %f58;
	sub.f32 	%f67, %f54, %f58;
	add.f32 	%f68, %f67, %f67;
	neg.f32 	%f69, %f58;
	fma.rn.f32 	%f70, %f69, %f54, %f68;
	mul.rn.f32 	%f71, %f56, %f70;
	add.f32 	%f72, %f66, %f58;
	sub.f32 	%f73, %f58, %f72;
	add.f32 	%f74, %f66, %f73;
	add.f32 	%f75, %f71, %f74;
	add.f32 	%f76, %f72, %f75;
	sub.f32 	%f77, %f72, %f76;
	add.f32 	%f78, %f75, %f77;
	mov.f32 	%f79, 0f3F317200;
	mul.rn.f32 	%f80, %f52, %f79;
	mov.f32 	%f81, 0f35BFBE8E;
	mul.rn.f32 	%f82, %f52, %f81;
	add.f32 	%f83, %f80, %f76;
	sub.f32 	%f84, %f80, %f83;
	add.f32 	%f85, %f76, %f84;
	add.f32 	%f86, %f78, %f85;
	add.f32 	%f87, %f82, %f86;
	add.f32 	%f88, %f83, %f87;
	sub.f32 	%f89, %f83, %f88;
	add.f32 	%f90, %f87, %f89;
	mul.rn.f32 	%f91, %f42, %f88;
	neg.f32 	%f92, %f91;
	fma.rn.f32 	%f93, %f42, %f88, %f92;
	fma.rn.f32 	%f94, %f42, %f90, %f93;
	mov.f32 	%f95, 0f00000000;
	fma.rn.f32 	%f96, %f95, %f88, %f94;
	add.rn.f32 	%f97, %f91, %f96;
	neg.f32 	%f98, %f97;
	add.rn.f32 	%f99, %f91, %f98;
	add.rn.f32 	%f100, %f99, %f96;
	mov.b32 	%r32, %f97;
	setp.eq.s32 	%p7, %r32, 1118925336;
	add.s32 	%r33, %r32, -1;
	mov.b32 	%f101, %r33;
	add.f32 	%f102, %f100, 0f37000000;
	selp.f32 	%f14, %f102, %f100, %p7;
	selp.f32 	%f103, %f101, %f97, %p7;
	mov.f32 	%f104, 0f3FB8AA3B;
	mul.rn.f32 	%f105, %f103, %f104;
	cvt.rzi.f32.f32 	%f106, %f105;
	abs.f32 	%f107, %f106;
	setp.gt.f32 	%p8, %f107, 0f42FC0000;
	mov.b32 	%r34, %f106;
	and.b32  	%r35, %r34, -2147483648;
	or.b32  	%r36, %r35, 1123811328;
	mov.b32 	%f108, %r36;
	selp.f32 	%f109, %f108, %f106, %p8;
	mov.f32 	%f110, 0fBF317218;
	fma.rn.f32 	%f111, %f109, %f110, %f103;
	mov.f32 	%f112, 0f3102E308;
	fma.rn.f32 	%f113, %f109, %f112, %f111;
	mul.f32 	%f114, %f113, 0f3FB8AA3B;
	add.f32 	%f115, %f109, 0f4B40007F;
	mov.b32 	%r37, %f115;
	shl.b32 	%r38, %r37, 23;
	mov.b32 	%f116, %r38;
	ex2.approx.ftz.f32 	%f117, %f114;
	mul.f32 	%f15, %f117, %f116;
	setp.eq.f32 	%p9, %f15, 0f7F800000;
	mov.f32 	%f149, 0f7F800000;
	@%p9 bra 	$L__BB0_3;

	fma.rn.f32 	%f149, %f15, %f14, %f15;

$L__BB0_3:
	setp.lt.f32 	%p10, %f11, 0f00000000;
	setp.eq.f32 	%p11, %f12, 0f3F800000;
	and.pred  	%p1, %p10, %p11;
	setp.eq.f32 	%p12, %f11, 0f00000000;
	@%p12 bra 	$L__BB0_7;
	bra.uni 	$L__BB0_4;

$L__BB0_7:
	add.f32 	%f122, %f11, %f11;
	selp.f32 	%f151, %f122, 0f00000000, %p11;
	bra.uni 	$L__BB0_8;

$L__BB0_4:
	mov.b32 	%r39, %f149;
	xor.b32  	%r40, %r39, -2147483648;
	mov.b32 	%f118, %r40;
	selp.f32 	%f151, %f118, %f149, %p1;
	setp.geu.f32 	%p13, %f11, 0f00000000;
	@%p13 bra 	$L__BB0_8;

	mov.f32 	%f119, 0f40800000;
	cvt.rzi.f32.f32 	%f120, %f119;
	setp.eq.f32 	%p14, %f120, 0f40800000;
	@%p14 bra 	$L__BB0_8;

	mov.f32 	%f151, 0f7FFFFFFF;

$L__BB0_8:
	add.f32 	%f123, %f13, 0f40800000;
	mov.b32 	%r41, %f123;
	setp.lt.s32 	%p16, %r41, 2139095040;
	@%p16 bra 	$L__BB0_13;

	setp.gtu.f32 	%p17, %f13, 0f7F800000;
	@%p17 bra 	$L__BB0_12;
	bra.uni 	$L__BB0_10;

$L__BB0_12:
	add.f32 	%f151, %f11, 0f40800000;
	bra.uni 	$L__BB0_13;

$L__BB0_10:
	setp.neu.f32 	%p18, %f13, 0f7F800000;
	@%p18 bra 	$L__BB0_13;

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

$L__BB0_13:
	mov.f32 	%f124, 0f3F800000;
	sub.f32 	%f125, %f124, %f151;
	setp.eq.f32 	%p19, %f11, 0f3F800000;
	selp.f32 	%f126, 0f00000000, %f125, %p19;
	cvt.sat.f32.f32 	%f127, %f126;
	ld.const.u64 	%rd2, [params+552];
	mul.lo.s32 	%r50, %r1, 10;
	cvt.rn.f32.s32 	%f128, %r50;
	sub.f32 	%f129, %f7, %f128;
	mul.f32 	%f130, %f129, 0fC1200000;
	mul.f32 	%f131, %f5, 0f41200000;
	tex.2d.v4.f32.f32 	{%f132, %f133, %f134, %f135}, [%rd2, {%f130, %f131}];
	mul.f32 	%f136, %f127, %f132;
	mul.f32 	%f137, %f127, %f133;
	mul.f32 	%f138, %f127, %f134;
	setp.gt.f32 	%p20, %f29, %f9;
	mov.b32 	%f139, %r12;
	selp.f32 	%f140, %f139, 0f3F800000, %p20;
	mov.b32 	%f141, %r10;
	selp.f32 	%f142, %f141, 0f3F800000, %p20;
	mov.b32 	%f143, %r8;
	selp.f32 	%f144, %f143, 0f3F800000, %p20;
	mul.f32 	%f145, %f136, %f144;
	mul.f32 	%f146, %f137, %f142;
	mul.f32 	%f147, %f138, %f140;
	mov.b32 	%r43, %f145;
	mov.u32 	%r42, 0;
	// begin inline asm
	call _optix_set_payload, (%r42, %r43);
	// end inline asm
	mov.b32 	%r45, %f146;
	mov.u32 	%r44, 1;
	// begin inline asm
	call _optix_set_payload, (%r44, %r45);
	// end inline asm
	mov.b32 	%r47, %f147;
	mov.u32 	%r46, 2;
	// begin inline asm
	call _optix_set_payload, (%r46, %r47);
	// end inline asm
	neg.f32 	%f148, %f29;
	mov.b32 	%r49, %f148;
	mov.u32 	%r48, 3;
	// begin inline asm
	call _optix_set_payload, (%r48, %r49);
	// end inline asm

$L__BB0_14:
	ret;

}