169 lines
4.0 KiB
Plaintext
169 lines
4.0 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<3>;
|
|
.reg .b16 %rs<12>;
|
|
.reg .f32 %f<55>;
|
|
.reg .b32 %r<24>;
|
|
.reg .b64 %rd<3>;
|
|
|
|
|
|
// begin inline asm
|
|
call (%r2), _optix_read_primitive_idx, ();
|
|
// end inline asm
|
|
// begin inline asm
|
|
call (%f10, %f11), _optix_get_triangle_barycentrics, ();
|
|
// end inline asm
|
|
mov.u32 %r3, 0;
|
|
// begin inline asm
|
|
call (%rd1), _optix_call_direct_callable,(%r3);
|
|
// end inline asm
|
|
{ // callseq 0, 0
|
|
.reg .b32 temp_param_reg;
|
|
.param .b32 param0;
|
|
st.param.b32 [param0+0], %r2;
|
|
.param .align 8 .b8 param1[8];
|
|
st.param.v2.f32 [param1+0], {%f10, %f11};
|
|
.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 {%f4, %f5}, [retval0+16];
|
|
ld.param.f32 %f6, [retval0+24];
|
|
ld.param.v4.b8 {%rs5, %rs6, %rs7, %rs8}, [retval0+28];
|
|
} // callseq 0
|
|
ld.const.u32 %r1, [params+112];
|
|
ld.const.f32 %f7, [params+700];
|
|
// begin inline asm
|
|
call (%f12), _optix_get_ray_tmax, ();
|
|
// end inline asm
|
|
// begin inline asm
|
|
call (%f13), _optix_get_world_ray_direction_x, ();
|
|
// end inline asm
|
|
// begin inline asm
|
|
call (%f14), _optix_get_world_ray_direction_y, ();
|
|
// end inline asm
|
|
// begin inline asm
|
|
call (%f15), _optix_get_world_ray_direction_z, ();
|
|
// end inline asm
|
|
mul.f32 %f16, %f2, %f14;
|
|
fma.rn.f32 %f17, %f1, %f13, %f16;
|
|
fma.rn.f32 %f18, %f3, %f15, %f17;
|
|
setp.ltu.f32 %p1, %f18, 0f00000000;
|
|
@%p1 bra $L__BB0_3;
|
|
bra.uni $L__BB0_1;
|
|
|
|
$L__BB0_3:
|
|
neg.f32 %f54, %f12;
|
|
mov.b32 %r23, %f54;
|
|
// begin inline asm
|
|
call _optix_set_payload, (%r3, %r23);
|
|
// end inline asm
|
|
mov.u32 %r20, 1;
|
|
// begin inline asm
|
|
call _optix_set_payload, (%r20, %r23);
|
|
// end inline asm
|
|
mov.u32 %r22, 2;
|
|
// begin inline asm
|
|
call _optix_set_payload, (%r22, %r23);
|
|
// end inline asm
|
|
bra.uni $L__BB0_4;
|
|
|
|
$L__BB0_1:
|
|
div.rn.f32 %f19, %f4, 0f41200000;
|
|
cvt.rzi.s32.f32 %r4, %f19;
|
|
setp.ne.s32 %p2, %r4, %r1;
|
|
@%p2 bra $L__BB0_4;
|
|
|
|
cvt.rzi.s32.f32 %r11, %f4;
|
|
cvt.rn.f32.s32 %f23, %r11;
|
|
sub.f32 %f24, %f4, %f23;
|
|
ld.const.u64 %rd2, [params+1120];
|
|
tex.2d.v4.u32.f32 {%r12, %r13, %r14, %r15}, [%rd2, {%f24, %f5}];
|
|
cvt.u16.u32 %rs9, %r12;
|
|
// begin inline asm
|
|
{ cvt.f32.f16 %f20, %rs9;}
|
|
|
|
// end inline asm
|
|
cvt.u16.u32 %rs10, %r13;
|
|
// begin inline asm
|
|
{ cvt.f32.f16 %f21, %rs10;}
|
|
|
|
// end inline asm
|
|
cvt.u16.u32 %rs11, %r14;
|
|
// begin inline asm
|
|
{ cvt.f32.f16 %f22, %rs11;}
|
|
|
|
// end inline asm
|
|
mul.f32 %f25, %f7, %f12;
|
|
neg.f32 %f26, %f25;
|
|
mov.f32 %f27, 0f3F000000;
|
|
mov.f32 %f28, 0f3BBB989D;
|
|
fma.rn.f32 %f29, %f26, %f28, %f27;
|
|
mov.f32 %f30, 0f3FB8AA3B;
|
|
mov.f32 %f31, 0f437C0000;
|
|
cvt.sat.f32.f32 %f32, %f29;
|
|
mov.f32 %f33, 0f4B400001;
|
|
fma.rm.f32 %f34, %f32, %f31, %f33;
|
|
add.f32 %f35, %f34, 0fCB40007F;
|
|
neg.f32 %f36, %f35;
|
|
fma.rn.f32 %f37, %f26, %f30, %f36;
|
|
mov.f32 %f38, 0f32A57060;
|
|
fma.rn.f32 %f39, %f26, %f38, %f37;
|
|
mov.b32 %r16, %f34;
|
|
shl.b32 %r17, %r16, 23;
|
|
mov.b32 %f40, %r17;
|
|
ex2.approx.ftz.f32 %f41, %f39;
|
|
mul.f32 %f42, %f41, %f40;
|
|
mul.f32 %f43, %f20, %f42;
|
|
mul.f32 %f44, %f21, %f42;
|
|
mul.f32 %f45, %f22, %f42;
|
|
ld.const.v2.f32 {%f46, %f47}, [params+648];
|
|
mov.u32 %r5, 0;
|
|
mov.u32 %r7, 1;
|
|
mul.f32 %f50, %f46, %f43;
|
|
mul.f32 %f51, %f47, %f44;
|
|
ld.const.f32 %f52, [params+656];
|
|
mul.f32 %f53, %f52, %f45;
|
|
mov.b32 %r6, %f50;
|
|
// begin inline asm
|
|
call _optix_set_payload, (%r5, %r6);
|
|
// end inline asm
|
|
mov.b32 %r8, %f51;
|
|
// begin inline asm
|
|
call _optix_set_payload, (%r7, %r8);
|
|
// end inline asm
|
|
mov.b32 %r10, %f53;
|
|
mov.u32 %r9, 2;
|
|
// begin inline asm
|
|
call _optix_set_payload, (%r9, %r10);
|
|
// end inline asm
|
|
|
|
$L__BB0_4:
|
|
ret;
|
|
|
|
}
|
|
|