Advertisement
Not a member of Pastebin yet?
Sign Up,
it unlocks many cool features!
- //
- // Generated by NVIDIA NVVM Compiler
- //
- // Compiler Build ID: CL-26218862
- // Cuda compilation tools, release 10.1, V10.1.168
- // Based on LLVM 3.4svn
- //
- .version 6.4
- .target sm_30
- .address_size 64
- .extern .func (.param .b32 func_retval0) _Z5powiffi
- (
- .param .b32 _Z5powiffi_param_0,
- .param .b32 _Z5powiffi_param_1
- )
- ;
- .func (.param .align 4 .b8 func_retval0[12]) _Z5traceRK4Vec3IfES2_PK6SpherejRKi(
- .param .b32 _Z5traceRK4Vec3IfES2_PK6SpherejRKi_param_0,
- .param .b32 _Z5traceRK4Vec3IfES2_PK6SpherejRKi_param_1,
- .param .b32 _Z5traceRK4Vec3IfES2_PK6SpherejRKi_param_2,
- .param .b64 _Z5traceRK4Vec3IfES2_PK6SpherejRKi_param_3,
- .param .b64 _Z5traceRK4Vec3IfES2_PK6SpherejRKi_param_4,
- .param .b32 _Z5traceRK4Vec3IfES2_PK6SpherejRKi_param_5,
- .param .b64 _Z5traceRK4Vec3IfES2_PK6SpherejRKi_param_6
- )
- {
- .local .align 4 .b8 __local_depot0[32];
- .reg .b64 %SP;
- .reg .b64 %SPL;
- .reg .pred %p<88>;
- .reg .b16 %rs<5>;
- .reg .f32 %f<480>;
- .reg .b32 %r<40>;
- .reg .b64 %rd<93>;
- mov.u64 %SPL, __local_depot0;
- cvta.local.u64 %SP, %SPL;
- ld.param.f32 %f179, [_Z5traceRK4Vec3IfES2_PK6SpherejRKi_param_0];
- ld.param.f32 %f180, [_Z5traceRK4Vec3IfES2_PK6SpherejRKi_param_1];
- ld.param.f32 %f181, [_Z5traceRK4Vec3IfES2_PK6SpherejRKi_param_2];
- ld.param.u64 %rd34, [_Z5traceRK4Vec3IfES2_PK6SpherejRKi_param_3];
- ld.param.u64 %rd35, [_Z5traceRK4Vec3IfES2_PK6SpherejRKi_param_4];
- ld.param.u32 %r13, [_Z5traceRK4Vec3IfES2_PK6SpherejRKi_param_5];
- ld.param.u64 %rd36, [_Z5traceRK4Vec3IfES2_PK6SpherejRKi_param_6];
- setp.eq.s32 %p8, %r13, 0;
- mov.f32 %f433, 0f4CBEBC20;
- mov.u64 %rd83, 0;
- @%p8 bra BB0_61;
- cvta.to.local.u64 %rd39, %rd34;
- ld.local.f32 %f1, [%rd39];
- ld.local.f32 %f2, [%rd39+4];
- ld.local.f32 %f3, [%rd39+8];
- and.b32 %r1, %r13, 3;
- setp.eq.s32 %p9, %r1, 0;
- mov.f32 %f439, 0f00000000;
- mov.f32 %f433, 0f4CBEBC20;
- mov.u64 %rd83, 0;
- mov.u32 %r36, 0;
- @%p9 bra BB0_30;
- setp.eq.s32 %p10, %r1, 1;
- mov.f32 %f435, 0f4CBEBC20;
- mov.u64 %rd83, 0;
- mov.u32 %r35, 0;
- mov.f32 %f433, %f435;
- @%p10 bra BB0_22;
- setp.eq.s32 %p11, %r1, 2;
- mov.f32 %f431, 0f4CBEBC20;
- mov.u64 %rd83, 0;
- mov.u32 %r34, 0;
- @%p11 bra BB0_4;
- bra.uni BB0_5;
- BB0_4:
- mov.f32 %f433, %f431;
- bra.uni BB0_14;
- BB0_5:
- ld.f32 %f189, [%rd35+8];
- sub.f32 %f4, %f189, %f179;
- ld.f32 %f190, [%rd35+12];
- sub.f32 %f5, %f190, %f180;
- ld.f32 %f191, [%rd35+16];
- sub.f32 %f6, %f191, %f181;
- mul.f32 %f192, %f5, %f2;
- fma.rn.f32 %f193, %f4, %f1, %f192;
- fma.rn.f32 %f7, %f6, %f3, %f193;
- setp.lt.f32 %p13, %f7, 0f00000000;
- mov.pred %p81, -1;
- mov.f32 %f188, 0f4CBEBC20;
- @%p13 bra BB0_6;
- bra.uni BB0_7;
- BB0_6:
- mov.f32 %f428, %f188;
- mov.f32 %f429, %f188;
- bra.uni BB0_9;
- BB0_7:
- mul.f32 %f196, %f5, %f5;
- fma.rn.f32 %f197, %f4, %f4, %f196;
- fma.rn.f32 %f198, %f6, %f6, %f197;
- mul.f32 %f199, %f7, %f7;
- sub.f32 %f8, %f198, %f199;
- ld.f32 %f9, [%rd35+24];
- setp.gt.f32 %p15, %f8, %f9;
- mov.f32 %f428, %f188;
- mov.f32 %f429, %f188;
- @%p15 bra BB0_9;
- sub.f32 %f200, %f9, %f8;
- sqrt.rn.f32 %f201, %f200;
- sub.f32 %f429, %f7, %f201;
- add.f32 %f428, %f7, %f201;
- mov.pred %p81, 0;
- BB0_9:
- mov.u32 %r34, 1;
- @%p81 bra BB0_10;
- bra.uni BB0_11;
- BB0_10:
- mov.f32 %f433, %f188;
- bra.uni BB0_14;
- BB0_11:
- setp.lt.f32 %p17, %f429, 0f00000000;
- selp.f32 %f433, %f428, %f429, %p17;
- setp.geu.f32 %p18, %f433, 0f4CBEBC20;
- @%p18 bra BB0_12;
- mov.u64 %rd83, %rd35;
- bra.uni BB0_14;
- BB0_12:
- mov.f32 %f433, %f188;
- BB0_14:
- mul.wide.u32 %rd44, %r34, 80;
- add.s64 %rd45, %rd35, %rd44;
- ld.f32 %f206, [%rd45+8];
- sub.f32 %f16, %f206, %f179;
- ld.f32 %f207, [%rd45+12];
- sub.f32 %f17, %f207, %f180;
- ld.f32 %f208, [%rd45+16];
- sub.f32 %f18, %f208, %f181;
- mul.f32 %f209, %f17, %f2;
- fma.rn.f32 %f210, %f16, %f1, %f209;
- fma.rn.f32 %f19, %f18, %f3, %f210;
- setp.lt.f32 %p20, %f19, 0f00000000;
- mov.pred %p82, -1;
- @%p20 bra BB0_15;
- bra.uni BB0_16;
- BB0_15:
- mov.f32 %f432, %f431;
- bra.uni BB0_18;
- BB0_16:
- mul.f32 %f213, %f17, %f17;
- fma.rn.f32 %f214, %f16, %f16, %f213;
- fma.rn.f32 %f215, %f18, %f18, %f214;
- mul.f32 %f216, %f19, %f19;
- sub.f32 %f20, %f215, %f216;
- ld.f32 %f21, [%rd45+24];
- setp.gt.f32 %p22, %f20, %f21;
- mov.f32 %f432, %f431;
- @%p22 bra BB0_18;
- sub.f32 %f217, %f21, %f20;
- sqrt.rn.f32 %f218, %f217;
- sub.f32 %f432, %f19, %f218;
- add.f32 %f431, %f19, %f218;
- mov.pred %p82, 0;
- BB0_18:
- @%p82 bra BB0_21;
- setp.lt.f32 %p24, %f432, 0f00000000;
- selp.f32 %f26, %f431, %f432, %p24;
- setp.geu.f32 %p25, %f26, %f433;
- @%p25 bra BB0_21;
- mov.f32 %f433, %f26;
- mov.u64 %rd83, %rd45;
- BB0_21:
- add.s32 %r35, %r34, 1;
- BB0_22:
- cvt.u64.u32 %rd5, %r35;
- mul.wide.u32 %rd49, %r35, 80;
- add.s64 %rd50, %rd35, %rd49;
- add.s64 %rd6, %rd50, 8;
- ld.f32 %f221, [%rd50+8];
- sub.f32 %f29, %f221, %f179;
- ld.f32 %f222, [%rd50+12];
- sub.f32 %f30, %f222, %f180;
- ld.f32 %f223, [%rd50+16];
- sub.f32 %f31, %f223, %f181;
- mul.f32 %f224, %f30, %f2;
- fma.rn.f32 %f225, %f29, %f1, %f224;
- fma.rn.f32 %f32, %f31, %f3, %f225;
- setp.lt.f32 %p27, %f32, 0f00000000;
- mov.pred %p83, -1;
- @%p27 bra BB0_23;
- bra.uni BB0_24;
- BB0_23:
- mov.f32 %f436, %f435;
- bra.uni BB0_26;
- BB0_24:
- mul.f32 %f228, %f30, %f30;
- fma.rn.f32 %f229, %f29, %f29, %f228;
- fma.rn.f32 %f230, %f31, %f31, %f229;
- mul.f32 %f231, %f32, %f32;
- sub.f32 %f33, %f230, %f231;
- ld.f32 %f34, [%rd6+16];
- setp.gt.f32 %p29, %f33, %f34;
- mov.f32 %f436, %f435;
- @%p29 bra BB0_26;
- sub.f32 %f232, %f34, %f33;
- sqrt.rn.f32 %f233, %f232;
- sub.f32 %f436, %f32, %f233;
- add.f32 %f435, %f32, %f233;
- mov.pred %p83, 0;
- BB0_26:
- @%p83 bra BB0_29;
- setp.lt.f32 %p31, %f436, 0f00000000;
- selp.f32 %f39, %f435, %f436, %p31;
- setp.geu.f32 %p32, %f39, %f433;
- @%p32 bra BB0_29;
- mul.lo.s64 %rd51, %rd5, 80;
- add.s64 %rd83, %rd35, %rd51;
- mov.f32 %f433, %f39;
- BB0_29:
- cvt.u32.u64 %r20, %rd5;
- add.s32 %r36, %r20, 1;
- mov.f32 %f439, %f433;
- BB0_30:
- setp.lt.u32 %p33, %r13, 4;
- @%p33 bra BB0_31;
- bra.uni BB0_32;
- BB0_31:
- mov.f32 %f433, %f439;
- bra.uni BB0_61;
- BB0_32:
- cvt.u64.u32 %rd11, %r36;
- mul.wide.u32 %rd52, %r36, 80;
- add.s64 %rd53, %rd35, %rd52;
- add.s64 %rd12, %rd53, 8;
- ld.f32 %f236, [%rd53+8];
- sub.f32 %f44, %f236, %f179;
- ld.f32 %f237, [%rd53+12];
- sub.f32 %f45, %f237, %f180;
- ld.f32 %f238, [%rd53+16];
- sub.f32 %f46, %f238, %f181;
- mul.f32 %f239, %f45, %f2;
- fma.rn.f32 %f240, %f44, %f1, %f239;
- fma.rn.f32 %f47, %f46, %f3, %f240;
- setp.lt.f32 %p35, %f47, 0f00000000;
- mov.pred %p84, -1;
- mov.f32 %f441, 0f4CBEBC20;
- @%p35 bra BB0_33;
- bra.uni BB0_34;
- BB0_33:
- mov.f32 %f442, %f441;
- bra.uni BB0_36;
- BB0_34:
- mul.f32 %f243, %f45, %f45;
- fma.rn.f32 %f244, %f44, %f44, %f243;
- fma.rn.f32 %f245, %f46, %f46, %f244;
- mul.f32 %f246, %f47, %f47;
- sub.f32 %f48, %f245, %f246;
- ld.f32 %f49, [%rd12+16];
- setp.gt.f32 %p37, %f48, %f49;
- mov.f32 %f442, %f441;
- @%p37 bra BB0_36;
- sub.f32 %f247, %f49, %f48;
- sqrt.rn.f32 %f248, %f247;
- sub.f32 %f442, %f47, %f248;
- add.f32 %f441, %f47, %f248;
- mov.pred %p84, 0;
- BB0_36:
- @%p84 bra BB0_39;
- setp.lt.f32 %p39, %f442, 0f00000000;
- selp.f32 %f54, %f441, %f442, %p39;
- setp.geu.f32 %p40, %f54, %f433;
- @%p40 bra BB0_39;
- mul.lo.s64 %rd54, %rd11, 80;
- add.s64 %rd83, %rd35, %rd54;
- mov.f32 %f433, %f54;
- BB0_39:
- add.s32 %r21, %r36, 1;
- cvt.u64.u32 %rd15, %r21;
- mul.wide.u32 %rd55, %r21, 80;
- add.s64 %rd56, %rd35, %rd55;
- add.s64 %rd16, %rd56, 8;
- ld.f32 %f251, [%rd56+8];
- sub.f32 %f56, %f251, %f179;
- ld.f32 %f252, [%rd56+12];
- sub.f32 %f57, %f252, %f180;
- ld.f32 %f253, [%rd56+16];
- sub.f32 %f58, %f253, %f181;
- mul.f32 %f254, %f57, %f2;
- fma.rn.f32 %f255, %f56, %f1, %f254;
- fma.rn.f32 %f59, %f58, %f3, %f255;
- setp.lt.f32 %p42, %f59, 0f00000000;
- mov.pred %p85, -1;
- mov.f32 %f444, 0f4CBEBC20;
- @%p42 bra BB0_40;
- bra.uni BB0_41;
- BB0_40:
- mov.f32 %f445, %f444;
- bra.uni BB0_43;
- BB0_41:
- mul.f32 %f258, %f57, %f57;
- fma.rn.f32 %f259, %f56, %f56, %f258;
- fma.rn.f32 %f260, %f58, %f58, %f259;
- mul.f32 %f261, %f59, %f59;
- sub.f32 %f60, %f260, %f261;
- ld.f32 %f61, [%rd16+16];
- setp.gt.f32 %p44, %f60, %f61;
- mov.f32 %f445, %f444;
- @%p44 bra BB0_43;
- sub.f32 %f262, %f61, %f60;
- sqrt.rn.f32 %f263, %f262;
- sub.f32 %f445, %f59, %f263;
- add.f32 %f444, %f59, %f263;
- mov.pred %p85, 0;
- BB0_43:
- @%p85 bra BB0_46;
- setp.lt.f32 %p46, %f445, 0f00000000;
- selp.f32 %f66, %f444, %f445, %p46;
- setp.geu.f32 %p47, %f66, %f433;
- @%p47 bra BB0_46;
- mul.lo.s64 %rd57, %rd15, 80;
- add.s64 %rd83, %rd35, %rd57;
- mov.f32 %f433, %f66;
- BB0_46:
- add.s32 %r22, %r36, 2;
- cvt.u64.u32 %rd19, %r22;
- mul.wide.u32 %rd58, %r22, 80;
- add.s64 %rd59, %rd35, %rd58;
- add.s64 %rd20, %rd59, 8;
- ld.f32 %f266, [%rd59+8];
- sub.f32 %f68, %f266, %f179;
- ld.f32 %f267, [%rd59+12];
- sub.f32 %f69, %f267, %f180;
- ld.f32 %f268, [%rd59+16];
- sub.f32 %f70, %f268, %f181;
- mul.f32 %f269, %f69, %f2;
- fma.rn.f32 %f270, %f68, %f1, %f269;
- fma.rn.f32 %f71, %f70, %f3, %f270;
- setp.lt.f32 %p49, %f71, 0f00000000;
- mov.pred %p86, -1;
- mov.f32 %f447, 0f4CBEBC20;
- @%p49 bra BB0_47;
- bra.uni BB0_48;
- BB0_47:
- mov.f32 %f448, %f447;
- bra.uni BB0_50;
- BB0_48:
- mul.f32 %f273, %f69, %f69;
- fma.rn.f32 %f274, %f68, %f68, %f273;
- fma.rn.f32 %f275, %f70, %f70, %f274;
- mul.f32 %f276, %f71, %f71;
- sub.f32 %f72, %f275, %f276;
- ld.f32 %f73, [%rd20+16];
- setp.gt.f32 %p51, %f72, %f73;
- mov.f32 %f448, %f447;
- @%p51 bra BB0_50;
- sub.f32 %f277, %f73, %f72;
- sqrt.rn.f32 %f278, %f277;
- sub.f32 %f448, %f71, %f278;
- add.f32 %f447, %f71, %f278;
- mov.pred %p86, 0;
- BB0_50:
- @%p86 bra BB0_53;
- setp.lt.f32 %p53, %f448, 0f00000000;
- selp.f32 %f78, %f447, %f448, %p53;
- setp.geu.f32 %p54, %f78, %f433;
- @%p54 bra BB0_53;
- mul.lo.s64 %rd60, %rd19, 80;
- add.s64 %rd83, %rd35, %rd60;
- mov.f32 %f433, %f78;
- BB0_53:
- add.s32 %r23, %r36, 3;
- cvt.u64.u32 %rd23, %r23;
- mul.wide.u32 %rd61, %r23, 80;
- add.s64 %rd62, %rd35, %rd61;
- add.s64 %rd24, %rd62, 8;
- ld.f32 %f281, [%rd62+8];
- sub.f32 %f80, %f281, %f179;
- ld.f32 %f282, [%rd62+12];
- sub.f32 %f81, %f282, %f180;
- ld.f32 %f283, [%rd62+16];
- sub.f32 %f82, %f283, %f181;
- mul.f32 %f284, %f81, %f2;
- fma.rn.f32 %f285, %f80, %f1, %f284;
- fma.rn.f32 %f83, %f82, %f3, %f285;
- setp.lt.f32 %p56, %f83, 0f00000000;
- mov.pred %p87, -1;
- mov.f32 %f450, 0f4CBEBC20;
- @%p56 bra BB0_54;
- bra.uni BB0_55;
- BB0_54:
- mov.f32 %f451, %f450;
- bra.uni BB0_57;
- BB0_55:
- mul.f32 %f288, %f81, %f81;
- fma.rn.f32 %f289, %f80, %f80, %f288;
- fma.rn.f32 %f290, %f82, %f82, %f289;
- mul.f32 %f291, %f83, %f83;
- sub.f32 %f84, %f290, %f291;
- ld.f32 %f85, [%rd24+16];
- setp.gt.f32 %p58, %f84, %f85;
- mov.f32 %f451, %f450;
- @%p58 bra BB0_57;
- sub.f32 %f292, %f85, %f84;
- sqrt.rn.f32 %f293, %f292;
- sub.f32 %f451, %f83, %f293;
- add.f32 %f450, %f83, %f293;
- mov.pred %p87, 0;
- BB0_57:
- @%p87 bra BB0_60;
- setp.lt.f32 %p60, %f451, 0f00000000;
- selp.f32 %f90, %f450, %f451, %p60;
- setp.geu.f32 %p61, %f90, %f433;
- @%p61 bra BB0_60;
- mul.lo.s64 %rd63, %rd23, 80;
- add.s64 %rd83, %rd35, %rd63;
- mov.f32 %f433, %f90;
- BB0_60:
- cvt.u32.u64 %r24, %rd11;
- add.s32 %r36, %r24, 4;
- setp.lt.u32 %p62, %r36, %r13;
- @%p62 bra BB0_32;
- BB0_61:
- setp.eq.s64 %p63, %rd83, 0;
- mov.f32 %f477, 0f40000000;
- mov.f32 %f478, %f477;
- mov.f32 %f479, %f477;
- @%p63 bra BB0_89;
- cvta.to.local.u64 %rd64, %rd34;
- ld.local.f32 %f93, [%rd64];
- ld.local.f32 %f94, [%rd64+4];
- ld.local.f32 %f95, [%rd64+8];
- fma.rn.f32 %f96, %f433, %f93, %f179;
- fma.rn.f32 %f97, %f433, %f94, %f180;
- fma.rn.f32 %f98, %f433, %f95, %f181;
- add.s64 %rd28, %rd83, 8;
- ld.f32 %f297, [%rd83+8];
- sub.f32 %f454, %f96, %f297;
- ld.f32 %f298, [%rd83+12];
- sub.f32 %f455, %f97, %f298;
- ld.f32 %f299, [%rd83+16];
- sub.f32 %f456, %f98, %f299;
- mul.f32 %f300, %f455, %f455;
- fma.rn.f32 %f301, %f454, %f454, %f300;
- fma.rn.f32 %f102, %f456, %f456, %f301;
- setp.leu.f32 %p64, %f102, 0f00000000;
- @%p64 bra BB0_64;
- sqrt.rn.f32 %f302, %f102;
- rcp.rn.f32 %f303, %f302;
- mul.f32 %f454, %f454, %f303;
- mul.f32 %f455, %f455, %f303;
- mul.f32 %f456, %f456, %f303;
- BB0_64:
- mul.f32 %f304, %f455, %f94;
- fma.rn.f32 %f305, %f454, %f93, %f304;
- fma.rn.f32 %f306, %f456, %f95, %f305;
- mov.u16 %rs4, 0;
- setp.leu.f32 %p65, %f306, 0f00000000;
- @%p65 bra BB0_66;
- neg.f32 %f454, %f454;
- neg.f32 %f455, %f455;
- neg.f32 %f456, %f456;
- mov.u16 %rs4, 1;
- BB0_66:
- ld.f32 %f307, [%rd28+44];
- setp.gt.f32 %p66, %f307, 0f00000000;
- @%p66 bra BB0_68;
- ld.f32 %f308, [%rd28+48];
- setp.leu.f32 %p67, %f308, 0f00000000;
- @%p67 bra BB0_69;
- BB0_68:
- cvta.to.local.u64 %rd65, %rd36;
- ld.local.u32 %r25, [%rd65];
- setp.lt.s32 %p68, %r25, 5;
- @%p68 bra BB0_81;
- bra.uni BB0_69;
- BB0_81:
- mul.f32 %f351, %f455, %f94;
- fma.rn.f32 %f352, %f454, %f93, %f351;
- fma.rn.f32 %f353, %f456, %f95, %f352;
- add.f32 %f354, %f353, 0f3F800000;
- mov.u32 %r29, 3;
- // Callseq Start 0
- {
- .reg .b32 temp_param_reg;
- // <end>}
- .param .b32 param0;
- st.param.f32 [param0+0], %f354;
- .param .b32 param1;
- st.param.b32 [param1+0], %r29;
- .param .b32 retval0;
- call.uni (retval0),
- _Z5powiffi,
- (
- param0,
- param1
- );
- ld.param.f32 %f355, [retval0+0];
- //{
- }// Callseq End 0
- fma.rn.f32 %f143, %f355, 0f3F666666, 0f3DCCCCCD;
- ld.local.f32 %f356, [%rd64];
- ld.local.f32 %f357, [%rd64+4];
- mul.f32 %f358, %f455, %f357;
- fma.rn.f32 %f359, %f454, %f356, %f358;
- ld.local.f32 %f360, [%rd64+8];
- fma.rn.f32 %f361, %f456, %f360, %f359;
- add.f32 %f362, %f454, %f454;
- mul.f32 %f363, %f362, %f361;
- add.f32 %f364, %f455, %f455;
- mul.f32 %f365, %f364, %f361;
- add.f32 %f366, %f456, %f456;
- mul.f32 %f367, %f366, %f361;
- sub.f32 %f144, %f356, %f363;
- sub.f32 %f145, %f357, %f365;
- sub.f32 %f146, %f360, %f367;
- add.u64 %rd71, %SP, 4;
- add.u64 %rd32, %SPL, 4;
- st.local.f32 [%rd32], %f144;
- st.local.f32 [%rd32+4], %f145;
- st.local.f32 [%rd32+8], %f146;
- mul.f32 %f368, %f145, %f145;
- fma.rn.f32 %f369, %f144, %f144, %f368;
- fma.rn.f32 %f147, %f146, %f146, %f369;
- setp.leu.f32 %p77, %f147, 0f00000000;
- @%p77 bra BB0_83;
- sqrt.rn.f32 %f370, %f147;
- rcp.rn.f32 %f371, %f370;
- mul.f32 %f372, %f371, %f144;
- st.local.f32 [%rd32], %f372;
- mul.f32 %f373, %f371, %f145;
- st.local.f32 [%rd32+4], %f373;
- mul.f32 %f374, %f371, %f146;
- st.local.f32 [%rd32+8], %f374;
- BB0_83:
- mul.f32 %f148, %f454, 0f38D1B717;
- add.f32 %f378, %f96, %f148;
- mul.f32 %f149, %f455, 0f38D1B717;
- add.f32 %f379, %f97, %f149;
- mul.f32 %f150, %f456, 0f38D1B717;
- add.f32 %f380, %f98, %f150;
- ld.local.u32 %r30, [%rd65];
- add.s32 %r31, %r30, 1;
- add.u64 %rd73, %SP, 0;
- add.u64 %rd74, %SPL, 0;
- st.local.u32 [%rd74], %r31;
- // Callseq Start 1
- {
- .reg .b32 temp_param_reg;
- // <end>}
- .param .b32 param0;
- st.param.f32 [param0+0], %f378;
- .param .b32 param1;
- st.param.f32 [param1+0], %f379;
- .param .b32 param2;
- st.param.f32 [param2+0], %f380;
- .param .b64 param3;
- st.param.b64 [param3+0], %rd71;
- .param .b64 param4;
- st.param.b64 [param4+0], %rd35;
- .param .b32 param5;
- st.param.b32 [param5+0], %r13;
- .param .b64 param6;
- st.param.b64 [param6+0], %rd73;
- .param .align 4 .b8 retval0[12];
- call.uni (retval0),
- _Z5traceRK4Vec3IfES2_PK6SpherejRKi,
- (
- param0,
- param1,
- param2,
- param3,
- param4,
- param5,
- param6
- );
- ld.param.f32 %f151, [retval0+0];
- ld.param.f32 %f152, [retval0+4];
- ld.param.f32 %f153, [retval0+8];
- //{
- }// Callseq End 1
- ld.f32 %f470, [%rd28+44];
- mov.f32 %f471, 0f00000000;
- setp.eq.f32 %p78, %f470, 0f00000000;
- mov.f32 %f472, %f471;
- mov.f32 %f473, %f471;
- @%p78 bra BB0_87;
- setp.eq.s16 %p79, %rs4, 0;
- selp.f32 %f381, 0f3F68BA2E, 0f3F8CCCCD, %p79;
- ld.local.f32 %f382, [%rd64];
- ld.local.f32 %f383, [%rd64+4];
- mul.f32 %f384, %f455, %f383;
- fma.rn.f32 %f385, %f454, %f382, %f384;
- ld.local.f32 %f386, [%rd64+8];
- fma.rn.f32 %f387, %f456, %f386, %f385;
- mul.f32 %f388, %f381, %f381;
- mul.f32 %f389, %f387, %f387;
- mov.f32 %f390, 0f3F800000;
- sub.f32 %f391, %f390, %f389;
- mul.f32 %f392, %f388, %f391;
- sub.f32 %f393, %f390, %f392;
- mul.f32 %f394, %f381, %f387;
- neg.f32 %f395, %f394;
- sqrt.rn.f32 %f396, %f393;
- sub.f32 %f397, %f395, %f396;
- mul.f32 %f398, %f454, %f397;
- mul.f32 %f399, %f455, %f397;
- mul.f32 %f400, %f456, %f397;
- fma.rn.f32 %f155, %f381, %f382, %f398;
- fma.rn.f32 %f156, %f381, %f383, %f399;
- fma.rn.f32 %f157, %f381, %f386, %f400;
- add.u64 %rd77, %SP, 20;
- add.u64 %rd33, %SPL, 20;
- st.local.f32 [%rd33], %f155;
- st.local.f32 [%rd33+4], %f156;
- st.local.f32 [%rd33+8], %f157;
- mul.f32 %f401, %f156, %f156;
- fma.rn.f32 %f402, %f155, %f155, %f401;
- fma.rn.f32 %f158, %f157, %f157, %f402;
- setp.leu.f32 %p80, %f158, 0f00000000;
- @%p80 bra BB0_86;
- sqrt.rn.f32 %f403, %f158;
- rcp.rn.f32 %f404, %f403;
- mul.f32 %f405, %f404, %f155;
- st.local.f32 [%rd33], %f405;
- mul.f32 %f406, %f404, %f156;
- st.local.f32 [%rd33+4], %f406;
- mul.f32 %f407, %f404, %f157;
- st.local.f32 [%rd33+8], %f407;
- BB0_86:
- sub.f32 %f408, %f96, %f148;
- ld.local.u32 %r32, [%rd65];
- add.s32 %r33, %r32, 1;
- add.u64 %rd79, %SP, 16;
- add.u64 %rd80, %SPL, 16;
- st.local.u32 [%rd80], %r33;
- sub.f32 %f409, %f98, %f150;
- sub.f32 %f410, %f97, %f149;
- // Callseq Start 2
- {
- .reg .b32 temp_param_reg;
- // <end>}
- .param .b32 param0;
- st.param.f32 [param0+0], %f408;
- .param .b32 param1;
- st.param.f32 [param1+0], %f410;
- .param .b32 param2;
- st.param.f32 [param2+0], %f409;
- .param .b64 param3;
- st.param.b64 [param3+0], %rd77;
- .param .b64 param4;
- st.param.b64 [param4+0], %rd35;
- .param .b32 param5;
- st.param.b32 [param5+0], %r13;
- .param .b64 param6;
- st.param.b64 [param6+0], %rd79;
- .param .align 4 .b8 retval0[12];
- call.uni (retval0),
- _Z5traceRK4Vec3IfES2_PK6SpherejRKi,
- (
- param0,
- param1,
- param2,
- param3,
- param4,
- param5,
- param6
- );
- ld.param.f32 %f471, [retval0+0];
- ld.param.f32 %f472, [retval0+4];
- ld.param.f32 %f473, [retval0+8];
- //{
- }// Callseq End 2
- ld.f32 %f470, [%rd28+44];
- BB0_87:
- mov.f32 %f411, 0f3F800000;
- sub.f32 %f412, %f411, %f143;
- mul.f32 %f413, %f412, %f471;
- mul.f32 %f414, %f412, %f472;
- mul.f32 %f415, %f412, %f473;
- mul.f32 %f416, %f413, %f470;
- mul.f32 %f417, %f414, %f470;
- mul.f32 %f418, %f415, %f470;
- fma.rn.f32 %f419, %f143, %f151, %f416;
- fma.rn.f32 %f420, %f143, %f152, %f417;
- fma.rn.f32 %f421, %f143, %f153, %f418;
- ld.f32 %f422, [%rd28+20];
- mul.f32 %f474, %f422, %f419;
- ld.f32 %f423, [%rd28+24];
- mul.f32 %f475, %f420, %f423;
- ld.f32 %f424, [%rd28+28];
- mul.f32 %f476, %f421, %f424;
- bra.uni BB0_88;
- BB0_69:
- mov.f32 %f474, 0f00000000;
- mov.f32 %f475, %f474;
- mov.f32 %f476, %f474;
- @%p8 bra BB0_88;
- fma.rn.f32 %f115, %f454, 0f38D1B717, %f96;
- fma.rn.f32 %f116, %f455, 0f38D1B717, %f97;
- fma.rn.f32 %f117, %f456, 0f38D1B717, %f98;
- mov.f32 %f476, 0f00000000;
- mov.u32 %r38, 0;
- mov.f32 %f475, %f476;
- mov.f32 %f474, %f476;
- BB0_71:
- cvt.u64.u32 %rd29, %r38;
- mul.wide.u32 %rd66, %r38, 80;
- add.s64 %rd67, %rd35, %rd66;
- add.s64 %rd30, %rd67, 40;
- ld.f32 %f121, [%rd67+40];
- setp.leu.f32 %p70, %f121, 0f00000000;
- @%p70 bra BB0_80;
- ld.f32 %f315, [%rd30+-32];
- sub.f32 %f465, %f315, %f96;
- ld.f32 %f316, [%rd30+-28];
- sub.f32 %f464, %f316, %f97;
- ld.f32 %f317, [%rd30+-24];
- sub.f32 %f463, %f317, %f98;
- mul.f32 %f318, %f464, %f464;
- fma.rn.f32 %f319, %f465, %f465, %f318;
- fma.rn.f32 %f125, %f463, %f463, %f319;
- setp.leu.f32 %p71, %f125, 0f00000000;
- @%p71 bra BB0_74;
- sqrt.rn.f32 %f320, %f125;
- rcp.rn.f32 %f321, %f320;
- mul.f32 %f465, %f465, %f321;
- mul.f32 %f464, %f464, %f321;
- mul.f32 %f463, %f463, %f321;
- BB0_74:
- mov.u32 %r39, 0;
- BB0_75:
- cvt.u32.u64 %r28, %rd29;
- setp.eq.s32 %p72, %r28, %r39;
- @%p72 bra BB0_78;
- mul.wide.u32 %rd68, %r39, 80;
- add.s64 %rd69, %rd35, %rd68;
- add.s64 %rd31, %rd69, 8;
- ld.f32 %f322, [%rd69+8];
- sub.f32 %f132, %f322, %f115;
- ld.f32 %f323, [%rd69+12];
- sub.f32 %f133, %f323, %f116;
- ld.f32 %f324, [%rd69+16];
- sub.f32 %f134, %f324, %f117;
- mul.f32 %f325, %f464, %f133;
- fma.rn.f32 %f326, %f465, %f132, %f325;
- fma.rn.f32 %f135, %f463, %f134, %f326;
- setp.lt.f32 %p73, %f135, 0f00000000;
- @%p73 bra BB0_78;
- mul.f32 %f328, %f133, %f133;
- fma.rn.f32 %f329, %f132, %f132, %f328;
- fma.rn.f32 %f330, %f134, %f134, %f329;
- mul.f32 %f331, %f135, %f135;
- sub.f32 %f332, %f330, %f331;
- ld.f32 %f333, [%rd31+16];
- mov.f32 %f466, 0f00000000;
- setp.leu.f32 %p74, %f332, %f333;
- @%p74 bra BB0_79;
- BB0_78:
- add.s32 %r39, %r39, 1;
- setp.lt.u32 %p75, %r39, %r13;
- mov.f32 %f466, 0f3F800000;
- @%p75 bra BB0_75;
- BB0_79:
- ld.f32 %f335, [%rd28+20];
- mul.f32 %f336, %f466, %f335;
- ld.f32 %f337, [%rd28+24];
- mul.f32 %f338, %f466, %f337;
- ld.f32 %f339, [%rd28+28];
- mul.f32 %f340, %f466, %f339;
- mul.f32 %f341, %f454, %f465;
- fma.rn.f32 %f342, %f455, %f464, %f341;
- fma.rn.f32 %f343, %f456, %f463, %f342;
- mov.f32 %f344, 0f00000000;
- max.f32 %f345, %f344, %f343;
- mul.f32 %f346, %f336, %f345;
- mul.f32 %f347, %f338, %f345;
- mul.f32 %f348, %f340, %f345;
- ld.f32 %f349, [%rd30+4];
- ld.f32 %f350, [%rd30+8];
- fma.rn.f32 %f474, %f346, %f121, %f474;
- fma.rn.f32 %f475, %f347, %f349, %f475;
- fma.rn.f32 %f476, %f348, %f350, %f476;
- BB0_80:
- add.s32 %r38, %r38, 1;
- setp.lt.u32 %p76, %r38, %r13;
- @%p76 bra BB0_71;
- BB0_88:
- ld.f32 %f425, [%rd28+32];
- add.f32 %f477, %f474, %f425;
- ld.f32 %f426, [%rd28+36];
- add.f32 %f478, %f475, %f426;
- ld.f32 %f427, [%rd28+40];
- add.f32 %f479, %f476, %f427;
- BB0_89:
- st.param.f32 [func_retval0+0], %f477;
- st.param.f32 [func_retval0+4], %f478;
- st.param.f32 [func_retval0+8], %f479;
- ret;
- }
- // .globl raytrace_kernel
- .visible .entry raytrace_kernel(
- .param .u32 raytrace_kernel_param_0,
- .param .u32 raytrace_kernel_param_1,
- .param .u64 raytrace_kernel_param_2,
- .param .u64 raytrace_kernel_param_3,
- .param .u32 raytrace_kernel_param_4,
- .param .f32 raytrace_kernel_param_5,
- .param .f32 raytrace_kernel_param_6,
- .param .f32 raytrace_kernel_param_7,
- .param .f32 raytrace_kernel_param_8
- )
- {
- .local .align 4 .b8 __local_depot1[56];
- .reg .b64 %SP;
- .reg .b64 %SPL;
- .reg .pred %p<263>;
- .reg .b16 %rs<13>;
- .reg .f32 %f<1399>;
- .reg .b32 %r<148>;
- .reg .f64 %fd<18>;
- .reg .b64 %rd<391>;
- mov.u64 %SPL, __local_depot1;
- cvta.local.u64 %SP, %SPL;
- ld.param.u32 %r36, [raytrace_kernel_param_0];
- ld.param.u32 %r38, [raytrace_kernel_param_1];
- ld.param.u64 %rd158, [raytrace_kernel_param_3];
- ld.param.u32 %r37, [raytrace_kernel_param_4];
- ld.param.f32 %f539, [raytrace_kernel_param_5];
- ld.param.f32 %f540, [raytrace_kernel_param_6];
- ld.param.f32 %f541, [raytrace_kernel_param_7];
- ld.param.f32 %f542, [raytrace_kernel_param_8];
- cvta.to.global.u64 %rd1, %rd158;
- mov.u32 %r39, %ntid.x;
- mov.u32 %r40, %ctaid.x;
- mov.u32 %r41, %tid.x;
- mad.lo.s32 %r1, %r39, %r40, %r41;
- mov.u32 %r42, %ntid.y;
- mov.u32 %r43, %ctaid.y;
- mov.u32 %r44, %tid.y;
- mad.lo.s32 %r2, %r42, %r43, %r44;
- setp.ge.u32 %p22, %r2, %r38;
- setp.ge.u32 %p23, %r1, %r36;
- or.pred %p24, %p22, %p23;
- @%p24 bra BB1_271;
- cvt.rn.f64.s32 %fd1, %r1;
- add.f64 %fd2, %fd1, 0d3FE0000000000000;
- cvt.f64.f32 %fd3, %f539;
- mul.f64 %fd4, %fd3, %fd2;
- fma.rn.f64 %fd5, %fd4, 0d4000000000000000, 0dBFF0000000000000;
- cvt.f64.f32 %fd6, %f542;
- mul.f64 %fd7, %fd6, %fd5;
- cvt.f64.f32 %fd8, %f541;
- mul.f64 %fd9, %fd8, %fd7;
- cvt.rn.f32.f64 %f9, %fd9;
- cvt.rn.f64.s32 %fd10, %r2;
- add.f64 %fd11, %fd10, 0d3FE0000000000000;
- cvt.f64.f32 %fd12, %f540;
- mul.f64 %fd13, %fd12, %fd11;
- fma.rn.f64 %fd14, %fd12, %fd11, %fd13;
- mov.f64 %fd15, 0d3FF0000000000000;
- sub.f64 %fd16, %fd15, %fd14;
- mul.f64 %fd17, %fd6, %fd16;
- cvt.rn.f32.f64 %f8, %fd17;
- mul.f32 %f544, %f8, %f8;
- fma.rn.f32 %f545, %f9, %f9, %f544;
- add.f32 %f3, %f545, 0f3F800000;
- mov.f32 %f7, 0fBF800000;
- setp.leu.f32 %p25, %f3, 0f00000000;
- @%p25 bra BB1_3;
- sqrt.rn.f32 %f546, %f3;
- rcp.rn.f32 %f547, %f546;
- mul.f32 %f9, %f9, %f547;
- mul.f32 %f8, %f8, %f547;
- neg.f32 %f7, %f547;
- BB1_3:
- setp.eq.s32 %p26, %r37, 0;
- mov.f32 %f1263, 0f4CBEBC20;
- mov.u64 %rd328, 0;
- mov.u64 %rd327, %rd328;
- @%p26 bra BB1_63;
- and.b32 %r48, %r37, 3;
- mov.f32 %f1263, 0f4CBEBC20;
- mov.u64 %rd328, 0;
- mov.u32 %r132, 0;
- mov.f32 %f1250, 0f00000000;
- setp.eq.s32 %p27, %r48, 0;
- mov.u64 %rd327, %rd328;
- @%p27 bra BB1_32;
- mov.f32 %f1263, 0f4CBEBC20;
- mov.u64 %rd328, 0;
- mov.u32 %r131, 0;
- and.b32 %r126, %r37, 3;
- setp.eq.s32 %p28, %r126, 1;
- mov.u64 %rd327, %rd328;
- @%p28 bra BB1_24;
- mov.f32 %f1263, 0f4CBEBC20;
- mov.u64 %rd328, 0;
- mov.u32 %r130, 0;
- and.b32 %r128, %r37, 3;
- setp.eq.s32 %p29, %r128, 2;
- mov.u64 %rd327, %rd328;
- @%p29 bra BB1_16;
- ld.global.f32 %f10, [%rd1+8];
- ld.global.f32 %f11, [%rd1+12];
- mul.f32 %f555, %f8, %f11;
- fma.rn.f32 %f556, %f9, %f10, %f555;
- ld.global.f32 %f12, [%rd1+16];
- fma.rn.f32 %f13, %f7, %f12, %f556;
- setp.lt.f32 %p31, %f13, 0f00000000;
- mov.pred %p242, -1;
- mov.f32 %f1239, 0f4CBEBC20;
- @%p31 bra BB1_8;
- bra.uni BB1_9;
- BB1_8:
- mov.f32 %f1240, %f1239;
- bra.uni BB1_11;
- BB1_9:
- mov.f32 %f1239, 0f4CBEBC20;
- add.s64 %rd286, %rd1, 8;
- mul.f32 %f559, %f11, %f11;
- fma.rn.f32 %f560, %f10, %f10, %f559;
- fma.rn.f32 %f561, %f12, %f12, %f560;
- mul.f32 %f562, %f13, %f13;
- sub.f32 %f14, %f561, %f562;
- ld.global.f32 %f15, [%rd286+16];
- setp.gt.f32 %p33, %f14, %f15;
- mov.f32 %f1240, %f1239;
- @%p33 bra BB1_11;
- sub.f32 %f563, %f15, %f14;
- sqrt.rn.f32 %f564, %f563;
- sub.f32 %f1240, %f13, %f564;
- add.f32 %f1239, %f13, %f564;
- mov.pred %p242, 0;
- BB1_11:
- mov.u64 %rd328, 0;
- mov.f32 %f1263, 0f4CBEBC20;
- mov.u32 %r130, 1;
- @%p242 bra BB1_12;
- bra.uni BB1_13;
- BB1_12:
- mov.u64 %rd327, %rd328;
- bra.uni BB1_16;
- BB1_13:
- mov.u64 %rd328, 0;
- mov.f32 %f1219, 0f4CBEBC20;
- setp.lt.f32 %p35, %f1240, 0f00000000;
- selp.f32 %f1263, %f1239, %f1240, %p35;
- setp.geu.f32 %p36, %f1263, 0f4CBEBC20;
- @%p36 bra BB1_14;
- mov.u64 %rd328, %rd158;
- mov.u64 %rd327, %rd1;
- bra.uni BB1_16;
- BB1_14:
- mov.u64 %rd327, %rd328;
- mov.f32 %f1263, %f1219;
- BB1_16:
- mov.f32 %f1242, 0f4CBEBC20;
- mul.wide.u32 %rd171, %r130, 80;
- add.s64 %rd172, %rd1, %rd171;
- ld.global.f32 %f22, [%rd172+8];
- ld.global.f32 %f23, [%rd172+12];
- mul.f32 %f569, %f8, %f23;
- fma.rn.f32 %f570, %f9, %f22, %f569;
- ld.global.f32 %f24, [%rd172+16];
- fma.rn.f32 %f25, %f7, %f24, %f570;
- setp.lt.f32 %p38, %f25, 0f00000000;
- mov.pred %p243, -1;
- @%p38 bra BB1_17;
- bra.uni BB1_18;
- BB1_17:
- mov.f32 %f1243, %f1242;
- bra.uni BB1_20;
- BB1_18:
- mul.wide.u32 %rd282, %r130, 80;
- add.s64 %rd281, %rd1, %rd282;
- add.s64 %rd280, %rd281, 8;
- mov.f32 %f1242, 0f4CBEBC20;
- mul.f32 %f573, %f23, %f23;
- fma.rn.f32 %f574, %f22, %f22, %f573;
- fma.rn.f32 %f575, %f24, %f24, %f574;
- mul.f32 %f576, %f25, %f25;
- sub.f32 %f26, %f575, %f576;
- ld.global.f32 %f27, [%rd280+16];
- setp.gt.f32 %p40, %f26, %f27;
- mov.f32 %f1243, %f1242;
- @%p40 bra BB1_20;
- sub.f32 %f577, %f27, %f26;
- sqrt.rn.f32 %f578, %f577;
- sub.f32 %f1243, %f25, %f578;
- add.f32 %f1242, %f25, %f578;
- mov.pred %p243, 0;
- BB1_20:
- @%p243 bra BB1_23;
- setp.lt.f32 %p42, %f1243, 0f00000000;
- selp.f32 %f32, %f1242, %f1243, %p42;
- setp.geu.f32 %p43, %f32, %f1263;
- @%p43 bra BB1_23;
- cvt.u64.u32 %rd283, %r130;
- mul.lo.s64 %rd173, %rd283, 80;
- add.s64 %rd327, %rd1, %rd173;
- add.s64 %rd328, %rd158, %rd173;
- mov.f32 %f1263, %f32;
- BB1_23:
- cvt.u64.u32 %rd284, %r130;
- cvt.u32.u64 %r52, %rd284;
- add.s32 %r131, %r52, 1;
- BB1_24:
- mov.f32 %f1246, 0f4CBEBC20;
- mul.wide.u32 %rd174, %r131, 80;
- add.s64 %rd175, %rd1, %rd174;
- ld.global.f32 %f35, [%rd175+8];
- ld.global.f32 %f36, [%rd175+12];
- mul.f32 %f581, %f8, %f36;
- fma.rn.f32 %f582, %f9, %f35, %f581;
- ld.global.f32 %f37, [%rd175+16];
- fma.rn.f32 %f38, %f7, %f37, %f582;
- setp.lt.f32 %p45, %f38, 0f00000000;
- mov.pred %p244, -1;
- @%p45 bra BB1_25;
- bra.uni BB1_26;
- BB1_25:
- mov.f32 %f1247, %f1246;
- bra.uni BB1_28;
- BB1_26:
- mul.wide.u32 %rd291, %r131, 80;
- add.s64 %rd290, %rd1, %rd291;
- add.s64 %rd289, %rd290, 8;
- mov.f32 %f1246, 0f4CBEBC20;
- mul.f32 %f585, %f36, %f36;
- fma.rn.f32 %f586, %f35, %f35, %f585;
- fma.rn.f32 %f587, %f37, %f37, %f586;
- mul.f32 %f588, %f38, %f38;
- sub.f32 %f39, %f587, %f588;
- ld.global.f32 %f40, [%rd289+16];
- setp.gt.f32 %p47, %f39, %f40;
- mov.f32 %f1247, %f1246;
- @%p47 bra BB1_28;
- sub.f32 %f589, %f40, %f39;
- sqrt.rn.f32 %f590, %f589;
- sub.f32 %f1247, %f38, %f590;
- add.f32 %f1246, %f38, %f590;
- mov.pred %p244, 0;
- BB1_28:
- @%p244 bra BB1_31;
- setp.lt.f32 %p49, %f1247, 0f00000000;
- selp.f32 %f45, %f1246, %f1247, %p49;
- setp.geu.f32 %p50, %f45, %f1263;
- @%p50 bra BB1_31;
- cvt.u64.u32 %rd292, %r131;
- mul.lo.s64 %rd176, %rd292, 80;
- add.s64 %rd327, %rd1, %rd176;
- add.s64 %rd328, %rd158, %rd176;
- mov.f32 %f1263, %f45;
- BB1_31:
- cvt.u64.u32 %rd293, %r131;
- cvt.u32.u64 %r53, %rd293;
- add.s32 %r132, %r53, 1;
- mov.f32 %f1250, %f1263;
- BB1_32:
- setp.lt.u32 %p51, %r37, 4;
- @%p51 bra BB1_33;
- bra.uni BB1_34;
- BB1_33:
- mov.f32 %f1263, %f1250;
- bra.uni BB1_63;
- BB1_34:
- mul.wide.u32 %rd177, %r132, 80;
- add.s64 %rd178, %rd1, %rd177;
- ld.global.f32 %f50, [%rd178+8];
- ld.global.f32 %f51, [%rd178+12];
- mul.f32 %f593, %f8, %f51;
- fma.rn.f32 %f594, %f9, %f50, %f593;
- ld.global.f32 %f52, [%rd178+16];
- fma.rn.f32 %f53, %f7, %f52, %f594;
- setp.lt.f32 %p53, %f53, 0f00000000;
- mov.pred %p245, -1;
- mov.f32 %f1252, 0f4CBEBC20;
- @%p53 bra BB1_35;
- bra.uni BB1_36;
- BB1_35:
- mov.f32 %f1253, %f1252;
- bra.uni BB1_38;
- BB1_36:
- mul.wide.u32 %rd296, %r132, 80;
- add.s64 %rd295, %rd1, %rd296;
- add.s64 %rd294, %rd295, 8;
- mul.f32 %f597, %f51, %f51;
- fma.rn.f32 %f598, %f50, %f50, %f597;
- fma.rn.f32 %f599, %f52, %f52, %f598;
- mul.f32 %f600, %f53, %f53;
- sub.f32 %f54, %f599, %f600;
- ld.global.f32 %f55, [%rd294+16];
- setp.gt.f32 %p55, %f54, %f55;
- mov.f32 %f1253, %f1252;
- @%p55 bra BB1_38;
- sub.f32 %f601, %f55, %f54;
- sqrt.rn.f32 %f602, %f601;
- sub.f32 %f1253, %f53, %f602;
- add.f32 %f1252, %f53, %f602;
- mov.pred %p245, 0;
- BB1_38:
- @%p245 bra BB1_41;
- setp.lt.f32 %p57, %f1253, 0f00000000;
- selp.f32 %f60, %f1252, %f1253, %p57;
- setp.geu.f32 %p58, %f60, %f1263;
- @%p58 bra BB1_41;
- cvt.u64.u32 %rd297, %r132;
- mul.lo.s64 %rd179, %rd297, 80;
- add.s64 %rd327, %rd1, %rd179;
- add.s64 %rd328, %rd158, %rd179;
- mov.f32 %f1263, %f60;
- BB1_41:
- add.s32 %r54, %r132, 1;
- mul.wide.u32 %rd180, %r54, 80;
- add.s64 %rd181, %rd1, %rd180;
- ld.global.f32 %f62, [%rd181+8];
- ld.global.f32 %f63, [%rd181+12];
- mul.f32 %f605, %f8, %f63;
- fma.rn.f32 %f606, %f9, %f62, %f605;
- ld.global.f32 %f64, [%rd181+16];
- fma.rn.f32 %f65, %f7, %f64, %f606;
- setp.lt.f32 %p60, %f65, 0f00000000;
- mov.pred %p246, -1;
- mov.f32 %f1255, 0f4CBEBC20;
- @%p60 bra BB1_42;
- bra.uni BB1_43;
- BB1_42:
- mov.f32 %f1256, %f1255;
- bra.uni BB1_45;
- BB1_43:
- add.s32 %r118, %r132, 1;
- mul.wide.u32 %rd301, %r118, 80;
- add.s64 %rd300, %rd1, %rd301;
- add.s64 %rd299, %rd300, 8;
- mul.f32 %f609, %f63, %f63;
- fma.rn.f32 %f610, %f62, %f62, %f609;
- fma.rn.f32 %f611, %f64, %f64, %f610;
- mul.f32 %f612, %f65, %f65;
- sub.f32 %f66, %f611, %f612;
- ld.global.f32 %f67, [%rd299+16];
- setp.gt.f32 %p62, %f66, %f67;
- mov.f32 %f1256, %f1255;
- @%p62 bra BB1_45;
- sub.f32 %f613, %f67, %f66;
- sqrt.rn.f32 %f614, %f613;
- sub.f32 %f1256, %f65, %f614;
- add.f32 %f1255, %f65, %f614;
- mov.pred %p246, 0;
- BB1_45:
- @%p246 bra BB1_48;
- setp.lt.f32 %p64, %f1256, 0f00000000;
- selp.f32 %f72, %f1255, %f1256, %p64;
- setp.geu.f32 %p65, %f72, %f1263;
- @%p65 bra BB1_48;
- add.s32 %r119, %r132, 1;
- cvt.u64.u32 %rd302, %r119;
- mul.lo.s64 %rd182, %rd302, 80;
- add.s64 %rd327, %rd1, %rd182;
- add.s64 %rd328, %rd158, %rd182;
- mov.f32 %f1263, %f72;
- BB1_48:
- add.s32 %r55, %r132, 2;
- mul.wide.u32 %rd183, %r55, 80;
- add.s64 %rd184, %rd1, %rd183;
- ld.global.f32 %f74, [%rd184+8];
- ld.global.f32 %f75, [%rd184+12];
- mul.f32 %f617, %f8, %f75;
- fma.rn.f32 %f618, %f9, %f74, %f617;
- ld.global.f32 %f76, [%rd184+16];
- fma.rn.f32 %f77, %f7, %f76, %f618;
- setp.lt.f32 %p67, %f77, 0f00000000;
- mov.pred %p247, -1;
- mov.f32 %f1258, 0f4CBEBC20;
- @%p67 bra BB1_49;
- bra.uni BB1_50;
- BB1_49:
- mov.f32 %f1259, %f1258;
- bra.uni BB1_52;
- BB1_50:
- add.s32 %r120, %r132, 2;
- mul.wide.u32 %rd305, %r120, 80;
- add.s64 %rd304, %rd1, %rd305;
- add.s64 %rd303, %rd304, 8;
- mul.f32 %f621, %f75, %f75;
- fma.rn.f32 %f622, %f74, %f74, %f621;
- fma.rn.f32 %f623, %f76, %f76, %f622;
- mul.f32 %f624, %f77, %f77;
- sub.f32 %f78, %f623, %f624;
- ld.global.f32 %f79, [%rd303+16];
- setp.gt.f32 %p69, %f78, %f79;
- mov.f32 %f1259, %f1258;
- @%p69 bra BB1_52;
- sub.f32 %f625, %f79, %f78;
- sqrt.rn.f32 %f626, %f625;
- sub.f32 %f1259, %f77, %f626;
- add.f32 %f1258, %f77, %f626;
- mov.pred %p247, 0;
- BB1_52:
- @%p247 bra BB1_55;
- setp.lt.f32 %p71, %f1259, 0f00000000;
- selp.f32 %f84, %f1258, %f1259, %p71;
- setp.geu.f32 %p72, %f84, %f1263;
- @%p72 bra BB1_55;
- add.s32 %r121, %r132, 2;
- cvt.u64.u32 %rd306, %r121;
- mul.lo.s64 %rd185, %rd306, 80;
- add.s64 %rd327, %rd1, %rd185;
- add.s64 %rd328, %rd158, %rd185;
- mov.f32 %f1263, %f84;
- BB1_55:
- add.s32 %r56, %r132, 3;
- mul.wide.u32 %rd186, %r56, 80;
- add.s64 %rd187, %rd1, %rd186;
- ld.global.f32 %f86, [%rd187+8];
- ld.global.f32 %f87, [%rd187+12];
- mul.f32 %f629, %f8, %f87;
- fma.rn.f32 %f630, %f9, %f86, %f629;
- ld.global.f32 %f88, [%rd187+16];
- fma.rn.f32 %f89, %f7, %f88, %f630;
- setp.lt.f32 %p74, %f89, 0f00000000;
- mov.pred %p248, -1;
- mov.f32 %f1261, 0f4CBEBC20;
- @%p74 bra BB1_56;
- bra.uni BB1_57;
- BB1_56:
- mov.f32 %f1262, %f1261;
- bra.uni BB1_59;
- BB1_57:
- add.s32 %r122, %r132, 3;
- mul.wide.u32 %rd309, %r122, 80;
- add.s64 %rd308, %rd1, %rd309;
- add.s64 %rd307, %rd308, 8;
- mul.f32 %f633, %f87, %f87;
- fma.rn.f32 %f634, %f86, %f86, %f633;
- fma.rn.f32 %f635, %f88, %f88, %f634;
- mul.f32 %f636, %f89, %f89;
- sub.f32 %f90, %f635, %f636;
- ld.global.f32 %f91, [%rd307+16];
- setp.gt.f32 %p76, %f90, %f91;
- mov.f32 %f1262, %f1261;
- @%p76 bra BB1_59;
- sub.f32 %f637, %f91, %f90;
- sqrt.rn.f32 %f638, %f637;
- sub.f32 %f1262, %f89, %f638;
- add.f32 %f1261, %f89, %f638;
- mov.pred %p248, 0;
- BB1_59:
- @%p248 bra BB1_62;
- setp.lt.f32 %p78, %f1262, 0f00000000;
- selp.f32 %f96, %f1261, %f1262, %p78;
- setp.geu.f32 %p79, %f96, %f1263;
- @%p79 bra BB1_62;
- add.s32 %r123, %r132, 3;
- cvt.u64.u32 %rd310, %r123;
- mul.lo.s64 %rd188, %rd310, 80;
- add.s64 %rd327, %rd1, %rd188;
- add.s64 %rd328, %rd158, %rd188;
- mov.f32 %f1263, %f96;
- BB1_62:
- cvt.u64.u32 %rd298, %r132;
- cvt.u32.u64 %r57, %rd298;
- add.s32 %r132, %r57, 4;
- setp.lt.u32 %p80, %r132, %r37;
- @%p80 bra BB1_34;
- BB1_63:
- setp.eq.s64 %p81, %rd328, 0;
- mov.f32 %f1396, 0f40000000;
- mov.f32 %f1397, %f1396;
- mov.f32 %f1398, %f1396;
- @%p81 bra BB1_270;
- fma.rn.f32 %f99, %f9, %f1263, 0f00000000;
- fma.rn.f32 %f100, %f8, %f1263, 0f00000000;
- fma.rn.f32 %f101, %f7, %f1263, 0f00000000;
- add.s64 %rd49, %rd327, 8;
- ld.global.f32 %f642, [%rd327+8];
- sub.f32 %f1268, %f99, %f642;
- ld.global.f32 %f643, [%rd327+12];
- sub.f32 %f1269, %f100, %f643;
- ld.global.f32 %f644, [%rd327+16];
- sub.f32 %f1270, %f101, %f644;
- mul.f32 %f645, %f1269, %f1269;
- fma.rn.f32 %f646, %f1268, %f1268, %f645;
- fma.rn.f32 %f105, %f1270, %f1270, %f646;
- setp.leu.f32 %p82, %f105, 0f00000000;
- @%p82 bra BB1_66;
- sqrt.rn.f32 %f647, %f105;
- rcp.rn.f32 %f648, %f647;
- mul.f32 %f1268, %f1268, %f648;
- mul.f32 %f1269, %f1269, %f648;
- mul.f32 %f1270, %f1270, %f648;
- BB1_66:
- mul.f32 %f649, %f8, %f1269;
- fma.rn.f32 %f650, %f9, %f1268, %f649;
- fma.rn.f32 %f651, %f7, %f1270, %f650;
- mov.u16 %rs1, 0;
- setp.leu.f32 %p83, %f651, 0f00000000;
- @%p83 bra BB1_68;
- neg.f32 %f1268, %f1268;
- neg.f32 %f1269, %f1269;
- neg.f32 %f1270, %f1270;
- mov.u16 %rs1, 1;
- BB1_68:
- ld.global.f32 %f652, [%rd49+44];
- setp.gt.f32 %p84, %f652, 0f00000000;
- @%p84 bra BB1_82;
- ld.global.f32 %f653, [%rd49+48];
- setp.gt.f32 %p85, %f653, 0f00000000;
- @%p85 bra BB1_82;
- bra.uni BB1_70;
- BB1_82:
- mul.f32 %f696, %f8, %f1269;
- fma.rn.f32 %f697, %f9, %f1268, %f696;
- fma.rn.f32 %f146, %f7, %f1270, %f697;
- add.f32 %f698, %f146, 0f3F800000;
- mov.u32 %r61, 3;
- // Callseq Start 3
- {
- .reg .b32 temp_param_reg;
- // <end>}
- .param .b32 param0;
- st.param.f32 [param0+0], %f698;
- .param .b32 param1;
- st.param.b32 [param1+0], %r61;
- .param .b32 retval0;
- call.uni (retval0),
- _Z5powiffi,
- (
- param0,
- param1
- );
- ld.param.f32 %f699, [retval0+0];
- //{
- }// Callseq End 3
- fma.rn.f32 %f147, %f699, 0f3F666666, 0f3DCCCCCD;
- add.f32 %f700, %f1268, %f1268;
- mul.f32 %f701, %f700, %f146;
- add.f32 %f702, %f1269, %f1269;
- mul.f32 %f703, %f702, %f146;
- add.f32 %f704, %f1270, %f1270;
- mul.f32 %f705, %f704, %f146;
- sub.f32 %f1281, %f9, %f701;
- sub.f32 %f1282, %f8, %f703;
- sub.f32 %f1283, %f7, %f705;
- mul.f32 %f706, %f1282, %f1282;
- fma.rn.f32 %f707, %f1281, %f1281, %f706;
- fma.rn.f32 %f151, %f1283, %f1283, %f707;
- setp.leu.f32 %p94, %f151, 0f00000000;
- @%p94 bra BB1_84;
- sqrt.rn.f32 %f708, %f151;
- rcp.rn.f32 %f709, %f708;
- mul.f32 %f1281, %f1281, %f709;
- mul.f32 %f1282, %f1282, %f709;
- mul.f32 %f1283, %f1283, %f709;
- BB1_84:
- mul.f32 %f158, %f1268, 0f38D1B717;
- add.f32 %f159, %f99, %f158;
- mul.f32 %f160, %f1269, 0f38D1B717;
- add.f32 %f161, %f100, %f160;
- mul.f32 %f162, %f1270, 0f38D1B717;
- add.f32 %f163, %f101, %f162;
- mov.f32 %f1308, 0f4CBEBC20;
- mov.u64 %rd365, 0;
- mov.u64 %rd350, %rd365;
- @%p26 bra BB1_146;
- and.b32 %r65, %r37, 3;
- mov.f32 %f1308, 0f4CBEBC20;
- mov.u64 %rd365, 0;
- mov.u32 %r138, 0;
- mov.f32 %f1295, 0f00000000;
- setp.eq.s32 %p96, %r65, 0;
- @%p96 bra BB1_86;
- setp.eq.s32 %p97, %r65, 1;
- @%p97 bra BB1_88;
- bra.uni BB1_89;
- BB1_88:
- mov.u64 %rd350, %rd365;
- bra.uni BB1_107;
- BB1_70:
- mov.f32 %f1393, 0f00000000;
- mov.f32 %f1394, %f1393;
- mov.f32 %f1395, %f1393;
- @%p26 bra BB1_269;
- fma.rn.f32 %f118, %f1268, 0f38D1B717, %f99;
- fma.rn.f32 %f119, %f1269, 0f38D1B717, %f100;
- fma.rn.f32 %f120, %f1270, 0f38D1B717, %f101;
- mov.f32 %f1395, 0f00000000;
- mov.u32 %r134, 0;
- mov.f32 %f1394, %f1395;
- mov.f32 %f1393, %f1395;
- BB1_72:
- cvt.u64.u32 %rd50, %r134;
- mul.wide.u32 %rd189, %r134, 80;
- add.s64 %rd190, %rd1, %rd189;
- add.s64 %rd51, %rd190, 40;
- ld.global.f32 %f124, [%rd190+40];
- setp.leu.f32 %p87, %f124, 0f00000000;
- @%p87 bra BB1_81;
- ld.global.f32 %f660, [%rd51+-32];
- sub.f32 %f1276, %f660, %f99;
- ld.global.f32 %f661, [%rd51+-28];
- sub.f32 %f1275, %f661, %f100;
- ld.global.f32 %f662, [%rd51+-24];
- sub.f32 %f1274, %f662, %f101;
- mul.f32 %f663, %f1275, %f1275;
- fma.rn.f32 %f664, %f1276, %f1276, %f663;
- fma.rn.f32 %f128, %f1274, %f1274, %f664;
- setp.leu.f32 %p88, %f128, 0f00000000;
- @%p88 bra BB1_75;
- sqrt.rn.f32 %f665, %f128;
- rcp.rn.f32 %f666, %f665;
- mul.f32 %f1276, %f1276, %f666;
- mul.f32 %f1275, %f1275, %f666;
- mul.f32 %f1274, %f1274, %f666;
- BB1_75:
- mov.u32 %r135, 0;
- BB1_76:
- setp.eq.s32 %p89, %r134, %r135;
- @%p89 bra BB1_79;
- mul.wide.u32 %rd191, %r135, 80;
- add.s64 %rd192, %rd1, %rd191;
- add.s64 %rd52, %rd192, 8;
- ld.global.f32 %f667, [%rd192+8];
- sub.f32 %f135, %f667, %f118;
- ld.global.f32 %f668, [%rd192+12];
- sub.f32 %f136, %f668, %f119;
- ld.global.f32 %f669, [%rd192+16];
- sub.f32 %f137, %f669, %f120;
- mul.f32 %f670, %f1275, %f136;
- fma.rn.f32 %f671, %f1276, %f135, %f670;
- fma.rn.f32 %f138, %f1274, %f137, %f671;
- setp.lt.f32 %p90, %f138, 0f00000000;
- @%p90 bra BB1_79;
- mul.f32 %f673, %f136, %f136;
- fma.rn.f32 %f674, %f135, %f135, %f673;
- fma.rn.f32 %f675, %f137, %f137, %f674;
- mul.f32 %f676, %f138, %f138;
- sub.f32 %f677, %f675, %f676;
- ld.global.f32 %f678, [%rd52+16];
- mov.f32 %f1277, 0f00000000;
- setp.leu.f32 %p91, %f677, %f678;
- @%p91 bra BB1_80;
- BB1_79:
- add.s32 %r135, %r135, 1;
- setp.lt.u32 %p92, %r135, %r37;
- mov.f32 %f1277, 0f3F800000;
- @%p92 bra BB1_76;
- BB1_80:
- ld.global.f32 %f680, [%rd49+20];
- mul.f32 %f681, %f1277, %f680;
- ld.global.f32 %f682, [%rd49+24];
- mul.f32 %f683, %f1277, %f682;
- ld.global.f32 %f684, [%rd49+28];
- mul.f32 %f685, %f1277, %f684;
- mul.f32 %f686, %f1268, %f1276;
- fma.rn.f32 %f687, %f1269, %f1275, %f686;
- fma.rn.f32 %f688, %f1270, %f1274, %f687;
- mov.f32 %f689, 0f00000000;
- max.f32 %f690, %f689, %f688;
- mul.f32 %f691, %f681, %f690;
- mul.f32 %f692, %f683, %f690;
- mul.f32 %f693, %f685, %f690;
- ld.global.f32 %f694, [%rd51+4];
- ld.global.f32 %f695, [%rd51+8];
- fma.rn.f32 %f1393, %f691, %f124, %f1393;
- fma.rn.f32 %f1394, %f692, %f694, %f1394;
- fma.rn.f32 %f1395, %f693, %f695, %f1395;
- BB1_81:
- cvt.u32.u64 %r60, %rd50;
- add.s32 %r134, %r60, 1;
- setp.lt.u32 %p93, %r134, %r37;
- @%p93 bra BB1_72;
- bra.uni BB1_269;
- BB1_86:
- mov.u64 %rd350, %rd365;
- bra.uni BB1_115;
- BB1_89:
- setp.eq.s32 %p98, %r65, 2;
- mov.u64 %rd350, %rd365;
- @%p98 bra BB1_99;
- ld.global.f32 %f717, [%rd1+8];
- sub.f32 %f164, %f717, %f159;
- ld.global.f32 %f718, [%rd1+12];
- sub.f32 %f165, %f718, %f161;
- ld.global.f32 %f719, [%rd1+16];
- sub.f32 %f166, %f719, %f163;
- mul.f32 %f720, %f1282, %f165;
- fma.rn.f32 %f721, %f1281, %f164, %f720;
- fma.rn.f32 %f167, %f1283, %f166, %f721;
- setp.lt.f32 %p100, %f167, 0f00000000;
- mov.pred %p249, -1;
- mov.f32 %f1284, 0f4CBEBC20;
- @%p100 bra BB1_91;
- bra.uni BB1_92;
- BB1_91:
- mov.f32 %f1285, %f1284;
- bra.uni BB1_94;
- BB1_92:
- mov.f32 %f1284, 0f4CBEBC20;
- mul.f32 %f724, %f165, %f165;
- fma.rn.f32 %f725, %f164, %f164, %f724;
- fma.rn.f32 %f726, %f166, %f166, %f725;
- mul.f32 %f727, %f167, %f167;
- sub.f32 %f168, %f726, %f727;
- ld.global.f32 %f169, [%rd1+24];
- setp.gt.f32 %p102, %f168, %f169;
- mov.f32 %f1285, %f1284;
- @%p102 bra BB1_94;
- sub.f32 %f728, %f169, %f168;
- sqrt.rn.f32 %f729, %f728;
- sub.f32 %f1285, %f167, %f729;
- add.f32 %f1284, %f167, %f729;
- mov.pred %p249, 0;
- BB1_94:
- mov.u64 %rd365, 0;
- mov.f32 %f1308, 0f4CBEBC20;
- mov.u32 %r138, 1;
- @%p249 bra BB1_95;
- bra.uni BB1_96;
- BB1_95:
- mov.u64 %rd350, %rd365;
- bra.uni BB1_99;
- BB1_96:
- mov.u64 %rd365, 0;
- mov.f32 %f1234, 0f4CBEBC20;
- setp.lt.f32 %p104, %f1285, 0f00000000;
- selp.f32 %f1308, %f1284, %f1285, %p104;
- setp.geu.f32 %p105, %f1308, 0f4CBEBC20;
- @%p105 bra BB1_97;
- mov.u64 %rd365, %rd1;
- mov.u64 %rd350, %rd158;
- bra.uni BB1_99;
- BB1_97:
- mov.u64 %rd350, %rd365;
- mov.f32 %f1308, %f1234;
- BB1_99:
- mov.f32 %f1287, 0f4CBEBC20;
- mul.wide.u32 %rd205, %r138, 80;
- add.s64 %rd206, %rd1, %rd205;
- ld.global.f32 %f734, [%rd206+8];
- sub.f32 %f176, %f734, %f159;
- ld.global.f32 %f735, [%rd206+12];
- sub.f32 %f177, %f735, %f161;
- ld.global.f32 %f736, [%rd206+16];
- sub.f32 %f178, %f736, %f163;
- mul.f32 %f737, %f1282, %f177;
- fma.rn.f32 %f738, %f1281, %f176, %f737;
- fma.rn.f32 %f179, %f1283, %f178, %f738;
- setp.lt.f32 %p107, %f179, 0f00000000;
- mov.pred %p250, -1;
- @%p107 bra BB1_100;
- bra.uni BB1_101;
- BB1_100:
- mov.f32 %f1288, %f1287;
- bra.uni BB1_103;
- BB1_101:
- mul.wide.u32 %rd313, %r138, 80;
- add.s64 %rd312, %rd1, %rd313;
- add.s64 %rd311, %rd312, 8;
- mov.f32 %f1287, 0f4CBEBC20;
- mul.f32 %f741, %f177, %f177;
- fma.rn.f32 %f742, %f176, %f176, %f741;
- fma.rn.f32 %f743, %f178, %f178, %f742;
- mul.f32 %f744, %f179, %f179;
- sub.f32 %f180, %f743, %f744;
- ld.global.f32 %f181, [%rd311+16];
- setp.gt.f32 %p109, %f180, %f181;
- mov.f32 %f1288, %f1287;
- @%p109 bra BB1_103;
- sub.f32 %f745, %f181, %f180;
- sqrt.rn.f32 %f746, %f745;
- sub.f32 %f1288, %f179, %f746;
- add.f32 %f1287, %f179, %f746;
- mov.pred %p250, 0;
- BB1_103:
- @%p250 bra BB1_106;
- setp.lt.f32 %p111, %f1288, 0f00000000;
- selp.f32 %f186, %f1287, %f1288, %p111;
- setp.geu.f32 %p112, %f186, %f1308;
- @%p112 bra BB1_106;
- cvt.u64.u32 %rd314, %r138;
- mul.lo.s64 %rd207, %rd314, 80;
- add.s64 %rd365, %rd1, %rd207;
- add.s64 %rd350, %rd158, %rd207;
- mov.f32 %f1308, %f186;
- BB1_106:
- cvt.u64.u32 %rd315, %r138;
- cvt.u32.u64 %r69, %rd315;
- add.s32 %r138, %r69, 1;
- BB1_107:
- mov.f32 %f1291, 0f4CBEBC20;
- mul.wide.u32 %rd208, %r138, 80;
- add.s64 %rd209, %rd1, %rd208;
- ld.global.f32 %f749, [%rd209+8];
- sub.f32 %f189, %f749, %f159;
- ld.global.f32 %f750, [%rd209+12];
- sub.f32 %f190, %f750, %f161;
- ld.global.f32 %f751, [%rd209+16];
- sub.f32 %f191, %f751, %f163;
- mul.f32 %f752, %f1282, %f190;
- fma.rn.f32 %f753, %f1281, %f189, %f752;
- fma.rn.f32 %f192, %f1283, %f191, %f753;
- setp.lt.f32 %p114, %f192, 0f00000000;
- mov.pred %p251, -1;
- @%p114 bra BB1_108;
- bra.uni BB1_109;
- BB1_108:
- mov.f32 %f1292, %f1291;
- bra.uni BB1_111;
- BB1_109:
- mul.wide.u32 %rd318, %r138, 80;
- add.s64 %rd317, %rd1, %rd318;
- add.s64 %rd316, %rd317, 8;
- mov.f32 %f1291, 0f4CBEBC20;
- mul.f32 %f756, %f190, %f190;
- fma.rn.f32 %f757, %f189, %f189, %f756;
- fma.rn.f32 %f758, %f191, %f191, %f757;
- mul.f32 %f759, %f192, %f192;
- sub.f32 %f193, %f758, %f759;
- ld.global.f32 %f194, [%rd316+16];
- setp.gt.f32 %p116, %f193, %f194;
- mov.f32 %f1292, %f1291;
- @%p116 bra BB1_111;
- sub.f32 %f760, %f194, %f193;
- sqrt.rn.f32 %f761, %f760;
- sub.f32 %f1292, %f192, %f761;
- add.f32 %f1291, %f192, %f761;
- mov.pred %p251, 0;
- BB1_111:
- @%p251 bra BB1_114;
- setp.lt.f32 %p118, %f1292, 0f00000000;
- selp.f32 %f199, %f1291, %f1292, %p118;
- setp.geu.f32 %p119, %f199, %f1308;
- @%p119 bra BB1_114;
- cvt.u64.u32 %rd319, %r138;
- mul.lo.s64 %rd210, %rd319, 80;
- add.s64 %rd365, %rd1, %rd210;
- add.s64 %rd350, %rd158, %rd210;
- mov.f32 %f1308, %f199;
- BB1_114:
- cvt.u64.u32 %rd320, %r138;
- cvt.u32.u64 %r70, %rd320;
- add.s32 %r138, %r70, 1;
- mov.f32 %f1295, %f1308;
- BB1_115:
- setp.lt.u32 %p120, %r37, 4;
- @%p120 bra BB1_116;
- bra.uni BB1_117;
- BB1_116:
- mov.f32 %f1308, %f1295;
- bra.uni BB1_146;
- BB1_117:
- cvt.u64.u32 %rd73, %r138;
- mul.wide.u32 %rd211, %r138, 80;
- add.s64 %rd212, %rd1, %rd211;
- add.s64 %rd74, %rd212, 8;
- ld.global.f32 %f764, [%rd212+8];
- sub.f32 %f204, %f764, %f159;
- ld.global.f32 %f765, [%rd212+12];
- sub.f32 %f205, %f765, %f161;
- ld.global.f32 %f766, [%rd212+16];
- sub.f32 %f206, %f766, %f163;
- mul.f32 %f767, %f1282, %f205;
- fma.rn.f32 %f768, %f1281, %f204, %f767;
- fma.rn.f32 %f207, %f1283, %f206, %f768;
- setp.lt.f32 %p122, %f207, 0f00000000;
- mov.pred %p252, -1;
- mov.f32 %f1297, 0f4CBEBC20;
- @%p122 bra BB1_118;
- bra.uni BB1_119;
- BB1_118:
- mov.f32 %f1298, %f1297;
- bra.uni BB1_121;
- BB1_119:
- mul.f32 %f771, %f205, %f205;
- fma.rn.f32 %f772, %f204, %f204, %f771;
- fma.rn.f32 %f773, %f206, %f206, %f772;
- mul.f32 %f774, %f207, %f207;
- sub.f32 %f208, %f773, %f774;
- ld.global.f32 %f209, [%rd74+16];
- setp.gt.f32 %p124, %f208, %f209;
- mov.f32 %f1298, %f1297;
- @%p124 bra BB1_121;
- sub.f32 %f775, %f209, %f208;
- sqrt.rn.f32 %f776, %f775;
- sub.f32 %f1298, %f207, %f776;
- add.f32 %f1297, %f207, %f776;
- mov.pred %p252, 0;
- BB1_121:
- @%p252 bra BB1_124;
- setp.lt.f32 %p126, %f1298, 0f00000000;
- selp.f32 %f214, %f1297, %f1298, %p126;
- setp.geu.f32 %p127, %f214, %f1308;
- @%p127 bra BB1_124;
- mul.lo.s64 %rd213, %rd73, 80;
- add.s64 %rd365, %rd1, %rd213;
- add.s64 %rd350, %rd158, %rd213;
- mov.f32 %f1308, %f214;
- BB1_124:
- add.s32 %r71, %r138, 1;
- cvt.u64.u32 %rd79, %r71;
- mul.wide.u32 %rd214, %r71, 80;
- add.s64 %rd215, %rd1, %rd214;
- add.s64 %rd80, %rd215, 8;
- ld.global.f32 %f779, [%rd215+8];
- sub.f32 %f216, %f779, %f159;
- ld.global.f32 %f780, [%rd215+12];
- sub.f32 %f217, %f780, %f161;
- ld.global.f32 %f781, [%rd215+16];
- sub.f32 %f218, %f781, %f163;
- mul.f32 %f782, %f1282, %f217;
- fma.rn.f32 %f783, %f1281, %f216, %f782;
- fma.rn.f32 %f219, %f1283, %f218, %f783;
- setp.lt.f32 %p129, %f219, 0f00000000;
- mov.pred %p253, -1;
- mov.f32 %f1300, 0f4CBEBC20;
- @%p129 bra BB1_125;
- bra.uni BB1_126;
- BB1_125:
- mov.f32 %f1301, %f1300;
- bra.uni BB1_128;
- BB1_126:
- mul.f32 %f786, %f217, %f217;
- fma.rn.f32 %f787, %f216, %f216, %f786;
- fma.rn.f32 %f788, %f218, %f218, %f787;
- mul.f32 %f789, %f219, %f219;
- sub.f32 %f220, %f788, %f789;
- ld.global.f32 %f221, [%rd80+16];
- setp.gt.f32 %p131, %f220, %f221;
- mov.f32 %f1301, %f1300;
- @%p131 bra BB1_128;
- sub.f32 %f790, %f221, %f220;
- sqrt.rn.f32 %f791, %f790;
- sub.f32 %f1301, %f219, %f791;
- add.f32 %f1300, %f219, %f791;
- mov.pred %p253, 0;
- BB1_128:
- @%p253 bra BB1_131;
- setp.lt.f32 %p133, %f1301, 0f00000000;
- selp.f32 %f226, %f1300, %f1301, %p133;
- setp.geu.f32 %p134, %f226, %f1308;
- @%p134 bra BB1_131;
- mul.lo.s64 %rd216, %rd79, 80;
- add.s64 %rd365, %rd1, %rd216;
- add.s64 %rd350, %rd158, %rd216;
- mov.f32 %f1308, %f226;
- BB1_131:
- add.s32 %r72, %r138, 2;
- cvt.u64.u32 %rd85, %r72;
- mul.wide.u32 %rd217, %r72, 80;
- add.s64 %rd218, %rd1, %rd217;
- add.s64 %rd86, %rd218, 8;
- ld.global.f32 %f794, [%rd218+8];
- sub.f32 %f228, %f794, %f159;
- ld.global.f32 %f795, [%rd218+12];
- sub.f32 %f229, %f795, %f161;
- ld.global.f32 %f796, [%rd218+16];
- sub.f32 %f230, %f796, %f163;
- mul.f32 %f797, %f1282, %f229;
- fma.rn.f32 %f798, %f1281, %f228, %f797;
- fma.rn.f32 %f231, %f1283, %f230, %f798;
- setp.lt.f32 %p136, %f231, 0f00000000;
- mov.pred %p254, -1;
- mov.f32 %f1303, 0f4CBEBC20;
- @%p136 bra BB1_132;
- bra.uni BB1_133;
- BB1_132:
- mov.f32 %f1304, %f1303;
- bra.uni BB1_135;
- BB1_133:
- mul.f32 %f801, %f229, %f229;
- fma.rn.f32 %f802, %f228, %f228, %f801;
- fma.rn.f32 %f803, %f230, %f230, %f802;
- mul.f32 %f804, %f231, %f231;
- sub.f32 %f232, %f803, %f804;
- ld.global.f32 %f233, [%rd86+16];
- setp.gt.f32 %p138, %f232, %f233;
- mov.f32 %f1304, %f1303;
- @%p138 bra BB1_135;
- sub.f32 %f805, %f233, %f232;
- sqrt.rn.f32 %f806, %f805;
- sub.f32 %f1304, %f231, %f806;
- add.f32 %f1303, %f231, %f806;
- mov.pred %p254, 0;
- BB1_135:
- @%p254 bra BB1_138;
- setp.lt.f32 %p140, %f1304, 0f00000000;
- selp.f32 %f238, %f1303, %f1304, %p140;
- setp.geu.f32 %p141, %f238, %f1308;
- @%p141 bra BB1_138;
- mul.lo.s64 %rd219, %rd85, 80;
- add.s64 %rd365, %rd1, %rd219;
- add.s64 %rd350, %rd158, %rd219;
- mov.f32 %f1308, %f238;
- BB1_138:
- add.s32 %r73, %r138, 3;
- cvt.u64.u32 %rd91, %r73;
- mul.wide.u32 %rd220, %r73, 80;
- add.s64 %rd221, %rd1, %rd220;
- add.s64 %rd92, %rd221, 8;
- ld.global.f32 %f809, [%rd221+8];
- sub.f32 %f240, %f809, %f159;
- ld.global.f32 %f810, [%rd221+12];
- sub.f32 %f241, %f810, %f161;
- ld.global.f32 %f811, [%rd221+16];
- sub.f32 %f242, %f811, %f163;
- mul.f32 %f812, %f1282, %f241;
- fma.rn.f32 %f813, %f1281, %f240, %f812;
- fma.rn.f32 %f243, %f1283, %f242, %f813;
- setp.lt.f32 %p143, %f243, 0f00000000;
- mov.pred %p255, -1;
- mov.f32 %f1306, 0f4CBEBC20;
- @%p143 bra BB1_139;
- bra.uni BB1_140;
- BB1_139:
- mov.f32 %f1307, %f1306;
- bra.uni BB1_142;
- BB1_140:
- mul.f32 %f816, %f241, %f241;
- fma.rn.f32 %f817, %f240, %f240, %f816;
- fma.rn.f32 %f818, %f242, %f242, %f817;
- mul.f32 %f819, %f243, %f243;
- sub.f32 %f244, %f818, %f819;
- ld.global.f32 %f245, [%rd92+16];
- setp.gt.f32 %p145, %f244, %f245;
- mov.f32 %f1307, %f1306;
- @%p145 bra BB1_142;
- sub.f32 %f820, %f245, %f244;
- sqrt.rn.f32 %f821, %f820;
- sub.f32 %f1307, %f243, %f821;
- add.f32 %f1306, %f243, %f821;
- mov.pred %p255, 0;
- BB1_142:
- @%p255 bra BB1_145;
- setp.lt.f32 %p147, %f1307, 0f00000000;
- selp.f32 %f250, %f1306, %f1307, %p147;
- setp.geu.f32 %p148, %f250, %f1308;
- @%p148 bra BB1_145;
- mul.lo.s64 %rd222, %rd91, 80;
- add.s64 %rd365, %rd1, %rd222;
- add.s64 %rd350, %rd158, %rd222;
- mov.f32 %f1308, %f250;
- BB1_145:
- cvt.u32.u64 %r74, %rd73;
- add.s32 %r138, %r74, 4;
- setp.lt.u32 %p149, %r138, %r37;
- @%p149 bra BB1_117;
- BB1_146:
- setp.eq.s64 %p150, %rd350, 0;
- mov.f32 %f1333, 0f40000000;
- mov.f32 %f1334, %f1333;
- mov.f32 %f1335, %f1333;
- @%p150 bra BB1_173;
- fma.rn.f32 %f253, %f1281, %f1308, %f159;
- fma.rn.f32 %f254, %f1282, %f1308, %f161;
- fma.rn.f32 %f255, %f1283, %f1308, %f163;
- add.s64 %rd99, %rd365, 8;
- ld.global.f32 %f825, [%rd365+8];
- sub.f32 %f1310, %f253, %f825;
- ld.global.f32 %f826, [%rd365+12];
- sub.f32 %f1311, %f254, %f826;
- ld.global.f32 %f827, [%rd365+16];
- sub.f32 %f1312, %f255, %f827;
- mul.f32 %f828, %f1311, %f1311;
- fma.rn.f32 %f829, %f1310, %f1310, %f828;
- fma.rn.f32 %f259, %f1312, %f1312, %f829;
- setp.leu.f32 %p151, %f259, 0f00000000;
- @%p151 bra BB1_149;
- sqrt.rn.f32 %f830, %f259;
- rcp.rn.f32 %f831, %f830;
- mul.f32 %f1310, %f1310, %f831;
- mul.f32 %f1311, %f1311, %f831;
- mul.f32 %f1312, %f1312, %f831;
- BB1_149:
- mul.f32 %f832, %f1282, %f1311;
- fma.rn.f32 %f833, %f1281, %f1310, %f832;
- fma.rn.f32 %f834, %f1283, %f1312, %f833;
- mov.u16 %rs11, 0;
- setp.leu.f32 %p152, %f834, 0f00000000;
- @%p152 bra BB1_151;
- neg.f32 %f1310, %f1310;
- neg.f32 %f1311, %f1311;
- neg.f32 %f1312, %f1312;
- mov.u16 %rs11, 1;
- BB1_151:
- ld.global.f32 %f835, [%rd99+44];
- setp.gt.f32 %p153, %f835, 0f00000000;
- @%p153 bra BB1_165;
- ld.global.f32 %f836, [%rd99+48];
- setp.gt.f32 %p154, %f836, 0f00000000;
- @%p154 bra BB1_165;
- bra.uni BB1_153;
- BB1_165:
- mov.u32 %r124, 3;
- mul.f32 %f879, %f1282, %f1311;
- fma.rn.f32 %f880, %f1281, %f1310, %f879;
- fma.rn.f32 %f300, %f1283, %f1312, %f880;
- add.f32 %f881, %f300, 0f3F800000;
- // Callseq Start 4
- {
- .reg .b32 temp_param_reg;
- // <end>}
- .param .b32 param0;
- st.param.f32 [param0+0], %f881;
- .param .b32 param1;
- st.param.b32 [param1+0], %r124;
- .param .b32 retval0;
- call.uni (retval0),
- _Z5powiffi,
- (
- param0,
- param1
- );
- ld.param.f32 %f882, [retval0+0];
- //{
- }// Callseq End 4
- fma.rn.f32 %f301, %f882, 0f3F666666, 0f3DCCCCCD;
- add.f32 %f883, %f1310, %f1310;
- mul.f32 %f884, %f883, %f300;
- add.f32 %f885, %f1311, %f1311;
- mul.f32 %f886, %f885, %f300;
- add.f32 %f887, %f1312, %f1312;
- mul.f32 %f888, %f887, %f300;
- sub.f32 %f302, %f1281, %f884;
- sub.f32 %f303, %f1282, %f886;
- sub.f32 %f304, %f1283, %f888;
- add.u64 %rd227, %SP, 36;
- add.u64 %rd103, %SPL, 36;
- st.local.f32 [%rd103], %f302;
- st.local.f32 [%rd103+4], %f303;
- st.local.f32 [%rd103+8], %f304;
- mul.f32 %f889, %f303, %f303;
- fma.rn.f32 %f890, %f302, %f302, %f889;
- fma.rn.f32 %f305, %f304, %f304, %f890;
- setp.leu.f32 %p163, %f305, 0f00000000;
- @%p163 bra BB1_167;
- sqrt.rn.f32 %f891, %f305;
- rcp.rn.f32 %f892, %f891;
- mul.f32 %f893, %f302, %f892;
- st.local.f32 [%rd103], %f893;
- mul.f32 %f894, %f303, %f892;
- st.local.f32 [%rd103+4], %f894;
- mul.f32 %f895, %f304, %f892;
- st.local.f32 [%rd103+8], %f895;
- BB1_167:
- mul.f32 %f306, %f1310, 0f38D1B717;
- add.f32 %f899, %f253, %f306;
- mul.f32 %f307, %f1311, 0f38D1B717;
- add.f32 %f900, %f254, %f307;
- mul.f32 %f308, %f1312, 0f38D1B717;
- add.f32 %f901, %f255, %f308;
- add.u64 %rd228, %SP, 52;
- add.u64 %rd229, %SPL, 52;
- mov.u32 %r79, 2;
- st.local.u32 [%rd229], %r79;
- // Callseq Start 5
- {
- .reg .b32 temp_param_reg;
- // <end>}
- .param .b32 param0;
- st.param.f32 [param0+0], %f899;
- .param .b32 param1;
- st.param.f32 [param1+0], %f900;
- .param .b32 param2;
- st.param.f32 [param2+0], %f901;
- .param .b64 param3;
- st.param.b64 [param3+0], %rd227;
- .param .b64 param4;
- st.param.b64 [param4+0], %rd158;
- .param .b32 param5;
- st.param.b32 [param5+0], %r37;
- .param .b64 param6;
- st.param.b64 [param6+0], %rd228;
- .param .align 4 .b8 retval0[12];
- call.uni (retval0),
- _Z5traceRK4Vec3IfES2_PK6SpherejRKi,
- (
- param0,
- param1,
- param2,
- param3,
- param4,
- param5,
- param6
- );
- ld.param.f32 %f309, [retval0+0];
- ld.param.f32 %f310, [retval0+4];
- ld.param.f32 %f311, [retval0+8];
- //{
- }// Callseq End 5
- ld.global.f32 %f1326, [%rd99+44];
- mov.f32 %f1327, 0f00000000;
- setp.eq.f32 %p164, %f1326, 0f00000000;
- mov.f32 %f1328, %f1327;
- mov.f32 %f1329, %f1327;
- @%p164 bra BB1_171;
- setp.eq.s16 %p165, %rs11, 0;
- selp.f32 %f902, 0f3F68BA2E, 0f3F8CCCCD, %p165;
- mul.f32 %f903, %f902, %f902;
- mul.f32 %f904, %f300, %f300;
- mov.f32 %f905, 0f3F800000;
- sub.f32 %f906, %f905, %f904;
- mul.f32 %f907, %f903, %f906;
- sub.f32 %f908, %f905, %f907;
- mul.f32 %f909, %f300, %f902;
- neg.f32 %f910, %f909;
- sqrt.rn.f32 %f911, %f908;
- sub.f32 %f912, %f910, %f911;
- mul.f32 %f913, %f1310, %f912;
- mul.f32 %f914, %f1311, %f912;
- mul.f32 %f915, %f1312, %f912;
- fma.rn.f32 %f313, %f1281, %f902, %f913;
- fma.rn.f32 %f314, %f1282, %f902, %f914;
- fma.rn.f32 %f315, %f1283, %f902, %f915;
- add.u64 %rd231, %SP, 24;
- add.u64 %rd104, %SPL, 24;
- st.local.f32 [%rd104], %f313;
- st.local.f32 [%rd104+4], %f314;
- st.local.f32 [%rd104+8], %f315;
- mul.f32 %f916, %f314, %f314;
- fma.rn.f32 %f917, %f313, %f313, %f916;
- fma.rn.f32 %f316, %f315, %f315, %f917;
- setp.leu.f32 %p166, %f316, 0f00000000;
- @%p166 bra BB1_170;
- sqrt.rn.f32 %f918, %f316;
- rcp.rn.f32 %f919, %f918;
- mul.f32 %f920, %f313, %f919;
- st.local.f32 [%rd104], %f920;
- mul.f32 %f921, %f314, %f919;
- st.local.f32 [%rd104+4], %f921;
- mul.f32 %f922, %f315, %f919;
- st.local.f32 [%rd104+8], %f922;
- BB1_170:
- sub.f32 %f923, %f253, %f306;
- add.u64 %rd232, %SP, 48;
- add.u64 %rd233, %SPL, 48;
- st.local.u32 [%rd233], %r79;
- sub.f32 %f924, %f255, %f308;
- sub.f32 %f925, %f254, %f307;
- // Callseq Start 6
- {
- .reg .b32 temp_param_reg;
- // <end>}
- .param .b32 param0;
- st.param.f32 [param0+0], %f923;
- .param .b32 param1;
- st.param.f32 [param1+0], %f925;
- .param .b32 param2;
- st.param.f32 [param2+0], %f924;
- .param .b64 param3;
- st.param.b64 [param3+0], %rd231;
- .param .b64 param4;
- st.param.b64 [param4+0], %rd158;
- .param .b32 param5;
- st.param.b32 [param5+0], %r37;
- .param .b64 param6;
- st.param.b64 [param6+0], %rd232;
- .param .align 4 .b8 retval0[12];
- call.uni (retval0),
- _Z5traceRK4Vec3IfES2_PK6SpherejRKi,
- (
- param0,
- param1,
- param2,
- param3,
- param4,
- param5,
- param6
- );
- ld.param.f32 %f1327, [retval0+0];
- ld.param.f32 %f1328, [retval0+4];
- ld.param.f32 %f1329, [retval0+8];
- //{
- }// Callseq End 6
- ld.global.f32 %f1326, [%rd99+44];
- BB1_171:
- mov.f32 %f926, 0f3F800000;
- sub.f32 %f927, %f926, %f301;
- mul.f32 %f928, %f927, %f1327;
- mul.f32 %f929, %f927, %f1328;
- mul.f32 %f930, %f927, %f1329;
- mul.f32 %f931, %f1326, %f928;
- mul.f32 %f932, %f1326, %f929;
- mul.f32 %f933, %f1326, %f930;
- fma.rn.f32 %f934, %f301, %f309, %f931;
- fma.rn.f32 %f935, %f301, %f310, %f932;
- fma.rn.f32 %f936, %f301, %f311, %f933;
- ld.global.f32 %f937, [%rd99+20];
- mul.f32 %f1330, %f937, %f934;
- ld.global.f32 %f938, [%rd99+24];
- mul.f32 %f1331, %f935, %f938;
- ld.global.f32 %f939, [%rd99+28];
- mul.f32 %f1332, %f936, %f939;
- bra.uni BB1_172;
- BB1_153:
- mov.f32 %f1330, 0f00000000;
- mov.f32 %f1331, %f1330;
- mov.f32 %f1332, %f1330;
- @%p26 bra BB1_172;
- fma.rn.f32 %f272, %f1310, 0f38D1B717, %f253;
- fma.rn.f32 %f273, %f1311, 0f38D1B717, %f254;
- fma.rn.f32 %f274, %f1312, 0f38D1B717, %f255;
- mov.f32 %f1332, 0f00000000;
- mov.u32 %r140, 0;
- mov.f32 %f1331, %f1332;
- mov.f32 %f1330, %f1332;
- BB1_155:
- cvt.u64.u32 %rd100, %r140;
- mul.wide.u32 %rd223, %r140, 80;
- add.s64 %rd224, %rd1, %rd223;
- add.s64 %rd101, %rd224, 40;
- ld.global.f32 %f278, [%rd224+40];
- setp.leu.f32 %p156, %f278, 0f00000000;
- @%p156 bra BB1_164;
- ld.global.f32 %f843, [%rd101+-32];
- sub.f32 %f1321, %f843, %f253;
- ld.global.f32 %f844, [%rd101+-28];
- sub.f32 %f1320, %f844, %f254;
- ld.global.f32 %f845, [%rd101+-24];
- sub.f32 %f1319, %f845, %f255;
- mul.f32 %f846, %f1320, %f1320;
- fma.rn.f32 %f847, %f1321, %f1321, %f846;
- fma.rn.f32 %f282, %f1319, %f1319, %f847;
- setp.leu.f32 %p157, %f282, 0f00000000;
- @%p157 bra BB1_158;
- sqrt.rn.f32 %f848, %f282;
- rcp.rn.f32 %f849, %f848;
- mul.f32 %f1321, %f1321, %f849;
- mul.f32 %f1320, %f1320, %f849;
- mul.f32 %f1319, %f1319, %f849;
- BB1_158:
- mov.u32 %r141, 0;
- BB1_159:
- cvt.u32.u64 %r77, %rd100;
- setp.eq.s32 %p158, %r77, %r141;
- @%p158 bra BB1_162;
- mul.wide.u32 %rd225, %r141, 80;
- add.s64 %rd226, %rd1, %rd225;
- add.s64 %rd102, %rd226, 8;
- ld.global.f32 %f850, [%rd226+8];
- sub.f32 %f289, %f850, %f272;
- ld.global.f32 %f851, [%rd226+12];
- sub.f32 %f290, %f851, %f273;
- ld.global.f32 %f852, [%rd226+16];
- sub.f32 %f291, %f852, %f274;
- mul.f32 %f853, %f1320, %f290;
- fma.rn.f32 %f854, %f1321, %f289, %f853;
- fma.rn.f32 %f292, %f1319, %f291, %f854;
- setp.lt.f32 %p159, %f292, 0f00000000;
- @%p159 bra BB1_162;
- mul.f32 %f856, %f290, %f290;
- fma.rn.f32 %f857, %f289, %f289, %f856;
- fma.rn.f32 %f858, %f291, %f291, %f857;
- mul.f32 %f859, %f292, %f292;
- sub.f32 %f860, %f858, %f859;
- ld.global.f32 %f861, [%rd102+16];
- mov.f32 %f1322, 0f00000000;
- setp.leu.f32 %p160, %f860, %f861;
- @%p160 bra BB1_163;
- BB1_162:
- add.s32 %r141, %r141, 1;
- setp.lt.u32 %p161, %r141, %r37;
- mov.f32 %f1322, 0f3F800000;
- @%p161 bra BB1_159;
- BB1_163:
- ld.global.f32 %f863, [%rd99+20];
- mul.f32 %f864, %f1322, %f863;
- ld.global.f32 %f865, [%rd99+24];
- mul.f32 %f866, %f1322, %f865;
- ld.global.f32 %f867, [%rd99+28];
- mul.f32 %f868, %f1322, %f867;
- mul.f32 %f869, %f1310, %f1321;
- fma.rn.f32 %f870, %f1311, %f1320, %f869;
- fma.rn.f32 %f871, %f1312, %f1319, %f870;
- mov.f32 %f872, 0f00000000;
- max.f32 %f873, %f872, %f871;
- mul.f32 %f874, %f864, %f873;
- mul.f32 %f875, %f866, %f873;
- mul.f32 %f876, %f868, %f873;
- ld.global.f32 %f877, [%rd101+4];
- ld.global.f32 %f878, [%rd101+8];
- fma.rn.f32 %f1330, %f278, %f874, %f1330;
- fma.rn.f32 %f1331, %f875, %f877, %f1331;
- fma.rn.f32 %f1332, %f876, %f878, %f1332;
- BB1_164:
- add.s32 %r140, %r140, 1;
- setp.lt.u32 %p162, %r140, %r37;
- @%p162 bra BB1_155;
- BB1_172:
- ld.global.f32 %f940, [%rd99+32];
- add.f32 %f1333, %f1330, %f940;
- ld.global.f32 %f941, [%rd99+36];
- add.f32 %f1334, %f1331, %f941;
- ld.global.f32 %f942, [%rd99+40];
- add.f32 %f1335, %f1332, %f942;
- BB1_173:
- ld.global.f32 %f1389, [%rd49+44];
- mov.f32 %f1390, 0f00000000;
- setp.eq.f32 %p167, %f1389, 0f00000000;
- @%p167 bra BB1_174;
- setp.eq.s16 %p168, %rs1, 0;
- selp.f32 %f946, 0f3F68BA2E, 0f3F8CCCCD, %p168;
- mul.f32 %f947, %f946, %f946;
- mul.f32 %f948, %f146, %f146;
- mov.f32 %f949, 0f3F800000;
- sub.f32 %f950, %f949, %f948;
- mul.f32 %f951, %f947, %f950;
- sub.f32 %f952, %f949, %f951;
- mul.f32 %f953, %f146, %f946;
- neg.f32 %f954, %f953;
- sqrt.rn.f32 %f955, %f952;
- sub.f32 %f956, %f954, %f955;
- mul.f32 %f957, %f1268, %f956;
- mul.f32 %f958, %f1269, %f956;
- mul.f32 %f959, %f1270, %f956;
- fma.rn.f32 %f347, %f9, %f946, %f957;
- fma.rn.f32 %f346, %f8, %f946, %f958;
- fma.rn.f32 %f345, %f7, %f946, %f959;
- mul.f32 %f960, %f346, %f346;
- fma.rn.f32 %f961, %f347, %f347, %f960;
- fma.rn.f32 %f341, %f345, %f345, %f961;
- setp.leu.f32 %p169, %f341, 0f00000000;
- @%p169 bra BB1_177;
- sqrt.rn.f32 %f962, %f341;
- rcp.rn.f32 %f963, %f962;
- mul.f32 %f347, %f347, %f963;
- mul.f32 %f346, %f346, %f963;
- mul.f32 %f345, %f345, %f963;
- BB1_177:
- fma.rn.f32 %f1229, %f7, %f1263, 0f00000000;
- fma.rn.f32 %f1228, %f8, %f1263, 0f00000000;
- fma.rn.f32 %f1227, %f9, %f1263, 0f00000000;
- mul.f32 %f1226, %f1270, 0f38D1B717;
- mul.f32 %f1225, %f1269, 0f38D1B717;
- mul.f32 %f1224, %f1268, 0f38D1B717;
- sub.f32 %f348, %f1227, %f1224;
- sub.f32 %f349, %f1228, %f1225;
- sub.f32 %f350, %f1229, %f1226;
- mov.f32 %f1344, 0f4CBEBC20;
- mov.u64 %rd387, 0;
- mov.u64 %rd372, %rd387;
- @%p26 bra BB1_239;
- and.b32 %r84, %r37, 3;
- mov.f32 %f1346, 0f4CBEBC20;
- mov.u64 %rd387, 0;
- mov.u32 %r144, 0;
- mov.f32 %f1350, 0f00000000;
- setp.eq.s32 %p171, %r84, 0;
- @%p171 bra BB1_179;
- setp.eq.s32 %p172, %r84, 1;
- @%p172 bra BB1_181;
- bra.uni BB1_182;
- BB1_181:
- mov.u64 %rd372, %rd387;
- mov.f32 %f1344, %f1346;
- bra.uni BB1_200;
- BB1_179:
- mov.u64 %rd372, %rd387;
- mov.f32 %f1344, %f1346;
- bra.uni BB1_208;
- BB1_182:
- setp.eq.s32 %p173, %r84, 2;
- @%p173 bra BB1_183;
- bra.uni BB1_184;
- BB1_183:
- mov.u64 %rd372, %rd387;
- mov.f32 %f1344, %f1346;
- bra.uni BB1_192;
- BB1_184:
- ld.global.f32 %f971, [%rd1+8];
- sub.f32 %f351, %f971, %f348;
- ld.global.f32 %f972, [%rd1+12];
- sub.f32 %f352, %f972, %f349;
- ld.global.f32 %f973, [%rd1+16];
- sub.f32 %f353, %f973, %f350;
- mul.f32 %f974, %f346, %f352;
- fma.rn.f32 %f975, %f347, %f351, %f974;
- fma.rn.f32 %f354, %f345, %f353, %f975;
- setp.lt.f32 %p175, %f354, 0f00000000;
- mov.pred %p256, -1;
- mov.f32 %f970, 0f4CBEBC20;
- @%p175 bra BB1_185;
- bra.uni BB1_186;
- BB1_185:
- mov.f32 %f1339, %f970;
- mov.f32 %f1340, %f970;
- bra.uni BB1_188;
- BB1_186:
- mul.f32 %f978, %f352, %f352;
- fma.rn.f32 %f979, %f351, %f351, %f978;
- fma.rn.f32 %f980, %f353, %f353, %f979;
- mul.f32 %f981, %f354, %f354;
- sub.f32 %f355, %f980, %f981;
- ld.global.f32 %f356, [%rd1+24];
- setp.gt.f32 %p177, %f355, %f356;
- mov.f32 %f1339, %f970;
- mov.f32 %f1340, %f970;
- @%p177 bra BB1_188;
- sub.f32 %f982, %f356, %f355;
- sqrt.rn.f32 %f983, %f982;
- sub.f32 %f1340, %f354, %f983;
- add.f32 %f1339, %f354, %f983;
- mov.pred %p256, 0;
- BB1_188:
- mov.u32 %r144, 1;
- @%p256 bra BB1_189;
- setp.lt.f32 %p179, %f1340, 0f00000000;
- selp.f32 %f1344, %f1339, %f1340, %p179;
- setp.geu.f32 %p180, %f1344, 0f4CBEBC20;
- @%p180 bra BB1_189;
- mov.u64 %rd387, %rd1;
- mov.u64 %rd372, %rd158;
- bra.uni BB1_192;
- BB1_189:
- mov.u64 %rd372, %rd387;
- mov.f32 %f1344, %f970;
- BB1_192:
- cvt.u64.u32 %rd107, %r144;
- mul.wide.u32 %rd247, %r144, 80;
- add.s64 %rd248, %rd1, %rd247;
- add.s64 %rd108, %rd248, 8;
- ld.global.f32 %f988, [%rd248+8];
- sub.f32 %f363, %f988, %f348;
- ld.global.f32 %f989, [%rd248+12];
- sub.f32 %f364, %f989, %f349;
- ld.global.f32 %f990, [%rd248+16];
- sub.f32 %f365, %f990, %f350;
- mul.f32 %f991, %f346, %f364;
- fma.rn.f32 %f992, %f347, %f363, %f991;
- fma.rn.f32 %f366, %f345, %f365, %f992;
- setp.lt.f32 %p182, %f366, 0f00000000;
- mov.pred %p257, -1;
- @%p182 bra BB1_193;
- bra.uni BB1_194;
- BB1_193:
- mov.f32 %f1342, %f1346;
- mov.f32 %f1343, %f1346;
- bra.uni BB1_196;
- BB1_194:
- mul.f32 %f995, %f364, %f364;
- fma.rn.f32 %f996, %f363, %f363, %f995;
- fma.rn.f32 %f997, %f365, %f365, %f996;
- mul.f32 %f998, %f366, %f366;
- sub.f32 %f367, %f997, %f998;
- ld.global.f32 %f368, [%rd108+16];
- setp.gt.f32 %p184, %f367, %f368;
- mov.f32 %f1342, %f1346;
- mov.f32 %f1343, %f1346;
- @%p184 bra BB1_196;
- sub.f32 %f999, %f368, %f367;
- sqrt.rn.f32 %f1000, %f999;
- sub.f32 %f1343, %f366, %f1000;
- add.f32 %f1342, %f366, %f1000;
- mov.pred %p257, 0;
- BB1_196:
- @%p257 bra BB1_199;
- setp.lt.f32 %p186, %f1343, 0f00000000;
- selp.f32 %f373, %f1342, %f1343, %p186;
- setp.geu.f32 %p187, %f373, %f1344;
- @%p187 bra BB1_199;
- mul.lo.s64 %rd249, %rd107, 80;
- add.s64 %rd387, %rd1, %rd249;
- add.s64 %rd372, %rd158, %rd249;
- mov.f32 %f1344, %f373;
- BB1_199:
- cvt.u32.u64 %r88, %rd107;
- add.s32 %r144, %r88, 1;
- BB1_200:
- cvt.u64.u32 %rd115, %r144;
- mul.wide.u32 %rd250, %r144, 80;
- add.s64 %rd251, %rd1, %rd250;
- add.s64 %rd116, %rd251, 8;
- ld.global.f32 %f1003, [%rd251+8];
- sub.f32 %f376, %f1003, %f348;
- ld.global.f32 %f1004, [%rd251+12];
- sub.f32 %f377, %f1004, %f349;
- ld.global.f32 %f1005, [%rd251+16];
- sub.f32 %f378, %f1005, %f350;
- mul.f32 %f1006, %f346, %f377;
- fma.rn.f32 %f1007, %f347, %f376, %f1006;
- fma.rn.f32 %f379, %f345, %f378, %f1007;
- setp.lt.f32 %p189, %f379, 0f00000000;
- mov.pred %p258, -1;
- @%p189 bra BB1_201;
- bra.uni BB1_202;
- BB1_201:
- mov.f32 %f1347, %f1346;
- bra.uni BB1_204;
- BB1_202:
- mul.f32 %f1010, %f377, %f377;
- fma.rn.f32 %f1011, %f376, %f376, %f1010;
- fma.rn.f32 %f1012, %f378, %f378, %f1011;
- mul.f32 %f1013, %f379, %f379;
- sub.f32 %f380, %f1012, %f1013;
- ld.global.f32 %f381, [%rd116+16];
- setp.gt.f32 %p191, %f380, %f381;
- mov.f32 %f1347, %f1346;
- @%p191 bra BB1_204;
- sub.f32 %f1014, %f381, %f380;
- sqrt.rn.f32 %f1015, %f1014;
- sub.f32 %f1347, %f379, %f1015;
- add.f32 %f1346, %f379, %f1015;
- mov.pred %p258, 0;
- BB1_204:
- @%p258 bra BB1_207;
- setp.lt.f32 %p193, %f1347, 0f00000000;
- selp.f32 %f386, %f1346, %f1347, %p193;
- setp.geu.f32 %p194, %f386, %f1344;
- @%p194 bra BB1_207;
- mul.lo.s64 %rd252, %rd115, 80;
- add.s64 %rd387, %rd1, %rd252;
- add.s64 %rd372, %rd158, %rd252;
- mov.f32 %f1344, %f386;
- BB1_207:
- cvt.u32.u64 %r89, %rd115;
- add.s32 %r144, %r89, 1;
- mov.f32 %f1350, %f1344;
- BB1_208:
- setp.lt.u32 %p195, %r37, 4;
- @%p195 bra BB1_209;
- bra.uni BB1_210;
- BB1_209:
- mov.f32 %f1344, %f1350;
- bra.uni BB1_239;
- BB1_210:
- cvt.u64.u32 %rd125, %r144;
- mul.wide.u32 %rd253, %r144, 80;
- add.s64 %rd254, %rd1, %rd253;
- add.s64 %rd126, %rd254, 8;
- ld.global.f32 %f1018, [%rd254+8];
- sub.f32 %f391, %f1018, %f348;
- ld.global.f32 %f1019, [%rd254+12];
- sub.f32 %f392, %f1019, %f349;
- ld.global.f32 %f1020, [%rd254+16];
- sub.f32 %f393, %f1020, %f350;
- mul.f32 %f1021, %f346, %f392;
- fma.rn.f32 %f1022, %f347, %f391, %f1021;
- fma.rn.f32 %f394, %f345, %f393, %f1022;
- setp.lt.f32 %p197, %f394, 0f00000000;
- mov.pred %p259, -1;
- mov.f32 %f1352, 0f4CBEBC20;
- @%p197 bra BB1_211;
- bra.uni BB1_212;
- BB1_211:
- mov.f32 %f1353, %f1352;
- bra.uni BB1_214;
- BB1_212:
- mul.f32 %f1025, %f392, %f392;
- fma.rn.f32 %f1026, %f391, %f391, %f1025;
- fma.rn.f32 %f1027, %f393, %f393, %f1026;
- mul.f32 %f1028, %f394, %f394;
- sub.f32 %f395, %f1027, %f1028;
- ld.global.f32 %f396, [%rd126+16];
- setp.gt.f32 %p199, %f395, %f396;
- mov.f32 %f1353, %f1352;
- @%p199 bra BB1_214;
- sub.f32 %f1029, %f396, %f395;
- sqrt.rn.f32 %f1030, %f1029;
- sub.f32 %f1353, %f394, %f1030;
- add.f32 %f1352, %f394, %f1030;
- mov.pred %p259, 0;
- BB1_214:
- @%p259 bra BB1_217;
- setp.lt.f32 %p201, %f1353, 0f00000000;
- selp.f32 %f401, %f1352, %f1353, %p201;
- setp.geu.f32 %p202, %f401, %f1344;
- @%p202 bra BB1_217;
- mul.lo.s64 %rd255, %rd125, 80;
- add.s64 %rd387, %rd1, %rd255;
- add.s64 %rd372, %rd158, %rd255;
- mov.f32 %f1344, %f401;
- BB1_217:
- add.s32 %r90, %r144, 1;
- cvt.u64.u32 %rd131, %r90;
- mul.wide.u32 %rd256, %r90, 80;
- add.s64 %rd257, %rd1, %rd256;
- add.s64 %rd132, %rd257, 8;
- ld.global.f32 %f1033, [%rd257+8];
- sub.f32 %f403, %f1033, %f348;
- ld.global.f32 %f1034, [%rd257+12];
- sub.f32 %f404, %f1034, %f349;
- ld.global.f32 %f1035, [%rd257+16];
- sub.f32 %f405, %f1035, %f350;
- mul.f32 %f1036, %f346, %f404;
- fma.rn.f32 %f1037, %f347, %f403, %f1036;
- fma.rn.f32 %f406, %f345, %f405, %f1037;
- setp.lt.f32 %p204, %f406, 0f00000000;
- mov.pred %p260, -1;
- mov.f32 %f1355, 0f4CBEBC20;
- @%p204 bra BB1_218;
- bra.uni BB1_219;
- BB1_218:
- mov.f32 %f1356, %f1355;
- bra.uni BB1_221;
- BB1_219:
- mul.f32 %f1040, %f404, %f404;
- fma.rn.f32 %f1041, %f403, %f403, %f1040;
- fma.rn.f32 %f1042, %f405, %f405, %f1041;
- mul.f32 %f1043, %f406, %f406;
- sub.f32 %f407, %f1042, %f1043;
- ld.global.f32 %f408, [%rd132+16];
- setp.gt.f32 %p206, %f407, %f408;
- mov.f32 %f1356, %f1355;
- @%p206 bra BB1_221;
- sub.f32 %f1044, %f408, %f407;
- sqrt.rn.f32 %f1045, %f1044;
- sub.f32 %f1356, %f406, %f1045;
- add.f32 %f1355, %f406, %f1045;
- mov.pred %p260, 0;
- BB1_221:
- @%p260 bra BB1_224;
- setp.lt.f32 %p208, %f1356, 0f00000000;
- selp.f32 %f413, %f1355, %f1356, %p208;
- setp.geu.f32 %p209, %f413, %f1344;
- @%p209 bra BB1_224;
- mul.lo.s64 %rd258, %rd131, 80;
- add.s64 %rd387, %rd1, %rd258;
- add.s64 %rd372, %rd158, %rd258;
- mov.f32 %f1344, %f413;
- BB1_224:
- add.s32 %r91, %r144, 2;
- cvt.u64.u32 %rd137, %r91;
- mul.wide.u32 %rd259, %r91, 80;
- add.s64 %rd260, %rd1, %rd259;
- add.s64 %rd138, %rd260, 8;
- ld.global.f32 %f1048, [%rd260+8];
- sub.f32 %f415, %f1048, %f348;
- ld.global.f32 %f1049, [%rd260+12];
- sub.f32 %f416, %f1049, %f349;
- ld.global.f32 %f1050, [%rd260+16];
- sub.f32 %f417, %f1050, %f350;
- mul.f32 %f1051, %f346, %f416;
- fma.rn.f32 %f1052, %f347, %f415, %f1051;
- fma.rn.f32 %f418, %f345, %f417, %f1052;
- setp.lt.f32 %p211, %f418, 0f00000000;
- mov.pred %p261, -1;
- mov.f32 %f1358, 0f4CBEBC20;
- @%p211 bra BB1_225;
- bra.uni BB1_226;
- BB1_225:
- mov.f32 %f1359, %f1358;
- bra.uni BB1_228;
- BB1_226:
- mul.f32 %f1055, %f416, %f416;
- fma.rn.f32 %f1056, %f415, %f415, %f1055;
- fma.rn.f32 %f1057, %f417, %f417, %f1056;
- mul.f32 %f1058, %f418, %f418;
- sub.f32 %f419, %f1057, %f1058;
- ld.global.f32 %f420, [%rd138+16];
- setp.gt.f32 %p213, %f419, %f420;
- mov.f32 %f1359, %f1358;
- @%p213 bra BB1_228;
- sub.f32 %f1059, %f420, %f419;
- sqrt.rn.f32 %f1060, %f1059;
- sub.f32 %f1359, %f418, %f1060;
- add.f32 %f1358, %f418, %f1060;
- mov.pred %p261, 0;
- BB1_228:
- @%p261 bra BB1_231;
- setp.lt.f32 %p215, %f1359, 0f00000000;
- selp.f32 %f425, %f1358, %f1359, %p215;
- setp.geu.f32 %p216, %f425, %f1344;
- @%p216 bra BB1_231;
- mul.lo.s64 %rd261, %rd137, 80;
- add.s64 %rd387, %rd1, %rd261;
- add.s64 %rd372, %rd158, %rd261;
- mov.f32 %f1344, %f425;
- BB1_231:
- add.s32 %r92, %r144, 3;
- cvt.u64.u32 %rd143, %r92;
- mul.wide.u32 %rd262, %r92, 80;
- add.s64 %rd263, %rd1, %rd262;
- add.s64 %rd144, %rd263, 8;
- ld.global.f32 %f1063, [%rd263+8];
- sub.f32 %f427, %f1063, %f348;
- ld.global.f32 %f1064, [%rd263+12];
- sub.f32 %f428, %f1064, %f349;
- ld.global.f32 %f1065, [%rd263+16];
- sub.f32 %f429, %f1065, %f350;
- mul.f32 %f1066, %f346, %f428;
- fma.rn.f32 %f1067, %f347, %f427, %f1066;
- fma.rn.f32 %f430, %f345, %f429, %f1067;
- setp.lt.f32 %p218, %f430, 0f00000000;
- mov.pred %p262, -1;
- mov.f32 %f1361, 0f4CBEBC20;
- @%p218 bra BB1_232;
- bra.uni BB1_233;
- BB1_232:
- mov.f32 %f1362, %f1361;
- bra.uni BB1_235;
- BB1_233:
- mul.f32 %f1070, %f428, %f428;
- fma.rn.f32 %f1071, %f427, %f427, %f1070;
- fma.rn.f32 %f1072, %f429, %f429, %f1071;
- mul.f32 %f1073, %f430, %f430;
- sub.f32 %f431, %f1072, %f1073;
- ld.global.f32 %f432, [%rd144+16];
- setp.gt.f32 %p220, %f431, %f432;
- mov.f32 %f1362, %f1361;
- @%p220 bra BB1_235;
- sub.f32 %f1074, %f432, %f431;
- sqrt.rn.f32 %f1075, %f1074;
- sub.f32 %f1362, %f430, %f1075;
- add.f32 %f1361, %f430, %f1075;
- mov.pred %p262, 0;
- BB1_235:
- @%p262 bra BB1_238;
- setp.lt.f32 %p222, %f1362, 0f00000000;
- selp.f32 %f437, %f1361, %f1362, %p222;
- setp.geu.f32 %p223, %f437, %f1344;
- @%p223 bra BB1_238;
- mul.lo.s64 %rd264, %rd143, 80;
- add.s64 %rd387, %rd1, %rd264;
- add.s64 %rd372, %rd158, %rd264;
- mov.f32 %f1344, %f437;
- BB1_238:
- cvt.u32.u64 %r93, %rd125;
- add.s32 %r144, %r93, 4;
- setp.lt.u32 %p224, %r144, %r37;
- @%p224 bra BB1_210;
- BB1_239:
- setp.eq.s64 %p225, %rd372, 0;
- mov.f32 %f1390, 0f40000000;
- @%p225 bra BB1_174;
- fma.rn.f32 %f440, %f347, %f1344, %f348;
- fma.rn.f32 %f441, %f346, %f1344, %f349;
- fma.rn.f32 %f442, %f345, %f1344, %f350;
- add.s64 %rd151, %rd387, 8;
- ld.global.f32 %f1079, [%rd387+8];
- sub.f32 %f456, %f440, %f1079;
- ld.global.f32 %f1080, [%rd387+12];
- sub.f32 %f457, %f441, %f1080;
- ld.global.f32 %f1081, [%rd387+16];
- sub.f32 %f458, %f442, %f1081;
- mul.f32 %f1082, %f457, %f457;
- fma.rn.f32 %f1083, %f456, %f456, %f1082;
- fma.rn.f32 %f446, %f458, %f458, %f1083;
- setp.leu.f32 %p226, %f446, 0f00000000;
- @%p226 bra BB1_242;
- sqrt.rn.f32 %f1084, %f446;
- rcp.rn.f32 %f1085, %f1084;
- mul.f32 %f456, %f456, %f1085;
- mul.f32 %f457, %f457, %f1085;
- mul.f32 %f458, %f458, %f1085;
- BB1_242:
- mul.f32 %f1086, %f346, %f457;
- fma.rn.f32 %f1087, %f347, %f456, %f1086;
- fma.rn.f32 %f1088, %f345, %f458, %f1087;
- mov.u16 %rs12, 0;
- setp.leu.f32 %p227, %f1088, 0f00000000;
- @%p227 bra BB1_244;
- neg.f32 %f456, %f456;
- neg.f32 %f457, %f457;
- neg.f32 %f458, %f458;
- mov.u16 %rs12, 1;
- BB1_244:
- ld.global.f32 %f1089, [%rd151+44];
- setp.gt.f32 %p228, %f1089, 0f00000000;
- @%p228 bra BB1_260;
- ld.global.f32 %f1090, [%rd151+48];
- setp.gt.f32 %p229, %f1090, 0f00000000;
- @%p229 bra BB1_260;
- bra.uni BB1_246;
- BB1_260:
- mov.u32 %r125, 3;
- mul.f32 %f1133, %f346, %f457;
- fma.rn.f32 %f1134, %f347, %f456, %f1133;
- fma.rn.f32 %f487, %f345, %f458, %f1134;
- add.f32 %f1135, %f487, 0f3F800000;
- // Callseq Start 7
- {
- .reg .b32 temp_param_reg;
- // <end>}
- .param .b32 param0;
- st.param.f32 [param0+0], %f1135;
- .param .b32 param1;
- st.param.b32 [param1+0], %r125;
- .param .b32 retval0;
- call.uni (retval0),
- _Z5powiffi,
- (
- param0,
- param1
- );
- ld.param.f32 %f1136, [retval0+0];
- //{
- }// Callseq End 7
- fma.rn.f32 %f488, %f1136, 0f3F666666, 0f3DCCCCCD;
- add.f32 %f1137, %f456, %f456;
- mul.f32 %f1138, %f1137, %f487;
- add.f32 %f1139, %f457, %f457;
- mul.f32 %f1140, %f1139, %f487;
- add.f32 %f1141, %f458, %f458;
- mul.f32 %f1142, %f1141, %f487;
- sub.f32 %f489, %f347, %f1138;
- sub.f32 %f490, %f346, %f1140;
- sub.f32 %f491, %f345, %f1142;
- add.u64 %rd269, %SP, 12;
- add.u64 %rd155, %SPL, 12;
- st.local.f32 [%rd155], %f489;
- st.local.f32 [%rd155+4], %f490;
- st.local.f32 [%rd155+8], %f491;
- mul.f32 %f1143, %f490, %f490;
- fma.rn.f32 %f1144, %f489, %f489, %f1143;
- fma.rn.f32 %f492, %f491, %f491, %f1144;
- setp.leu.f32 %p238, %f492, 0f00000000;
- @%p238 bra BB1_262;
- sqrt.rn.f32 %f1145, %f492;
- rcp.rn.f32 %f1146, %f1145;
- mul.f32 %f1147, %f489, %f1146;
- st.local.f32 [%rd155], %f1147;
- mul.f32 %f1148, %f490, %f1146;
- st.local.f32 [%rd155+4], %f1148;
- mul.f32 %f1149, %f491, %f1146;
- st.local.f32 [%rd155+8], %f1149;
- BB1_262:
- mul.f32 %f493, %f456, 0f38D1B717;
- add.f32 %f1153, %f440, %f493;
- mul.f32 %f494, %f457, 0f38D1B717;
- add.f32 %f1154, %f441, %f494;
- mul.f32 %f495, %f458, 0f38D1B717;
- add.f32 %f1155, %f442, %f495;
- add.u64 %rd270, %SP, 52;
- add.u64 %rd271, %SPL, 52;
- mov.u32 %r98, 2;
- st.local.u32 [%rd271], %r98;
- // Callseq Start 8
- {
- .reg .b32 temp_param_reg;
- // <end>}
- .param .b32 param0;
- st.param.f32 [param0+0], %f1153;
- .param .b32 param1;
- st.param.f32 [param1+0], %f1154;
- .param .b32 param2;
- st.param.f32 [param2+0], %f1155;
- .param .b64 param3;
- st.param.b64 [param3+0], %rd269;
- .param .b64 param4;
- st.param.b64 [param4+0], %rd158;
- .param .b32 param5;
- st.param.b32 [param5+0], %r37;
- .param .b64 param6;
- st.param.b64 [param6+0], %rd270;
- .param .align 4 .b8 retval0[12];
- call.uni (retval0),
- _Z5traceRK4Vec3IfES2_PK6SpherejRKi,
- (
- param0,
- param1,
- param2,
- param3,
- param4,
- param5,
- param6
- );
- ld.param.f32 %f496, [retval0+0];
- ld.param.f32 %f497, [retval0+4];
- ld.param.f32 %f498, [retval0+8];
- //{
- }// Callseq End 8
- ld.global.f32 %f1381, [%rd151+44];
- mov.f32 %f1382, 0f00000000;
- setp.eq.f32 %p239, %f1381, 0f00000000;
- mov.f32 %f1383, %f1382;
- mov.f32 %f1384, %f1382;
- @%p239 bra BB1_266;
- setp.eq.s16 %p240, %rs12, 0;
- selp.f32 %f1156, 0f3F68BA2E, 0f3F8CCCCD, %p240;
- mul.f32 %f1157, %f1156, %f1156;
- mul.f32 %f1158, %f487, %f487;
- mov.f32 %f1159, 0f3F800000;
- sub.f32 %f1160, %f1159, %f1158;
- mul.f32 %f1161, %f1157, %f1160;
- sub.f32 %f1162, %f1159, %f1161;
- mul.f32 %f1163, %f487, %f1156;
- neg.f32 %f1164, %f1163;
- sqrt.rn.f32 %f1165, %f1162;
- sub.f32 %f1166, %f1164, %f1165;
- mul.f32 %f1167, %f456, %f1166;
- mul.f32 %f1168, %f457, %f1166;
- mul.f32 %f1169, %f458, %f1166;
- fma.rn.f32 %f500, %f347, %f1156, %f1167;
- fma.rn.f32 %f501, %f346, %f1156, %f1168;
- fma.rn.f32 %f502, %f345, %f1156, %f1169;
- add.u64 %rd273, %SP, 0;
- add.u64 %rd156, %SPL, 0;
- st.local.f32 [%rd156], %f500;
- st.local.f32 [%rd156+4], %f501;
- st.local.f32 [%rd156+8], %f502;
- mul.f32 %f1170, %f501, %f501;
- fma.rn.f32 %f1171, %f500, %f500, %f1170;
- fma.rn.f32 %f503, %f502, %f502, %f1171;
- setp.leu.f32 %p241, %f503, 0f00000000;
- @%p241 bra BB1_265;
- sqrt.rn.f32 %f1172, %f503;
- rcp.rn.f32 %f1173, %f1172;
- mul.f32 %f1174, %f500, %f1173;
- st.local.f32 [%rd156], %f1174;
- mul.f32 %f1175, %f501, %f1173;
- st.local.f32 [%rd156+4], %f1175;
- mul.f32 %f1176, %f502, %f1173;
- st.local.f32 [%rd156+8], %f1176;
- BB1_265:
- sub.f32 %f1177, %f440, %f493;
- add.u64 %rd274, %SP, 48;
- add.u64 %rd275, %SPL, 48;
- st.local.u32 [%rd275], %r98;
- sub.f32 %f1178, %f442, %f495;
- sub.f32 %f1179, %f441, %f494;
- // Callseq Start 9
- {
- .reg .b32 temp_param_reg;
- // <end>}
- .param .b32 param0;
- st.param.f32 [param0+0], %f1177;
- .param .b32 param1;
- st.param.f32 [param1+0], %f1179;
- .param .b32 param2;
- st.param.f32 [param2+0], %f1178;
- .param .b64 param3;
- st.param.b64 [param3+0], %rd273;
- .param .b64 param4;
- st.param.b64 [param4+0], %rd158;
- .param .b32 param5;
- st.param.b32 [param5+0], %r37;
- .param .b64 param6;
- st.param.b64 [param6+0], %rd274;
- .param .align 4 .b8 retval0[12];
- call.uni (retval0),
- _Z5traceRK4Vec3IfES2_PK6SpherejRKi,
- (
- param0,
- param1,
- param2,
- param3,
- param4,
- param5,
- param6
- );
- ld.param.f32 %f1382, [retval0+0];
- ld.param.f32 %f1383, [retval0+4];
- ld.param.f32 %f1384, [retval0+8];
- //{
- }// Callseq End 9
- ld.global.f32 %f1381, [%rd151+44];
- BB1_266:
- mov.f32 %f1180, 0f3F800000;
- sub.f32 %f1181, %f1180, %f488;
- mul.f32 %f1182, %f1181, %f1382;
- mul.f32 %f1183, %f1181, %f1383;
- mul.f32 %f1184, %f1181, %f1384;
- mul.f32 %f1185, %f1381, %f1182;
- mul.f32 %f1186, %f1381, %f1183;
- mul.f32 %f1187, %f1381, %f1184;
- fma.rn.f32 %f1188, %f488, %f496, %f1185;
- fma.rn.f32 %f1189, %f488, %f497, %f1186;
- fma.rn.f32 %f1190, %f488, %f498, %f1187;
- ld.global.f32 %f1191, [%rd151+20];
- mul.f32 %f1386, %f1191, %f1188;
- ld.global.f32 %f1192, [%rd151+24];
- mul.f32 %f1387, %f1189, %f1192;
- ld.global.f32 %f1193, [%rd151+28];
- mul.f32 %f1388, %f1190, %f1193;
- ld.global.f32 %f1389, [%rd49+44];
- bra.uni BB1_267;
- BB1_174:
- mov.f32 %f1391, %f1390;
- mov.f32 %f1392, %f1390;
- bra.uni BB1_268;
- BB1_246:
- mov.f32 %f1386, 0f00000000;
- @%p26 bra BB1_247;
- fma.rn.f32 %f459, %f456, 0f38D1B717, %f440;
- fma.rn.f32 %f460, %f457, 0f38D1B717, %f441;
- fma.rn.f32 %f461, %f458, 0f38D1B717, %f442;
- mov.f32 %f1388, 0f00000000;
- mov.u32 %r146, 0;
- mov.f32 %f1387, %f1388;
- mov.f32 %f1386, %f1388;
- BB1_249:
- cvt.u64.u32 %rd152, %r146;
- mul.wide.u32 %rd265, %r146, 80;
- add.s64 %rd266, %rd1, %rd265;
- add.s64 %rd153, %rd266, 40;
- ld.global.f32 %f465, [%rd266+40];
- setp.leu.f32 %p231, %f465, 0f00000000;
- @%p231 bra BB1_259;
- ld.global.f32 %f1097, [%rd153+-32];
- sub.f32 %f1376, %f1097, %f440;
- ld.global.f32 %f1098, [%rd153+-28];
- sub.f32 %f1375, %f1098, %f441;
- ld.global.f32 %f1099, [%rd153+-24];
- sub.f32 %f1374, %f1099, %f442;
- mul.f32 %f1100, %f1375, %f1375;
- fma.rn.f32 %f1101, %f1376, %f1376, %f1100;
- fma.rn.f32 %f469, %f1374, %f1374, %f1101;
- setp.leu.f32 %p232, %f469, 0f00000000;
- @%p232 bra BB1_252;
- sqrt.rn.f32 %f1102, %f469;
- rcp.rn.f32 %f1103, %f1102;
- mul.f32 %f1376, %f1376, %f1103;
- mul.f32 %f1375, %f1375, %f1103;
- mul.f32 %f1374, %f1374, %f1103;
- BB1_252:
- mov.u32 %r147, 0;
- BB1_253:
- cvt.u32.u64 %r96, %rd152;
- setp.eq.s32 %p233, %r96, %r147;
- @%p233 bra BB1_256;
- mul.wide.u32 %rd267, %r147, 80;
- add.s64 %rd268, %rd1, %rd267;
- add.s64 %rd154, %rd268, 8;
- ld.global.f32 %f1104, [%rd268+8];
- sub.f32 %f476, %f1104, %f459;
- ld.global.f32 %f1105, [%rd268+12];
- sub.f32 %f477, %f1105, %f460;
- ld.global.f32 %f1106, [%rd268+16];
- sub.f32 %f478, %f1106, %f461;
- mul.f32 %f1107, %f1375, %f477;
- fma.rn.f32 %f1108, %f1376, %f476, %f1107;
- fma.rn.f32 %f479, %f1374, %f478, %f1108;
- setp.lt.f32 %p234, %f479, 0f00000000;
- @%p234 bra BB1_256;
- mul.f32 %f1110, %f477, %f477;
- fma.rn.f32 %f1111, %f476, %f476, %f1110;
- fma.rn.f32 %f1112, %f478, %f478, %f1111;
- mul.f32 %f1113, %f479, %f479;
- sub.f32 %f1114, %f1112, %f1113;
- ld.global.f32 %f1115, [%rd154+16];
- mov.f32 %f1377, 0f00000000;
- setp.leu.f32 %p235, %f1114, %f1115;
- @%p235 bra BB1_258;
- BB1_256:
- add.s32 %r147, %r147, 1;
- setp.lt.u32 %p236, %r147, %r37;
- @%p236 bra BB1_253;
- mov.f32 %f1377, %f949;
- BB1_258:
- ld.global.f32 %f1117, [%rd151+20];
- mul.f32 %f1118, %f1377, %f1117;
- ld.global.f32 %f1119, [%rd151+24];
- mul.f32 %f1120, %f1377, %f1119;
- ld.global.f32 %f1121, [%rd151+28];
- mul.f32 %f1122, %f1377, %f1121;
- mul.f32 %f1123, %f456, %f1376;
- fma.rn.f32 %f1124, %f457, %f1375, %f1123;
- fma.rn.f32 %f1125, %f458, %f1374, %f1124;
- mov.f32 %f1126, 0f00000000;
- max.f32 %f1127, %f1126, %f1125;
- mul.f32 %f1128, %f1118, %f1127;
- mul.f32 %f1129, %f1120, %f1127;
- mul.f32 %f1130, %f1122, %f1127;
- ld.global.f32 %f1131, [%rd153+4];
- ld.global.f32 %f1132, [%rd153+8];
- fma.rn.f32 %f1386, %f465, %f1128, %f1386;
- fma.rn.f32 %f1387, %f1129, %f1131, %f1387;
- fma.rn.f32 %f1388, %f1130, %f1132, %f1388;
- BB1_259:
- add.s32 %r146, %r146, 1;
- setp.lt.u32 %p237, %r146, %r37;
- @%p237 bra BB1_249;
- bra.uni BB1_267;
- BB1_247:
- mov.f32 %f1387, %f1386;
- mov.f32 %f1388, %f1386;
- BB1_267:
- ld.global.f32 %f1194, [%rd151+32];
- add.f32 %f1390, %f1386, %f1194;
- ld.global.f32 %f1195, [%rd151+36];
- add.f32 %f1391, %f1387, %f1195;
- ld.global.f32 %f1196, [%rd151+40];
- add.f32 %f1392, %f1388, %f1196;
- BB1_268:
- mov.f32 %f1197, 0f3F800000;
- sub.f32 %f1198, %f1197, %f147;
- mul.f32 %f1199, %f1198, %f1390;
- mul.f32 %f1200, %f1198, %f1391;
- mul.f32 %f1201, %f1198, %f1392;
- mul.f32 %f1202, %f1199, %f1389;
- mul.f32 %f1203, %f1200, %f1389;
- mul.f32 %f1204, %f1201, %f1389;
- fma.rn.f32 %f1205, %f147, %f1333, %f1202;
- fma.rn.f32 %f1206, %f147, %f1334, %f1203;
- fma.rn.f32 %f1207, %f147, %f1335, %f1204;
- ld.global.f32 %f1208, [%rd49+20];
- mul.f32 %f1393, %f1208, %f1205;
- ld.global.f32 %f1209, [%rd49+24];
- mul.f32 %f1394, %f1206, %f1209;
- ld.global.f32 %f1210, [%rd49+28];
- mul.f32 %f1395, %f1207, %f1210;
- BB1_269:
- ld.global.f32 %f1211, [%rd49+32];
- add.f32 %f1396, %f1393, %f1211;
- ld.global.f32 %f1212, [%rd49+36];
- add.f32 %f1397, %f1394, %f1212;
- ld.global.f32 %f1213, [%rd49+40];
- add.f32 %f1398, %f1395, %f1213;
- BB1_270:
- ld.param.u64 %rd285, [raytrace_kernel_param_2];
- ld.param.u32 %r117, [raytrace_kernel_param_0];
- mov.u32 %r116, %tid.y;
- mov.u32 %r115, %ctaid.y;
- mov.u32 %r114, %ntid.y;
- mad.lo.s32 %r113, %r114, %r115, %r116;
- mov.u32 %r112, %tid.x;
- mov.u32 %r111, %ctaid.x;
- mov.u32 %r110, %ntid.x;
- mad.lo.s32 %r109, %r110, %r111, %r112;
- mad.lo.s32 %r108, %r113, %r117, %r109;
- cvta.to.global.u64 %rd277, %rd285;
- mul.wide.u32 %rd278, %r108, 12;
- add.s64 %rd279, %rd277, %rd278;
- st.global.f32 [%rd279], %f1396;
- st.global.f32 [%rd279+4], %f1397;
- st.global.f32 [%rd279+8], %f1398;
- BB1_271:
- ret;
- }
Advertisement
Add Comment
Please, Sign In to add comment
Advertisement