520 lines
18 KiB
Plaintext
520 lines
18 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 __raygen__oxMain
|
|
.const .align 16 .b8 params[1184];
|
|
|
|
.visible .entry __raygen__oxMain()
|
|
{
|
|
.reg .pred %p<34>;
|
|
.reg .b16 %rs<20>;
|
|
.reg .f32 %f<418>;
|
|
.reg .b32 %r<672>;
|
|
.reg .b64 %rd<42>;
|
|
|
|
|
|
// begin inline asm
|
|
call (%r7), _optix_get_launch_index_x, ();
|
|
// end inline asm
|
|
// begin inline asm
|
|
call (%r8), _optix_get_launch_index_y, ();
|
|
// end inline asm
|
|
ld.const.u64 %rd6, [params+416];
|
|
cvta.to.global.u64 %rd7, %rd6;
|
|
ld.const.u32 %r13, [params+408];
|
|
mad.lo.s32 %r14, %r13, %r8, %r7;
|
|
mul.wide.u32 %rd8, %r14, 4;
|
|
add.s64 %rd1, %rd7, %rd8;
|
|
ld.global.v2.u8 {%rs7, %rs19}, [%rd1];
|
|
or.b16 %rs9, %rs7, %rs19;
|
|
and.b16 %rs10, %rs9, 255;
|
|
setp.eq.s16 %p1, %rs10, 0;
|
|
@%p1 bra $L__BB0_2;
|
|
|
|
ld.global.u8 %rs18, [%rd1+2];
|
|
bra.uni $L__BB0_3;
|
|
|
|
$L__BB0_2:
|
|
ld.global.u8 %rs18, [%rd1+2];
|
|
setp.eq.s16 %p2, %rs18, 0;
|
|
mov.f32 %f396, 0f00000000;
|
|
mov.u16 %rs19, 0;
|
|
mov.f32 %f397, %f396;
|
|
mov.f32 %f398, %f396;
|
|
@%p2 bra $L__BB0_4;
|
|
|
|
$L__BB0_3:
|
|
cvt.rn.f32.u16 %f124, %rs7;
|
|
div.rn.f32 %f125, %f124, 0f437F0000;
|
|
fma.rn.f32 %f126, %f125, 0f40000000, 0fBF800000;
|
|
and.b16 %rs13, %rs19, 255;
|
|
cvt.rn.f32.u16 %f127, %rs13;
|
|
div.rn.f32 %f128, %f127, 0f437F0000;
|
|
fma.rn.f32 %f129, %f128, 0f40000000, 0fBF800000;
|
|
cvt.rn.f32.u16 %f130, %rs18;
|
|
div.rn.f32 %f131, %f130, 0f437F0000;
|
|
fma.rn.f32 %f132, %f131, 0f40000000, 0fBF800000;
|
|
mul.f32 %f133, %f129, %f129;
|
|
fma.rn.f32 %f134, %f126, %f126, %f133;
|
|
fma.rn.f32 %f135, %f132, %f132, %f134;
|
|
sqrt.rn.f32 %f136, %f135;
|
|
rcp.rn.f32 %f137, %f136;
|
|
mul.f32 %f398, %f137, %f132;
|
|
mul.f32 %f397, %f137, %f129;
|
|
mul.f32 %f396, %f126, %f137;
|
|
|
|
$L__BB0_4:
|
|
ld.const.v2.u32 {%r15, %r16}, [params];
|
|
add.s32 %r3, %r15, %r7;
|
|
add.s32 %r4, %r16, %r8;
|
|
setp.eq.f32 %p3, %f396, 0f00000000;
|
|
setp.eq.f32 %p4, %f397, 0f00000000;
|
|
and.pred %p5, %p3, %p4;
|
|
setp.eq.f32 %p6, %f398, 0f00000000;
|
|
and.pred %p7, %p6, %p5;
|
|
@%p7 bra $L__BB0_26;
|
|
bra.uni $L__BB0_5;
|
|
|
|
$L__BB0_26:
|
|
ld.const.u64 %rd37, [params+224];
|
|
cvta.to.global.u64 %rd38, %rd37;
|
|
ld.const.u32 %r669, [params+216];
|
|
mad.lo.s32 %r670, %r669, %r4, %r3;
|
|
mul.wide.u32 %rd39, %r670, 16;
|
|
add.s64 %rd40, %rd38, %rd39;
|
|
mov.f32 %f391, 0f00000000;
|
|
st.global.v4.f32 [%rd40], {%f391, %f391, %f391, %f391};
|
|
bra.uni $L__BB0_27;
|
|
|
|
$L__BB0_5:
|
|
ld.const.u64 %rd9, [params+432];
|
|
cvta.to.global.u64 %rd10, %rd9;
|
|
ld.const.u32 %r19, [params+424];
|
|
mad.lo.s32 %r20, %r19, %r8, %r7;
|
|
mul.wide.u32 %rd11, %r20, 16;
|
|
add.s64 %rd12, %rd10, %rd11;
|
|
ld.global.v4.f32 {%f138, %f139, %f140, %f141}, [%rd12];
|
|
abs.f32 %f146, %f138;
|
|
setp.le.f32 %p8, %f146, 0f7F800000;
|
|
selp.f32 %f10, %f138, 0f00000000, %p8;
|
|
abs.f32 %f147, %f139;
|
|
setp.le.f32 %p9, %f147, 0f7F800000;
|
|
selp.f32 %f11, %f139, 0f00000000, %p9;
|
|
abs.f32 %f148, %f140;
|
|
setp.le.f32 %p10, %f148, 0f7F800000;
|
|
selp.f32 %f12, %f140, 0f00000000, %p10;
|
|
add.f32 %f13, %f141, 0f38D1B717;
|
|
mov.f32 %f149, 0f38D1B717;
|
|
mul.f32 %f150, %f10, 0f3456BF95;
|
|
mul.f32 %f151, %f11, 0f3456BF95;
|
|
mul.f32 %f152, %f12, 0f3456BF95;
|
|
abs.f32 %f14, %f396;
|
|
div.rn.f32 %f153, %f150, %f14;
|
|
abs.f32 %f154, %f397;
|
|
div.rn.f32 %f155, %f151, %f154;
|
|
abs.f32 %f15, %f398;
|
|
div.rn.f32 %f156, %f152, %f15;
|
|
abs.f32 %f157, %f153;
|
|
abs.f32 %f158, %f155;
|
|
abs.f32 %f159, %f156;
|
|
max.f32 %f160, %f157, %f149;
|
|
max.f32 %f161, %f158, %f149;
|
|
max.f32 %f162, %f159, %f149;
|
|
mul.f32 %f16, %f396, %f160;
|
|
mul.f32 %f17, %f397, %f161;
|
|
mul.f32 %f18, %f398, %f162;
|
|
add.f32 %f19, %f10, %f16;
|
|
add.f32 %f20, %f11, %f17;
|
|
add.f32 %f21, %f12, %f18;
|
|
ld.const.u64 %rd13, [params+448];
|
|
cvta.to.global.u64 %rd14, %rd13;
|
|
ld.const.u32 %r21, [params+440];
|
|
mad.lo.s32 %r22, %r21, %r8, %r7;
|
|
mul.wide.u32 %rd15, %r22, 16;
|
|
add.s64 %rd16, %rd14, %rd15;
|
|
ld.global.v4.f32 {%f163, %f164, %f165, %f166}, [%rd16];
|
|
cvt.rzi.u32.f32 %r23, %f166;
|
|
ld.const.u64 %rd17, [params+712];
|
|
cvta.to.global.u64 %rd18, %rd17;
|
|
cvt.u64.u32 %rd19, %r23;
|
|
add.s64 %rd2, %rd18, %rd19;
|
|
ld.global.u8 %r671, [%rd2];
|
|
sub.f32 %f25, %f163, %f10;
|
|
sub.f32 %f26, %f164, %f11;
|
|
sub.f32 %f27, %f165, %f12;
|
|
mul.f32 %f168, %f26, %f26;
|
|
fma.rn.f32 %f169, %f25, %f25, %f168;
|
|
fma.rn.f32 %f170, %f27, %f27, %f169;
|
|
sqrt.rn.f32 %f28, %f170;
|
|
setp.ne.s32 %p11, %r671, 255;
|
|
setp.gt.f32 %p12, %f28, 0f3727C5AC;
|
|
and.pred %p13, %p11, %p12;
|
|
@%p13 bra $L__BB0_7;
|
|
bra.uni $L__BB0_6;
|
|
|
|
$L__BB0_7:
|
|
rcp.rn.f32 %f180, %f28;
|
|
mul.f32 %f174, %f25, %f180;
|
|
mul.f32 %f175, %f26, %f180;
|
|
mul.f32 %f176, %f27, %f180;
|
|
ld.const.u64 %rd41, [params+96];
|
|
mov.f32 %f179, 0f00000000;
|
|
mov.u32 %r60, 2;
|
|
mov.u32 %r62, 1;
|
|
mov.u32 %r63, 1065353216;
|
|
mov.u32 %r94, 0;
|
|
// begin inline asm
|
|
call(%r24,%r25,%r26,%r27,%r28,%r29,%r30,%r31,%r32,%r33,%r34,%r35,%r36,%r37,%r38,%r39,%r40,%r41,%r42,%r43,%r44,%r45,%r46,%r47,%r48,%r49,%r50,%r51,%r52,%r53,%r54,%r55),_optix_trace_typed_32,(%r94,%rd41,%f19,%f20,%f21,%f174,%f175,%f176,%f149,%f28,%f179,%r62,%r94,%r62,%r60,%r62,%r62,%r63,%r94,%r94,%r94,%r94,%r94,%r94,%r94,%r94,%r94,%r94,%r94,%r94,%r94,%r94,%r94,%r94,%r94,%r94,%r94,%r94,%r94,%r94,%r94,%r94,%r94,%r94,%r94,%r94,%r94,%r94,%r94);
|
|
// end inline asm
|
|
mov.b32 %f181, %r24;
|
|
setp.neu.f32 %p14, %f181, 0f00000000;
|
|
@%p14 bra $L__BB0_9;
|
|
|
|
mov.u16 %rs15, 255;
|
|
st.global.u8 [%rd2], %rs15;
|
|
mov.u32 %r671, 255;
|
|
bra.uni $L__BB0_9;
|
|
|
|
$L__BB0_6:
|
|
ld.const.u64 %rd41, [params+96];
|
|
|
|
$L__BB0_9:
|
|
setp.eq.s32 %p15, %r671, 0;
|
|
mov.u32 %r166, 0;
|
|
selp.f32 %f404, %f163, %f10, %p15;
|
|
add.f32 %f182, %f16, %f404;
|
|
selp.f32 %f405, %f164, %f11, %p15;
|
|
add.f32 %f183, %f17, %f405;
|
|
selp.f32 %f406, %f165, %f12, %p15;
|
|
add.f32 %f184, %f18, %f406;
|
|
neg.f32 %f192, %f397;
|
|
setp.gt.f32 %p16, %f14, %f15;
|
|
selp.f32 %f193, %f192, 0f00000000, %p16;
|
|
mov.f32 %f190, 0f00000000;
|
|
neg.f32 %f194, %f398;
|
|
selp.f32 %f195, %f396, %f194, %p16;
|
|
selp.f32 %f196, 0f00000000, %f397, %p16;
|
|
mul.f32 %f197, %f195, %f195;
|
|
fma.rn.f32 %f198, %f193, %f193, %f197;
|
|
fma.rn.f32 %f199, %f196, %f196, %f198;
|
|
sqrt.rn.f32 %f200, %f199;
|
|
rcp.rn.f32 %f201, %f200;
|
|
mul.f32 %f32, %f193, %f201;
|
|
mul.f32 %f33, %f195, %f201;
|
|
mul.f32 %f34, %f196, %f201;
|
|
mul.f32 %f202, %f398, %f33;
|
|
mul.f32 %f203, %f397, %f34;
|
|
sub.f32 %f35, %f202, %f203;
|
|
mul.f32 %f204, %f396, %f34;
|
|
mul.f32 %f205, %f398, %f32;
|
|
sub.f32 %f36, %f204, %f205;
|
|
mul.f32 %f206, %f397, %f32;
|
|
mul.f32 %f207, %f396, %f33;
|
|
sub.f32 %f37, %f206, %f207;
|
|
mov.f32 %f208, 0f3F8147AE;
|
|
sqrt.rn.f32 %f209, %f208;
|
|
rcp.rn.f32 %f38, %f209;
|
|
mul.f32 %f40, %f38, 0f00000000;
|
|
mul.f32 %f210, %f38, 0f3DCCCCCD;
|
|
mul.f32 %f211, %f182, 0f3456BF95;
|
|
mul.f32 %f212, %f183, 0f3456BF95;
|
|
mul.f32 %f213, %f184, 0f3456BF95;
|
|
abs.f32 %f214, %f211;
|
|
abs.f32 %f215, %f212;
|
|
abs.f32 %f216, %f213;
|
|
max.f32 %f217, %f214, %f215;
|
|
max.f32 %f218, %f217, %f216;
|
|
mov.f32 %f219, 0f38D1B717;
|
|
max.f32 %f41, %f218, %f219;
|
|
mul.f32 %f220, %f35, %f38;
|
|
mul.f32 %f221, %f36, %f38;
|
|
mul.f32 %f222, %f37, %f38;
|
|
mul.f32 %f45, %f32, %f40;
|
|
mul.f32 %f46, %f33, %f40;
|
|
mul.f32 %f47, %f34, %f40;
|
|
sub.f32 %f223, %f45, %f220;
|
|
sub.f32 %f224, %f46, %f221;
|
|
sub.f32 %f225, %f47, %f222;
|
|
mul.f32 %f48, %f396, %f210;
|
|
mul.f32 %f49, %f397, %f210;
|
|
mul.f32 %f50, %f398, %f210;
|
|
add.f32 %f185, %f223, %f48;
|
|
add.f32 %f186, %f224, %f49;
|
|
add.f32 %f187, %f225, %f50;
|
|
mul.f32 %f189, %f13, 0f3FB504F3;
|
|
mov.u32 %r129, 1;
|
|
mov.u32 %r132, 2;
|
|
mov.u32 %r134, 3;
|
|
// begin inline asm
|
|
call(%r96,%r97,%r98,%r99,%r100,%r101,%r102,%r103,%r104,%r105,%r106,%r107,%r108,%r109,%r110,%r111,%r112,%r113,%r114,%r115,%r116,%r117,%r118,%r119,%r120,%r121,%r122,%r123,%r124,%r125,%r126,%r127),_optix_trace_typed_32,(%r166,%rd41,%f182,%f183,%f184,%f185,%f186,%f187,%f190,%f189,%f190,%r129,%r166,%r166,%r132,%r166,%r134,%r166,%r166,%r166,%r166,%r166,%r166,%r166,%r166,%r166,%r166,%r166,%r166,%r166,%r166,%r166,%r166,%r166,%r166,%r166,%r166,%r166,%r166,%r166,%r166,%r166,%r166,%r166,%r166,%r166,%r166,%r166,%r166);
|
|
// end inline asm
|
|
mov.b32 %f55, %r96;
|
|
mov.b32 %f56, %r97;
|
|
mov.b32 %f57, %r98;
|
|
add.f32 %f226, %f55, %f56;
|
|
add.f32 %f227, %f226, %f57;
|
|
setp.eq.f32 %p17, %f227, 0f00000000;
|
|
mov.f32 %f403, 0f47C34F80;
|
|
@%p17 bra $L__BB0_11;
|
|
|
|
mul.f32 %f228, %f56, %f56;
|
|
fma.rn.f32 %f229, %f55, %f55, %f228;
|
|
fma.rn.f32 %f230, %f57, %f57, %f229;
|
|
sqrt.rn.f32 %f231, %f230;
|
|
rcp.rn.f32 %f232, %f231;
|
|
mul.f32 %f233, %f232, %f55;
|
|
mul.f32 %f234, %f232, %f56;
|
|
mul.f32 %f235, %f232, %f57;
|
|
fma.rn.f32 %f236, %f41, %f233, %f182;
|
|
fma.rn.f32 %f237, %f41, %f234, %f183;
|
|
fma.rn.f32 %f238, %f41, %f235, %f184;
|
|
setp.lt.f32 %p18, %f231, 0f47C34F80;
|
|
fma.rn.f32 %f239, %f231, %f185, %f236;
|
|
fma.rn.f32 %f240, %f231, %f186, %f237;
|
|
fma.rn.f32 %f241, %f231, %f187, %f238;
|
|
min.f32 %f403, %f231, 0f47C34F80;
|
|
selp.f32 %f406, %f241, %f406, %p18;
|
|
selp.f32 %f405, %f240, %f405, %p18;
|
|
selp.f32 %f404, %f239, %f404, %p18;
|
|
|
|
$L__BB0_11:
|
|
fma.rn.f32 %f251, %f35, %f38, %f45;
|
|
fma.rn.f32 %f252, %f36, %f38, %f46;
|
|
fma.rn.f32 %f253, %f37, %f38, %f47;
|
|
add.f32 %f245, %f251, %f48;
|
|
add.f32 %f246, %f252, %f49;
|
|
add.f32 %f247, %f253, %f50;
|
|
// begin inline asm
|
|
call(%r167,%r168,%r169,%r170,%r171,%r172,%r173,%r174,%r175,%r176,%r177,%r178,%r179,%r180,%r181,%r182,%r183,%r184,%r185,%r186,%r187,%r188,%r189,%r190,%r191,%r192,%r193,%r194,%r195,%r196,%r197,%r198),_optix_trace_typed_32,(%r166,%rd41,%f182,%f183,%f184,%f245,%f246,%f247,%f190,%f189,%f190,%r129,%r166,%r166,%r132,%r166,%r134,%r166,%r166,%r166,%r166,%r166,%r166,%r166,%r166,%r166,%r166,%r166,%r166,%r166,%r166,%r166,%r166,%r166,%r166,%r166,%r166,%r166,%r166,%r166,%r166,%r166,%r166,%r166,%r166,%r166,%r166,%r166,%r166);
|
|
// end inline asm
|
|
mov.b32 %f69, %r167;
|
|
mov.b32 %f70, %r168;
|
|
mov.b32 %f71, %r169;
|
|
add.f32 %f254, %f69, %f70;
|
|
add.f32 %f255, %f254, %f71;
|
|
setp.eq.f32 %p19, %f255, 0f00000000;
|
|
@%p19 bra $L__BB0_13;
|
|
|
|
mul.f32 %f256, %f70, %f70;
|
|
fma.rn.f32 %f257, %f69, %f69, %f256;
|
|
fma.rn.f32 %f258, %f71, %f71, %f257;
|
|
sqrt.rn.f32 %f259, %f258;
|
|
rcp.rn.f32 %f260, %f259;
|
|
mul.f32 %f261, %f260, %f69;
|
|
mul.f32 %f262, %f260, %f70;
|
|
mul.f32 %f263, %f260, %f71;
|
|
fma.rn.f32 %f264, %f41, %f261, %f182;
|
|
fma.rn.f32 %f265, %f41, %f262, %f183;
|
|
fma.rn.f32 %f266, %f41, %f263, %f184;
|
|
setp.lt.f32 %p20, %f259, %f403;
|
|
fma.rn.f32 %f267, %f259, %f245, %f264;
|
|
fma.rn.f32 %f268, %f259, %f246, %f265;
|
|
fma.rn.f32 %f269, %f259, %f247, %f266;
|
|
selp.f32 %f403, %f259, %f403, %p20;
|
|
selp.f32 %f406, %f269, %f406, %p20;
|
|
selp.f32 %f405, %f268, %f405, %p20;
|
|
selp.f32 %f404, %f267, %f404, %p20;
|
|
|
|
$L__BB0_13:
|
|
neg.f32 %f394, %f38;
|
|
mul.f32 %f393, %f38, 0f00000000;
|
|
mul.f32 %f80, %f35, %f393;
|
|
fma.rn.f32 %f279, %f32, %f394, %f80;
|
|
mul.f32 %f81, %f36, %f393;
|
|
fma.rn.f32 %f280, %f33, %f394, %f81;
|
|
mul.f32 %f82, %f37, %f393;
|
|
fma.rn.f32 %f281, %f34, %f394, %f82;
|
|
add.f32 %f273, %f279, %f48;
|
|
add.f32 %f274, %f280, %f49;
|
|
add.f32 %f275, %f281, %f50;
|
|
mov.f32 %f278, 0f00000000;
|
|
mov.u32 %r271, 1;
|
|
mov.u32 %r274, 2;
|
|
mov.u32 %r276, 3;
|
|
mov.u32 %r308, 0;
|
|
// begin inline asm
|
|
call(%r238,%r239,%r240,%r241,%r242,%r243,%r244,%r245,%r246,%r247,%r248,%r249,%r250,%r251,%r252,%r253,%r254,%r255,%r256,%r257,%r258,%r259,%r260,%r261,%r262,%r263,%r264,%r265,%r266,%r267,%r268,%r269),_optix_trace_typed_32,(%r308,%rd41,%f182,%f183,%f184,%f273,%f274,%f275,%f278,%f189,%f278,%r271,%r308,%r308,%r274,%r308,%r276,%r308,%r308,%r308,%r308,%r308,%r308,%r308,%r308,%r308,%r308,%r308,%r308,%r308,%r308,%r308,%r308,%r308,%r308,%r308,%r308,%r308,%r308,%r308,%r308,%r308,%r308,%r308,%r308,%r308,%r308,%r308,%r308);
|
|
// end inline asm
|
|
mov.b32 %f86, %r238;
|
|
mov.b32 %f87, %r239;
|
|
mov.b32 %f88, %r240;
|
|
add.f32 %f282, %f86, %f87;
|
|
add.f32 %f283, %f282, %f88;
|
|
setp.eq.f32 %p21, %f283, 0f00000000;
|
|
@%p21 bra $L__BB0_15;
|
|
|
|
mul.f32 %f284, %f87, %f87;
|
|
fma.rn.f32 %f285, %f86, %f86, %f284;
|
|
fma.rn.f32 %f286, %f88, %f88, %f285;
|
|
sqrt.rn.f32 %f287, %f286;
|
|
rcp.rn.f32 %f288, %f287;
|
|
mul.f32 %f289, %f288, %f86;
|
|
mul.f32 %f290, %f288, %f87;
|
|
mul.f32 %f291, %f288, %f88;
|
|
fma.rn.f32 %f292, %f41, %f289, %f182;
|
|
fma.rn.f32 %f293, %f41, %f290, %f183;
|
|
fma.rn.f32 %f294, %f41, %f291, %f184;
|
|
setp.lt.f32 %p22, %f287, %f403;
|
|
fma.rn.f32 %f295, %f287, %f273, %f292;
|
|
fma.rn.f32 %f296, %f287, %f274, %f293;
|
|
fma.rn.f32 %f297, %f287, %f275, %f294;
|
|
selp.f32 %f403, %f287, %f403, %p22;
|
|
selp.f32 %f406, %f297, %f406, %p22;
|
|
selp.f32 %f405, %f296, %f405, %p22;
|
|
selp.f32 %f404, %f295, %f404, %p22;
|
|
|
|
$L__BB0_15:
|
|
fma.rn.f32 %f307, %f32, %f38, %f80;
|
|
fma.rn.f32 %f308, %f33, %f38, %f81;
|
|
fma.rn.f32 %f309, %f34, %f38, %f82;
|
|
add.f32 %f301, %f307, %f48;
|
|
add.f32 %f302, %f308, %f49;
|
|
add.f32 %f303, %f309, %f50;
|
|
// begin inline asm
|
|
call(%r309,%r310,%r311,%r312,%r313,%r314,%r315,%r316,%r317,%r318,%r319,%r320,%r321,%r322,%r323,%r324,%r325,%r326,%r327,%r328,%r329,%r330,%r331,%r332,%r333,%r334,%r335,%r336,%r337,%r338,%r339,%r340),_optix_trace_typed_32,(%r308,%rd41,%f182,%f183,%f184,%f301,%f302,%f303,%f278,%f189,%f278,%r271,%r308,%r308,%r274,%r308,%r276,%r308,%r308,%r308,%r308,%r308,%r308,%r308,%r308,%r308,%r308,%r308,%r308,%r308,%r308,%r308,%r308,%r308,%r308,%r308,%r308,%r308,%r308,%r308,%r308,%r308,%r308,%r308,%r308,%r308,%r308,%r308,%r308);
|
|
// end inline asm
|
|
mov.b32 %f100, %r309;
|
|
mov.b32 %f101, %r310;
|
|
mov.b32 %f102, %r311;
|
|
add.f32 %f310, %f100, %f101;
|
|
add.f32 %f311, %f310, %f102;
|
|
setp.eq.f32 %p23, %f311, 0f00000000;
|
|
@%p23 bra $L__BB0_17;
|
|
|
|
mul.f32 %f312, %f101, %f101;
|
|
fma.rn.f32 %f313, %f100, %f100, %f312;
|
|
fma.rn.f32 %f314, %f102, %f102, %f313;
|
|
sqrt.rn.f32 %f315, %f314;
|
|
rcp.rn.f32 %f316, %f315;
|
|
mul.f32 %f317, %f316, %f100;
|
|
mul.f32 %f318, %f316, %f101;
|
|
mul.f32 %f319, %f316, %f102;
|
|
fma.rn.f32 %f320, %f41, %f317, %f182;
|
|
fma.rn.f32 %f321, %f41, %f318, %f183;
|
|
fma.rn.f32 %f322, %f41, %f319, %f184;
|
|
setp.lt.f32 %p24, %f315, %f403;
|
|
fma.rn.f32 %f323, %f315, %f301, %f320;
|
|
fma.rn.f32 %f324, %f315, %f302, %f321;
|
|
fma.rn.f32 %f325, %f315, %f303, %f322;
|
|
selp.f32 %f406, %f325, %f406, %p24;
|
|
selp.f32 %f405, %f324, %f405, %p24;
|
|
selp.f32 %f404, %f323, %f404, %p24;
|
|
|
|
$L__BB0_17:
|
|
abs.f32 %f326, %f404;
|
|
setp.le.f32 %p25, %f326, 0f7F800000;
|
|
selp.f32 %f414, %f404, 0f00000000, %p25;
|
|
abs.f32 %f327, %f405;
|
|
setp.le.f32 %p26, %f327, 0f7F800000;
|
|
selp.f32 %f415, %f405, 0f00000000, %p26;
|
|
abs.f32 %f328, %f406;
|
|
setp.le.f32 %p27, %f328, 0f7F800000;
|
|
selp.f32 %f416, %f406, 0f00000000, %p27;
|
|
ld.const.f32 %f112, [params+720];
|
|
setp.eq.f32 %p28, %f112, 0f00000000;
|
|
@%p28 bra $L__BB0_19;
|
|
|
|
fma.rn.f32 %f414, %f396, %f112, %f414;
|
|
fma.rn.f32 %f415, %f397, %f112, %f415;
|
|
fma.rn.f32 %f416, %f398, %f112, %f416;
|
|
|
|
$L__BB0_19:
|
|
ld.const.u64 %rd25, [params+224];
|
|
cvta.to.global.u64 %rd26, %rd25;
|
|
ld.const.u32 %r380, [params+216];
|
|
mad.lo.s32 %r381, %r380, %r4, %r3;
|
|
mul.wide.u32 %rd27, %r381, 16;
|
|
add.s64 %rd28, %rd26, %rd27;
|
|
mov.f32 %f329, 0f3F800000;
|
|
st.global.v4.f32 [%rd28], {%f414, %f415, %f416, %f329};
|
|
ld.const.u8 %rs16, [params+104];
|
|
and.b16 %rs17, %rs16, 32;
|
|
setp.eq.s16 %p29, %rs17, 0;
|
|
mov.f32 %f417, 0f437F0000;
|
|
@%p29 bra $L__BB0_27;
|
|
|
|
add.f32 %f395, %f141, 0f38D1B717;
|
|
mul.f32 %f337, %f395, 0f41000000;
|
|
mov.f32 %f336, 0f38D1B717;
|
|
mov.f32 %f338, 0f00000000;
|
|
mov.u32 %r415, 1;
|
|
mov.u32 %r418, 2;
|
|
mov.u32 %r420, 3;
|
|
mov.u32 %r452, 0;
|
|
// begin inline asm
|
|
call(%r382,%r383,%r384,%r385,%r386,%r387,%r388,%r389,%r390,%r391,%r392,%r393,%r394,%r395,%r396,%r397,%r398,%r399,%r400,%r401,%r402,%r403,%r404,%r405,%r406,%r407,%r408,%r409,%r410,%r411,%r412,%r413),_optix_trace_typed_32,(%r452,%rd41,%f182,%f183,%f184,%f185,%f186,%f187,%f336,%f337,%f338,%r415,%r452,%r452,%r418,%r452,%r420,%r452,%r452,%r452,%r452,%r452,%r452,%r452,%r452,%r452,%r452,%r452,%r452,%r452,%r452,%r452,%r452,%r452,%r452,%r452,%r452,%r452,%r452,%r452,%r452,%r452,%r452,%r452,%r452,%r452,%r452,%r452,%r452);
|
|
// end inline asm
|
|
mov.b32 %f340, %r382;
|
|
mov.b32 %f341, %r383;
|
|
mov.b32 %f342, %r384;
|
|
add.f32 %f343, %f340, %f341;
|
|
add.f32 %f344, %f343, %f342;
|
|
setp.neu.f32 %p30, %f344, 0f00000000;
|
|
@%p30 bra $L__BB0_25;
|
|
|
|
// begin inline asm
|
|
call(%r453,%r454,%r455,%r456,%r457,%r458,%r459,%r460,%r461,%r462,%r463,%r464,%r465,%r466,%r467,%r468,%r469,%r470,%r471,%r472,%r473,%r474,%r475,%r476,%r477,%r478,%r479,%r480,%r481,%r482,%r483,%r484),_optix_trace_typed_32,(%r452,%rd41,%f182,%f183,%f184,%f245,%f246,%f247,%f336,%f337,%f338,%r415,%r452,%r452,%r418,%r452,%r420,%r452,%r452,%r452,%r452,%r452,%r452,%r452,%r452,%r452,%r452,%r452,%r452,%r452,%r452,%r452,%r452,%r452,%r452,%r452,%r452,%r452,%r452,%r452,%r452,%r452,%r452,%r452,%r452,%r452,%r452,%r452,%r452);
|
|
// end inline asm
|
|
mov.b32 %f355, %r453;
|
|
mov.b32 %f356, %r454;
|
|
mov.b32 %f357, %r455;
|
|
add.f32 %f358, %f355, %f356;
|
|
add.f32 %f359, %f358, %f357;
|
|
setp.neu.f32 %p31, %f359, 0f00000000;
|
|
@%p31 bra $L__BB0_25;
|
|
|
|
mov.f32 %f366, 0f38D1B717;
|
|
mov.f32 %f368, 0f00000000;
|
|
mov.u32 %r557, 1;
|
|
mov.u32 %r560, 2;
|
|
mov.u32 %r562, 3;
|
|
mov.u32 %r594, 0;
|
|
// begin inline asm
|
|
call(%r524,%r525,%r526,%r527,%r528,%r529,%r530,%r531,%r532,%r533,%r534,%r535,%r536,%r537,%r538,%r539,%r540,%r541,%r542,%r543,%r544,%r545,%r546,%r547,%r548,%r549,%r550,%r551,%r552,%r553,%r554,%r555),_optix_trace_typed_32,(%r594,%rd41,%f182,%f183,%f184,%f273,%f274,%f275,%f366,%f337,%f368,%r557,%r594,%r594,%r560,%r594,%r562,%r594,%r594,%r594,%r594,%r594,%r594,%r594,%r594,%r594,%r594,%r594,%r594,%r594,%r594,%r594,%r594,%r594,%r594,%r594,%r594,%r594,%r594,%r594,%r594,%r594,%r594,%r594,%r594,%r594,%r594,%r594,%r594);
|
|
// end inline asm
|
|
mov.b32 %f370, %r524;
|
|
mov.b32 %f371, %r525;
|
|
mov.b32 %f372, %r526;
|
|
add.f32 %f373, %f370, %f371;
|
|
add.f32 %f374, %f373, %f372;
|
|
setp.neu.f32 %p32, %f374, 0f00000000;
|
|
@%p32 bra $L__BB0_25;
|
|
|
|
// begin inline asm
|
|
call(%r595,%r596,%r597,%r598,%r599,%r600,%r601,%r602,%r603,%r604,%r605,%r606,%r607,%r608,%r609,%r610,%r611,%r612,%r613,%r614,%r615,%r616,%r617,%r618,%r619,%r620,%r621,%r622,%r623,%r624,%r625,%r626),_optix_trace_typed_32,(%r594,%rd41,%f182,%f183,%f184,%f301,%f302,%f303,%f366,%f337,%f368,%r557,%r594,%r594,%r560,%r594,%r562,%r594,%r594,%r594,%r594,%r594,%r594,%r594,%r594,%r594,%r594,%r594,%r594,%r594,%r594,%r594,%r594,%r594,%r594,%r594,%r594,%r594,%r594,%r594,%r594,%r594,%r594,%r594,%r594,%r594,%r594,%r594,%r594);
|
|
// end inline asm
|
|
mov.b32 %f385, %r595;
|
|
mov.b32 %f386, %r596;
|
|
mov.b32 %f387, %r597;
|
|
add.f32 %f388, %f385, %f386;
|
|
add.f32 %f389, %f388, %f387;
|
|
setp.neu.f32 %p33, %f389, 0f00000000;
|
|
@%p33 bra $L__BB0_25;
|
|
|
|
mov.f32 %f417, 0f00000000;
|
|
|
|
$L__BB0_25:
|
|
ld.const.u64 %rd33, [params+192];
|
|
cvta.to.global.u64 %rd34, %rd33;
|
|
ld.const.u32 %r666, [params+184];
|
|
mad.lo.s32 %r667, %r666, %r4, %r3;
|
|
cvt.u64.u32 %rd35, %r667;
|
|
cvt.rzi.u32.f32 %r668, %f417;
|
|
add.s64 %rd36, %rd34, %rd35;
|
|
st.global.u8 [%rd36], %r668;
|
|
|
|
$L__BB0_27:
|
|
ret;
|
|
|
|
}
|
|
|