322 lines
8.4 KiB
Plaintext
322 lines
8.4 KiB
Plaintext
//
|
|
// 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;
|
|
|
|
}
|
|
|