137 lines
3.0 KiB
Plaintext
137 lines
3.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__dummy
|
|
.const .align 16 .b8 params[1184];
|
|
|
|
.visible .entry __closesthit__dummy()
|
|
{
|
|
|
|
|
|
|
|
ret;
|
|
|
|
}
|
|
// .globl __anyhit__oxMain
|
|
.visible .entry __anyhit__oxMain()
|
|
{
|
|
.reg .pred %p<3>;
|
|
.reg .b16 %rs<9>;
|
|
.reg .f32 %f<24>;
|
|
.reg .b32 %r<24>;
|
|
.reg .b64 %rd<2>;
|
|
|
|
|
|
// begin inline asm
|
|
call (%r1), _optix_read_primitive_idx, ();
|
|
// end inline asm
|
|
// begin inline asm
|
|
call (%f2, %f3), _optix_get_triangle_barycentrics, ();
|
|
// end inline asm
|
|
mov.u32 %r2, 0;
|
|
// begin inline asm
|
|
call (%rd1), _optix_call_direct_callable,(%r2);
|
|
// end inline asm
|
|
{ // callseq 0, 0
|
|
.reg .b32 temp_param_reg;
|
|
.param .b32 param0;
|
|
st.param.b32 [param0+0], %r1;
|
|
.param .align 8 .b8 param1[8];
|
|
st.param.v2.f32 [param1+0], {%f2, %f3};
|
|
.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 {%f4, %f5}, [retval0+0];
|
|
ld.param.f32 %f6, [retval0+8];
|
|
ld.param.v4.b8 {%rs1, %rs2, %rs3, %rs4}, [retval0+12];
|
|
ld.param.v2.f32 {%f7, %f8}, [retval0+16];
|
|
ld.param.f32 %f1, [retval0+24];
|
|
ld.param.v4.b8 {%rs5, %rs6, %rs7, %rs8}, [retval0+28];
|
|
} // callseq 0
|
|
setp.gtu.f32 %p1, %f1, 0f00000000;
|
|
@%p1 bra $L__BB1_4;
|
|
|
|
setp.geu.f32 %p2, %f1, 0f00000000;
|
|
@%p2 bra $L__BB1_3;
|
|
|
|
cvt.rzi.s32.f32 %r19, %f1;
|
|
not.b32 %r20, %r19;
|
|
and.b32 %r21, %r20, 255;
|
|
bfe.u32 %r22, %r20, 8, 8;
|
|
bfe.u32 %r23, %r20, 16, 8;
|
|
cvt.rn.f32.s32 %f10, %r21;
|
|
cvt.rn.f32.s32 %f11, %r22;
|
|
cvt.rn.f32.s32 %f12, %r23;
|
|
mul.f32 %f13, %f10, 0f3B808081;
|
|
mul.f32 %f14, %f11, 0f3B808081;
|
|
mul.f32 %f15, %f12, 0f3B808081;
|
|
mov.u32 %r11, 0;
|
|
// begin inline asm
|
|
call (%r3), _optix_get_payload, (%r11);
|
|
// end inline asm
|
|
mov.b32 %f16, %r3;
|
|
mov.u32 %r13, 1;
|
|
// begin inline asm
|
|
call (%r5), _optix_get_payload, (%r13);
|
|
// end inline asm
|
|
mov.b32 %f17, %r5;
|
|
mov.u32 %r15, 2;
|
|
// begin inline asm
|
|
call (%r7), _optix_get_payload, (%r15);
|
|
// end inline asm
|
|
mov.b32 %f18, %r7;
|
|
mov.u32 %r17, 3;
|
|
// begin inline asm
|
|
call (%r9), _optix_get_payload, (%r17);
|
|
// end inline asm
|
|
mov.b32 %f19, %r9;
|
|
mul.f32 %f20, %f13, %f16;
|
|
mul.f32 %f21, %f14, %f17;
|
|
mul.f32 %f22, %f15, %f18;
|
|
// begin inline asm
|
|
call (%f9), _optix_get_ray_tmax, ();
|
|
// end inline asm
|
|
min.f32 %f23, %f19, %f9;
|
|
mov.b32 %r12, %f20;
|
|
// begin inline asm
|
|
call _optix_set_payload, (%r11, %r12);
|
|
// end inline asm
|
|
mov.b32 %r14, %f21;
|
|
// begin inline asm
|
|
call _optix_set_payload, (%r13, %r14);
|
|
// end inline asm
|
|
mov.b32 %r16, %f22;
|
|
// begin inline asm
|
|
call _optix_set_payload, (%r15, %r16);
|
|
// end inline asm
|
|
mov.b32 %r18, %f23;
|
|
// begin inline asm
|
|
call _optix_set_payload, (%r17, %r18);
|
|
// end inline asm
|
|
|
|
$L__BB1_3:
|
|
// begin inline asm
|
|
call _optix_ignore_intersection, ();
|
|
// end inline asm
|
|
|
|
$L__BB1_4:
|
|
ret;
|
|
|
|
}
|
|
|