// // Generated by NVIDIA NVVM Compiler // // Compiler Build ID: CL-31833905 // Cuda compilation tools, release 11.8, V11.8.89 // Based on NVVM 7.0.1 // .version 7.8 .target sm_52 .address_size 64 // .globl __raygen__rg .const .align 8 .b8 params[32]; .global .align 4 .f32 _ZZN4tcnn19gaussian_cdf_approxEffE20MAGIC_SIGMOID_FACTOR = 0f3F4ABDDD; .global .align 4 .f32 _ZZN4tcnn30gaussian_cdf_approx_derivativeEffE20MAGIC_SIGMOID_FACTOR = 0f3F4ABDDD; .global .align 4 .b8 _ZZN3ngp5sobolEjjE10directions[640] = {0, 0, 0, 128, 0, 0, 0, 64, 0, 0, 0, 32, 0, 0, 0, 16, 0, 0, 0, 8, 0, 0, 0, 4, 0, 0, 0, 2, 0, 0, 0, 1, 0, 0, 128, 0, 0, 0, 64, 0, 0, 0, 32, 0, 0, 0, 16, 0, 0, 0, 8, 0, 0, 0, 4, 0, 0, 0, 2, 0, 0, 0, 1, 0, 0, 128, 0, 0, 0, 64, 0, 0, 0, 32, 0, 0, 0, 16, 0, 0, 0, 8, 0, 0, 0, 4, 0, 0, 0, 2, 0, 0, 0, 1, 0, 0, 128, 0, 0, 0, 64, 0, 0, 0, 32, 0, 0, 0, 16, 0, 0, 0, 8, 0, 0, 0, 4, 0, 0, 0, 2, 0, 0, 0, 1, 0, 0, 0, 0, 0, 0, 128, 0, 0, 0, 192, 0, 0, 0, 160, 0, 0, 0, 240, 0, 0, 0, 136, 0, 0, 0, 204, 0, 0, 0, 170, 0, 0, 0, 255, 0, 0, 128, 128, 0, 0, 192, 192, 0, 0, 160, 160, 0, 0, 240, 240, 0, 0, 136, 136, 0, 0, 204, 204, 0, 0, 170, 170, 0, 0, 255, 255, 0, 128, 0, 128, 0, 192, 0, 192, 0, 160, 0, 160, 0, 240, 0, 240, 0, 136, 0, 136, 0, 204, 0, 204, 0, 170, 0, 170, 0, 255, 0, 255, 128, 128, 128, 128, 192, 192, 192, 192, 160, 160, 160, 160, 240, 240, 240, 240, 136, 136, 136, 136, 204, 204, 204, 204, 170, 170, 170, 170, 255, 255, 255, 255, 0, 0, 0, 128, 0, 0, 0, 192, 0, 0, 0, 96, 0, 0, 0, 144, 0, 0, 0, 232, 0, 0, 0, 92, 0, 0, 0, 142, 0, 0, 0, 197, 0, 0, 128, 104, 0, 0, 192, 156, 0, 0, 96, 238, 0, 0, 144, 85, 0, 0, 104, 128, 0, 0, 156, 192, 0, 0, 238, 96, 0, 0, 85, 144, 0, 128, 128, 232, 0, 192, 192, 92, 0, 96, 96, 142, 0, 144, 144, 197, 0, 232, 104, 104, 0, 92, 156, 156, 0, 142, 238, 238, 0, 197, 85, 85, 128, 232, 0, 128, 192, 92, 0, 192, 96, 142, 0, 96, 144, 197, 0, 144, 104, 104, 0, 232, 156, 156, 0, 92, 238, 238, 0, 142, 85, 85, 0, 197, 0, 0, 0, 128, 0, 0, 0, 192, 0, 0, 0, 32, 0, 0, 0, 80, 0, 0, 0, 248, 0, 0, 0, 116, 0, 0, 0, 162, 0, 0, 0, 147, 0, 0, 128, 216, 0, 0, 64, 37, 0, 0, 224, 89, 0, 0, 208, 230, 0, 0, 8, 120, 0, 0, 12, 180, 0, 0, 2, 130, 0, 0, 5, 195, 0, 128, 143, 32, 0, 64, 71, 81, 0, 32, 234, 251, 0, 48, 217, 117, 0, 136, 133, 160, 0, 84, 78, 145, 0, 158, 231, 219, 0, 109, 219, 37, 128, 0, 128, 88, 192, 0, 64, 229, 32, 0, 224, 121, 80, 0, 208, 182, 248, 0, 8, 128, 116, 0, 12, 192, 162, 0, 2, 32, 147, 0, 5, 80, 0, 0, 0, 128, 0, 0, 0, 64, 0, 0, 0, 32, 0, 0, 0, 176, 0, 0, 0, 248, 0, 0, 0, 220, 0, 0, 0, 122, 0, 0, 0, 157, 0, 0, 128, 90, 0, 0, 192, 47, 0, 0, 96, 161, 0, 0, 176, 240, 0, 0, 136, 218, 0, 0, 196, 111, 0, 0, 98, 129, 0, 0, 187, 64, 0, 128, 135, 34, 0, 192, 201, 179, 0, 160, 101, 251, 0, 208, 178, 221, 0, 40, 2, 120, 0, 60, 11, 156, 0, 182, 15, 90, 0, 219, 13, 45, 128, 128, 135, 162, 64, 192, 201, 243, 32, 160, 101, 219, 176, 208, 178, 109, 248, 40, 2, 128, 220, 60, 11, 64, 122, 182, 15, 32, 157, 219, 13, 176}; .visible .entry __raygen__rg() { .reg .pred %p<2>; .reg .f32 %f<46>; .reg .b32 %r<81>; .reg .b64 %rd<13>; // begin inline asm call (%r2), _optix_get_launch_index_x, (); // end inline asm ld.const.u64 %rd3, [params]; cvta.to.global.u64 %rd4, %rd3; mul.wide.u32 %rd5, %r2, 12; add.s64 %rd6, %rd4, %rd5; ld.global.f32 %f1, [%rd6]; ld.global.f32 %f2, [%rd6+4]; ld.global.f32 %f3, [%rd6+8]; ld.const.u64 %rd7, [params+8]; cvta.to.global.u64 %rd8, %rd7; add.s64 %rd1, %rd8, %rd5; ld.global.f32 %f4, [%rd1]; ld.global.f32 %f5, [%rd1+4]; ld.global.f32 %f6, [%rd1+8]; ld.const.u64 %rd2, [params+24]; mov.f32 %f8, 0f5A0E1BCA; mov.f32 %f9, 0f00000000; mov.u32 %r41, 255; mov.u32 %r44, 1; mov.u32 %r46, 2; mov.u32 %r78, 0; // begin inline asm call(%r8,%r9,%r10,%r11,%r12,%r13,%r14,%r15,%r16,%r17,%r18,%r19,%r20,%r21,%r22,%r23,%r24,%r25,%r26,%r27,%r28,%r29,%r30,%r31,%r32,%r33,%r34,%r35,%r36,%r37,%r38,%r39),_optix_trace_typed_32,(%r78,%rd2,%f1,%f2,%f3,%f4,%f5,%f6,%f9,%f8,%f9,%r41,%r44,%r78,%r44,%r78,%r46,%r79,%r80,%r78,%r78,%r78,%r78,%r78,%r78,%r78,%r78,%r78,%r78,%r78,%r78,%r78,%r78,%r78,%r78,%r78,%r78,%r78,%r78,%r78,%r78,%r78,%r78,%r78,%r78,%r78,%r78,%r78,%r78); // end inline asm mov.b32 %f10, %r9; fma.rn.f32 %f11, %f4, %f10, %f1; fma.rn.f32 %f12, %f5, %f10, %f2; fma.rn.f32 %f13, %f6, %f10, %f3; st.global.f32 [%rd6], %f11; st.global.f32 [%rd6+4], %f12; st.global.f32 [%rd6+8], %f13; setp.eq.s32 %p1, %r8, -1; @%p1 bra $L__BB0_2; ld.const.u64 %rd9, [params+16]; cvta.to.global.u64 %rd10, %rd9; mul.wide.u32 %rd11, %r8, 36; add.s64 %rd12, %rd10, %rd11; ld.global.f32 %f14, [%rd12]; ld.global.f32 %f15, [%rd12+12]; sub.f32 %f16, %f15, %f14; ld.global.f32 %f17, [%rd12+4]; ld.global.f32 %f18, [%rd12+16]; sub.f32 %f19, %f18, %f17; ld.global.f32 %f20, [%rd12+8]; ld.global.f32 %f21, [%rd12+20]; sub.f32 %f22, %f21, %f20; ld.global.f32 %f23, [%rd12+24]; sub.f32 %f24, %f23, %f14; ld.global.f32 %f25, [%rd12+28]; sub.f32 %f26, %f25, %f17; ld.global.f32 %f27, [%rd12+32]; sub.f32 %f28, %f27, %f20; mul.f32 %f29, %f19, %f28; mul.f32 %f30, %f22, %f26; sub.f32 %f31, %f29, %f30; mul.f32 %f32, %f22, %f24; mul.f32 %f33, %f16, %f28; sub.f32 %f34, %f32, %f33; mul.f32 %f35, %f16, %f26; mul.f32 %f36, %f19, %f24; sub.f32 %f37, %f35, %f36; mul.f32 %f38, %f34, %f34; fma.rn.f32 %f39, %f31, %f31, %f38; fma.rn.f32 %f40, %f37, %f37, %f39; sqrt.rn.f32 %f41, %f40; rcp.rn.f32 %f42, %f41; mul.f32 %f43, %f42, %f31; mul.f32 %f44, %f42, %f34; mul.f32 %f45, %f37, %f42; st.global.f32 [%rd1], %f43; st.global.f32 [%rd1+4], %f44; st.global.f32 [%rd1+8], %f45; $L__BB0_2: ret; } // .globl __miss__ms .visible .entry __miss__ms() { .reg .f32 %f<2>; .reg .b32 %r<5>; mov.u32 %r1, 0; mov.u32 %r2, -1; // begin inline asm call _optix_set_payload, (%r1, %r2); // end inline asm // begin inline asm call (%f1), _optix_get_ray_tmax, (); // end inline asm mov.b32 %r4, %f1; mov.u32 %r3, 1; // begin inline asm call _optix_set_payload, (%r3, %r4); // end inline asm ret; } // .globl __closesthit__ch .visible .entry __closesthit__ch() { .reg .f32 %f<2>; .reg .b32 %r<6>; // begin inline asm call (%r1), _optix_read_primitive_idx, (); // end inline asm mov.u32 %r2, 0; // begin inline asm call _optix_set_payload, (%r2, %r1); // end inline asm // begin inline asm call (%f1), _optix_get_ray_tmax, (); // end inline asm mov.b32 %r5, %f1; mov.u32 %r4, 1; // begin inline asm call _optix_set_payload, (%r4, %r5); // end inline asm ret; }