// // Generated by NVIDIA NVVM Compiler // Compiler built on Sat Apr 7 09:27:48 2012 (1333783668) // Driver 295.41 // .version 3.0 .target sm_13, texmode_independent .address_size 32 entry butterfC( .param .u32 .ptr .global .align 8 butterfC_param_0, .param .u32 .ptr .global .align 8 butterfC_param_1, .param .u32 .ptr .global .align 8 butterfC_param_2, .param .u32 butterfC_param_3, .param .u32 butterfC_param_4, .param .u32 butterfC_param_5, .param .u32 butterfC_param_6, .param .u32 butterfC_param_7, .param .u32 butterfC_param_8 ) { .reg .f32 %f<285>; .reg .f64 %fd<244>; .reg .pred %p<58>; .reg .s32 %r<137>; ld.param.u32 %r1, [butterfC_param_6]; // inline asm mov.u32 %r39, %envreg3; // inline asm // inline asm mov.u32 %r40, %ntid.x; // inline asm // inline asm mov.u32 %r41, %ctaid.x; // inline asm // inline asm mov.u32 %r42, %tid.x; // inline asm add.s32 %r51, %r42, %r39; mad.lo.s32 %r3, %r41, %r40, %r51; // inline asm mov.u32 %r43, %envreg4; // inline asm // inline asm mov.u32 %r44, %ntid.y; // inline asm // inline asm mov.u32 %r45, %ctaid.y; // inline asm // inline asm mov.u32 %r46, %tid.y; // inline asm add.s32 %r52, %r46, %r43; mad.lo.s32 %r4, %r45, %r44, %r52; // inline asm mov.u32 %r47, %envreg5; // inline asm // inline asm mov.u32 %r48, %ntid.z; // inline asm // inline asm mov.u32 %r49, %ctaid.z; // inline asm // inline asm mov.u32 %r50, %tid.z; // inline asm add.s32 %r53, %r50, %r47; mad.lo.s32 %r5, %r49, %r48, %r53; mov.f32 %f44, 0f40800000; // inline asm abs.f32 %f43, %f44; // inline asm cvt.rn.f32.u32 %f275, %r1; setp.eq.f32 %p5, %f275, 0f00000000; @%p5 bra BB5_22; setp.nan.f32 %p6, %f275, %f275; @%p6 bra BB5_21; mov.f32 %f3, 0f7F800000; setp.eq.f32 %p7, %f275, 0f7F800000; setp.eq.f32 %p8, %f275, 0fFF800000; or.pred %p9, %p7, %p8; @%p9 bra BB5_18; mov.f32 %f49, 0f3F000000; mul.rn.f32 %f46, %f49, %f275; // inline asm cvt.rmi.f32.f32 %f45, %f46; // inline asm // inline asm cvt.rzi.f32.f32 %f47, %f275; // inline asm setp.eq.f32 %p12, %f43, 0f00000000; @%p12 bra BB5_15; // inline asm abs.f32 %f53, %f44; // inline asm mov.b32 %r6, %f53; shr.u32 %r54, %r6, 23; and.b32 %r55, %r54, 255; add.s32 %r130, %r55, -127; setp.eq.s32 %p13, %r55, 0; mov.f32 %f273, %f53; @%p13 bra BB5_5; bra.uni BB5_6; BB5_5: and.b32 %r56, %r6, -2139095041; or.b32 %r57, %r56, 1065353216; mov.b32 %f55, %r57; add.f32 %f56, %f55, 0fBF800000; mov.b32 %r58, %f56; shr.u32 %r59, %r58, 23; and.b32 %r60, %r59, 255; add.s32 %r130, %r60, -253; and.b32 %r61, %r58, -2139095041; or.b32 %r62, %r61, 1065353216; mov.b32 %f273, %r62; BB5_6: mov.b32 %r63, %f273; and.b32 %r64, %r63, -2139095041; or.b32 %r65, %r64, 1065353216; mov.b32 %f274, %r65; setp.gt.f32 %p14, %f274, 0f3FB504F3; @%p14 bra BB5_7; bra.uni BB5_8; BB5_7: mul.rn.f32 %f274, %f274, %f49; add.s32 %r130, %r130, 1; BB5_8: add.f32 %f66, %f274, 0f3F800000; rcp.approx.f32 %f60, %f66; add.f32 %f59, %f274, 0fBF800000; // inline asm mul.rz.f32 %f58, %f59, %f60; // inline asm mov.f32 %f67, 0f40000000; mul.rn.f32 %f68, %f67, %f58; mul.rn.f32 %f69, %f68, %f68; mov.f32 %f70, 0f3B18F0FE; mul.rn.f32 %f71, %f70, %f69; add.f32 %f72, %f71, 0f3C4CAF63; mul.rn.f32 %f73, %f72, %f69; add.f32 %f74, %f73, 0f3DAAAABD; mul.rn.f32 %f75, %f74, %f69; mul.rn.f32 %f63, %f75, %f68; mov.b32 %r66, %f68; and.b32 %r67, %r66, -4096; mov.b32 %f76, %r67; mov.b32 %r68, %f59; and.b32 %r69, %r68, -4096; mov.b32 %f77, %r69; sub.f32 %f78, %f59, %f76; mul.rn.f32 %f79, %f67, %f78; sub.f32 %f80, %f59, %f77; mul.rn.f32 %f81, %f76, %f77; sub.f32 %f82, %f79, %f81; mul.rn.f32 %f83, %f76, %f80; sub.f32 %f84, %f82, %f83; mul.rn.f32 %f85, %f60, %f84; add.f32 %f86, %f76, %f85; sub.f32 %f87, %f86, %f76; sub.f32 %f88, %f85, %f87; add.f32 %f89, %f86, %f63; sub.f32 %f62, %f86, %f89; // inline asm add.rz.f32 %f61, %f62, %f63; // inline asm add.f32 %f90, %f61, %f88; add.f32 %f91, %f89, %f90; sub.f32 %f92, %f89, %f91; add.f32 %f93, %f92, %f90; cvt.rn.f32.s32 %f94, %r130; mov.f32 %f95, 0f3F317200; mul.rn.f32 %f96, %f94, %f95; mov.f32 %f97, 0f35BFBE8E; mul.rn.f32 %f98, %f94, %f97; add.f32 %f99, %f96, %f91; sub.f32 %f100, %f96, %f99; add.f32 %f101, %f100, %f91; add.f32 %f102, %f101, %f93; add.f32 %f103, %f102, %f98; add.f32 %f10, %f99, %f103; sub.f32 %f104, %f99, %f10; add.f32 %f11, %f104, %f103; // inline asm abs.f32 %f64, %f275; // inline asm setp.gt.f32 %p15, %f64, 0f77F684DF; @%p15 bra BB5_9; bra.uni BB5_10; BB5_9: mov.f32 %f105, 0f39000000; mul.rn.f32 %f275, %f275, %f105; BB5_10: mov.f32 %f106, 0f45800800; mul.rn.f32 %f107, %f10, %f106; sub.f32 %f108, %f10, %f107; add.f32 %f109, %f108, %f107; sub.f32 %f110, %f10, %f109; mul.rn.f32 %f111, %f275, %f106; sub.f32 %f112, %f275, %f111; add.f32 %f113, %f112, %f111; sub.f32 %f114, %f275, %f113; mul.rn.f32 %f115, %f109, %f113; mul.rn.f32 %f116, %f10, %f275; sub.f32 %f117, %f115, %f116; mul.rn.f32 %f118, %f109, %f114; add.f32 %f119, %f117, %f118; mul.rn.f32 %f120, %f110, %f113; add.f32 %f121, %f119, %f120; mul.rn.f32 %f122, %f110, %f114; add.f32 %f123, %f121, %f122; mul.rn.f32 %f124, %f11, %f275; add.f32 %f125, %f124, %f123; add.f32 %f126, %f116, %f125; sub.f32 %f127, %f116, %f126; add.f32 %f14, %f127, %f125; mov.f32 %f283, %f14; mov.f32 %f284, %f126; mov.b32 %r12, %f126; setp.eq.s32 %p16, %r12, 1118925336; @%p16 bra BB5_11; bra.uni BB5_12; BB5_11: add.s32 %r70, %r12, -1; mov.b32 %f128, %r70; add.f32 %f129, %f14, 0f37000000; mov.f32 %f283, %f129; mov.f32 %f284, %f128; BB5_12: mov.f32 %f137, 0f3FB8AA3B; mul.rn.f32 %f131, %f284, %f137; // inline asm cvt.rzi.f32.f32 %f130, %f131; // inline asm mul.rn.f32 %f139, %f130, %f95; sub.f32 %f140, %f284, %f139; mul.rn.f32 %f142, %f130, %f97; sub.f32 %f143, %f140, %f142; mul.rn.f32 %f133, %f143, %f137; // inline asm ex2.approx.f32 %f132, %f133; // inline asm add.f32 %f135, %f130, 0f00000000; // inline asm ex2.approx.f32 %f134, %f135; // inline asm mul.rn.f32 %f144, %f132, %f134; setp.lt.f32 %p17, %f284, 0fC2D20000; selp.f32 %f145, 0f00000000, %f144, %p17; setp.gt.f32 %p18, %f284, 0f42D20000; selp.f32 %f15, %f3, %f145, %p18; setp.neu.f32 %p19, %f15, %f3; @%p19 bra BB5_14; mov.f32 %f276, %f15; bra.uni BB5_23; BB5_14: // inline asm mad.f32 %f146, %f15, %f283, %f15; // inline asm mov.f32 %f276, %f146; bra.uni BB5_23; BB5_15: setp.lt.f32 %p20, %f275, 0f00000000; @%p20 bra BB5_17; mov.f32 %f276, 0f00000000; bra.uni BB5_23; BB5_17: mov.f32 %f276, %f3; bra.uni BB5_23; BB5_18: setp.lt.f32 %p21, %f43, 0f3F800000; mov.b32 %r71, %f275; setp.lt.s32 %p2, %r71, 0; @%p21 bra BB5_20; selp.f32 %f276, 0f00000000, %f3, %p2; bra.uni BB5_23; BB5_20: selp.f32 %f276, %f3, 0f00000000, %p2; bra.uni BB5_23; BB5_21: add.f32 %f276, %f275, 0f40800000; bra.uni BB5_23; BB5_22: mov.f32 %f276, 0f3F800000; BB5_23: cvt.rzi.s32.f32 %r13, %f276; ld.param.u32 %r126, [butterfC_param_6]; add.s32 %r72, %r126, -1; cvt.rn.f32.u32 %f279, %r72; // inline asm abs.f32 %f152, %f44; // inline asm setp.eq.f32 %p22, %f279, 0f00000000; @%p22 bra BB5_45; setp.nan.f32 %p23, %f279, %f279; @%p23 bra BB5_44; mov.f32 %f24, 0f7F800000; setp.eq.f32 %p24, %f279, 0f7F800000; setp.eq.f32 %p25, %f279, 0fFF800000; or.pred %p26, %p24, %p25; @%p26 bra BB5_41; mov.f32 %f158, 0f3F000000; mul.rn.f32 %f155, %f158, %f279; // inline asm cvt.rmi.f32.f32 %f154, %f155; // inline asm // inline asm cvt.rzi.f32.f32 %f156, %f279; // inline asm setp.eq.f32 %p29, %f152, 0f00000000; @%p29 bra BB5_38; // inline asm abs.f32 %f162, %f44; // inline asm mov.b32 %r14, %f162; shr.u32 %r73, %r14, 23; and.b32 %r74, %r73, 255; add.s32 %r131, %r74, -127; setp.eq.s32 %p30, %r74, 0; mov.f32 %f277, %f162; @%p30 bra BB5_28; bra.uni BB5_29; BB5_28: and.b32 %r75, %r14, -2139095041; or.b32 %r76, %r75, 1065353216; mov.b32 %f164, %r76; add.f32 %f165, %f164, 0fBF800000; mov.b32 %r77, %f165; shr.u32 %r78, %r77, 23; and.b32 %r79, %r78, 255; add.s32 %r131, %r79, -253; and.b32 %r80, %r77, -2139095041; or.b32 %r81, %r80, 1065353216; mov.b32 %f277, %r81; BB5_29: mov.b32 %r82, %f277; and.b32 %r83, %r82, -2139095041; or.b32 %r84, %r83, 1065353216; mov.b32 %f278, %r84; setp.gt.f32 %p31, %f278, 0f3FB504F3; @%p31 bra BB5_30; bra.uni BB5_31; BB5_30: mul.rn.f32 %f278, %f278, %f158; add.s32 %r131, %r131, 1; BB5_31: add.f32 %f175, %f278, 0f3F800000; rcp.approx.f32 %f169, %f175; add.f32 %f168, %f278, 0fBF800000; // inline asm mul.rz.f32 %f167, %f168, %f169; // inline asm mov.f32 %f176, 0f40000000; mul.rn.f32 %f177, %f176, %f167; mul.rn.f32 %f178, %f177, %f177; mov.f32 %f179, 0f3B18F0FE; mul.rn.f32 %f180, %f179, %f178; add.f32 %f181, %f180, 0f3C4CAF63; mul.rn.f32 %f182, %f181, %f178; add.f32 %f183, %f182, 0f3DAAAABD; mul.rn.f32 %f184, %f183, %f178; mul.rn.f32 %f172, %f184, %f177; mov.b32 %r85, %f177; and.b32 %r86, %r85, -4096; mov.b32 %f185, %r86; mov.b32 %r87, %f168; and.b32 %r88, %r87, -4096; mov.b32 %f186, %r88; sub.f32 %f187, %f168, %f185; mul.rn.f32 %f188, %f176, %f187; sub.f32 %f189, %f168, %f186; mul.rn.f32 %f190, %f185, %f186; sub.f32 %f191, %f188, %f190; mul.rn.f32 %f192, %f185, %f189; sub.f32 %f193, %f191, %f192; mul.rn.f32 %f194, %f169, %f193; add.f32 %f195, %f185, %f194; sub.f32 %f196, %f195, %f185; sub.f32 %f197, %f194, %f196; add.f32 %f198, %f195, %f172; sub.f32 %f171, %f195, %f198; // inline asm add.rz.f32 %f170, %f171, %f172; // inline asm add.f32 %f199, %f170, %f197; add.f32 %f200, %f198, %f199; sub.f32 %f201, %f198, %f200; add.f32 %f202, %f201, %f199; cvt.rn.f32.s32 %f203, %r131; mov.f32 %f204, 0f3F317200; mul.rn.f32 %f205, %f203, %f204; mov.f32 %f206, 0f35BFBE8E; mul.rn.f32 %f207, %f203, %f206; add.f32 %f208, %f205, %f200; sub.f32 %f209, %f205, %f208; add.f32 %f210, %f209, %f200; add.f32 %f211, %f210, %f202; add.f32 %f212, %f211, %f207; add.f32 %f31, %f208, %f212; sub.f32 %f213, %f208, %f31; add.f32 %f32, %f213, %f212; // inline asm abs.f32 %f173, %f279; // inline asm setp.gt.f32 %p32, %f173, 0f77F684DF; @%p32 bra BB5_32; bra.uni BB5_33; BB5_32: mov.f32 %f214, 0f39000000; mul.rn.f32 %f279, %f279, %f214; BB5_33: mov.f32 %f215, 0f45800800; mul.rn.f32 %f216, %f31, %f215; sub.f32 %f217, %f31, %f216; add.f32 %f218, %f217, %f216; sub.f32 %f219, %f31, %f218; mul.rn.f32 %f220, %f279, %f215; sub.f32 %f221, %f279, %f220; add.f32 %f222, %f221, %f220; sub.f32 %f223, %f279, %f222; mul.rn.f32 %f224, %f218, %f222; mul.rn.f32 %f225, %f31, %f279; sub.f32 %f226, %f224, %f225; mul.rn.f32 %f227, %f218, %f223; add.f32 %f228, %f226, %f227; mul.rn.f32 %f229, %f219, %f222; add.f32 %f230, %f228, %f229; mul.rn.f32 %f231, %f219, %f223; add.f32 %f232, %f230, %f231; mul.rn.f32 %f233, %f32, %f279; add.f32 %f234, %f233, %f232; add.f32 %f235, %f225, %f234; sub.f32 %f236, %f225, %f235; add.f32 %f35, %f236, %f234; mov.f32 %f281, %f35; mov.f32 %f282, %f235; mov.b32 %r20, %f235; setp.eq.s32 %p33, %r20, 1118925336; @%p33 bra BB5_34; bra.uni BB5_35; BB5_34: add.s32 %r89, %r20, -1; mov.b32 %f237, %r89; add.f32 %f238, %f35, 0f37000000; mov.f32 %f281, %f238; mov.f32 %f282, %f237; BB5_35: mov.f32 %f246, 0f3FB8AA3B; mul.rn.f32 %f240, %f282, %f246; // inline asm cvt.rzi.f32.f32 %f239, %f240; // inline asm mul.rn.f32 %f248, %f239, %f204; sub.f32 %f249, %f282, %f248; mul.rn.f32 %f251, %f239, %f206; sub.f32 %f252, %f249, %f251; mul.rn.f32 %f242, %f252, %f246; // inline asm ex2.approx.f32 %f241, %f242; // inline asm add.f32 %f244, %f239, 0f00000000; // inline asm ex2.approx.f32 %f243, %f244; // inline asm mul.rn.f32 %f253, %f241, %f243; setp.lt.f32 %p34, %f282, 0fC2D20000; selp.f32 %f254, 0f00000000, %f253, %p34; setp.gt.f32 %p35, %f282, 0f42D20000; selp.f32 %f36, %f24, %f254, %p35; setp.neu.f32 %p36, %f36, %f24; @%p36 bra BB5_37; mov.f32 %f280, %f36; bra.uni BB5_46; BB5_37: // inline asm mad.f32 %f255, %f36, %f281, %f36; // inline asm mov.f32 %f280, %f255; bra.uni BB5_46; BB5_38: setp.lt.f32 %p37, %f279, 0f00000000; @%p37 bra BB5_40; mov.f32 %f280, 0f00000000; bra.uni BB5_46; BB5_40: mov.f32 %f280, %f24; bra.uni BB5_46; BB5_41: setp.lt.f32 %p38, %f152, 0f3F800000; mov.b32 %r90, %f279; setp.lt.s32 %p4, %r90, 0; @%p38 bra BB5_43; selp.f32 %f280, 0f00000000, %f24, %p4; bra.uni BB5_46; BB5_43: selp.f32 %f280, %f24, 0f00000000, %p4; bra.uni BB5_46; BB5_44: add.f32 %f280, %f279, 0f40800000; bra.uni BB5_46; BB5_45: mov.f32 %f280, 0f3F800000; BB5_46: cvt.rzi.s32.f32 %r21, %f280; ld.param.u32 %r129, [butterfC_param_8]; setp.eq.s32 %p39, %r129, 1; @%p39 bra BB5_51; ld.param.u32 %r128, [butterfC_param_8]; setp.eq.s32 %p40, %r128, 2; @%p40 bra BB5_50; ld.param.u32 %r127, [butterfC_param_8]; setp.ne.s32 %p41, %r127, 3; @%p41 bra BB5_52; div.s32 %r91, %r5, %r21; rem.s32 %r92, %r5, %r21; mad.lo.s32 %r132, %r91, %r13, %r92; bra.uni BB5_52; BB5_50: div.s32 %r93, %r4, %r21; rem.s32 %r94, %r4, %r21; mad.lo.s32 %r132, %r93, %r13, %r94; bra.uni BB5_52; BB5_51: div.s32 %r95, %r3, %r21; rem.s32 %r96, %r3, %r21; mad.lo.s32 %r132, %r95, %r13, %r96; BB5_52: mov.f64 %fd1, 0d7FF0000000000000; setp.eq.s32 %p42, %r132, 0; @%p42 bra BB5_66; mov.f64 %fd2, 0dFFF0000000000000; neg.f64 %fd24, %fd1; setp.eq.f64 %p43, %fd24, 0d4000000000000000; @%p43 bra BB5_66; mov.f64 %fd241, 0d4000000000000000; // inline asm abs.f64 %fd25, %fd241; // inline asm // inline asm { .reg .b32 lo; mov.b64 {lo, %r98}, %fd25; } // inline asm // inline asm { .reg .b32 hi; mov.b64 {%r99, hi}, %fd25; } // inline asm shr.u32 %r100, %r98, 20; and.b32 %r135, %r100, 2047; setp.eq.s32 %p44, %r135, 0; mov.u32 %r133, %r99; mov.u32 %r134, %r98; @%p44 bra BB5_55; bra.uni BB5_56; BB5_55: mov.f64 %fd31, 0d4350000000000000; mul.rn.f64 %fd30, %fd25, %fd31; // inline asm { .reg .b32 lo; mov.b64 {lo, %r101}, %fd30; } // inline asm // inline asm { .reg .b32 hi; mov.b64 {%r102, hi}, %fd30; } // inline asm shr.u32 %r103, %r101, 20; and.b32 %r104, %r103, 2047; add.s32 %r135, %r104, -54; mov.u32 %r133, %r102; mov.u32 %r134, %r101; BB5_56: add.s32 %r136, %r135, -1023; and.b32 %r107, %r134, -2146435073; or.b32 %r106, %r107, 1072693248; // inline asm mov.b64 %fd32, {%r133, %r106}; // inline asm setp.gt.u32 %p45, %r106, 1073127582; mov.f64 %fd240, %fd32; @%p45 bra BB5_57; bra.uni BB5_58; BB5_57: // inline asm { .reg .b32 hi; mov.b64 {%r108, hi}, %fd32; } // inline asm // inline asm { .reg .b32 lo; mov.b64 {lo, %r109}, %fd32; } // inline asm add.s32 %r111, %r109, -1048576; // inline asm mov.b64 %fd35, {%r108, %r111}; // inline asm add.s32 %r136, %r135, -1022; mov.f64 %fd240, %fd35; BB5_58: add.f64 %fd114, %fd240, 0d3FF0000000000000; rcp.rn.f64 %fd115, %fd114; add.f64 %fd62, %fd240, 0dBFF0000000000000; mul.rn.f64 %fd116, %fd62, %fd115; add.f64 %fd110, %fd116, %fd116; mul.rn.f64 %fd58, %fd110, %fd110; mov.f64 %fd37, 0d3EB0F5FF7D2CAFE2; mov.f64 %fd39, 0d3ED0F5D241AD3B5A; // inline asm fma.rn.f64 %fd36, %fd37, %fd58, %fd39; // inline asm mov.f64 %fd43, 0d3EF3B20A75488A3F; // inline asm fma.rn.f64 %fd40, %fd36, %fd58, %fd43; // inline asm mov.f64 %fd47, 0d3F1745CDE4FAECD5; // inline asm fma.rn.f64 %fd44, %fd40, %fd58, %fd47; // inline asm mov.f64 %fd51, 0d3F3C71C7258A578B; // inline asm fma.rn.f64 %fd48, %fd44, %fd58, %fd51; // inline asm mov.f64 %fd55, 0d3F6249249242B910; // inline asm fma.rn.f64 %fd52, %fd48, %fd58, %fd55; // inline asm mov.f64 %fd59, 0d3F89999999999DFB; // inline asm fma.rn.f64 %fd56, %fd52, %fd58, %fd59; // inline asm mul.rn.f64 %fd117, %fd56, %fd58; sub.f64 %fd118, %fd62, %fd110; mul.rn.f64 %fd63, %fd241, %fd118; neg.f64 %fd61, %fd110; // inline asm fma.rn.f64 %fd60, %fd61, %fd62, %fd63; // inline asm mul.rn.f64 %fd106, %fd115, %fd60; mov.f64 %fd119, 0d3FB5555555555555; add.f64 %fd120, %fd117, 0d3FB5555555555555; sub.f64 %fd121, %fd119, %fd120; add.f64 %fd122, %fd121, %fd117; add.f64 %fd123, %fd122, 0d0000000000000000; add.f64 %fd124, %fd123, 0dBC46A4CB00B9E7B0; add.f64 %fd73, %fd120, %fd124; sub.f64 %fd125, %fd120, %fd73; add.f64 %fd77, %fd125, %fd124; mul.rn.f64 %fd126, %fd73, %fd110; neg.f64 %fd67, %fd126; // inline asm fma.rn.f64 %fd64, %fd73, %fd110, %fd67; // inline asm // inline asm fma.rn.f64 %fd68, %fd77, %fd106, %fd64; // inline asm // inline asm fma.rn.f64 %fd72, %fd73, %fd106, %fd68; // inline asm // inline asm fma.rn.f64 %fd76, %fd77, %fd110, %fd72; // inline asm add.f64 %fd89, %fd126, %fd76; sub.f64 %fd127, %fd126, %fd89; add.f64 %fd93, %fd127, %fd76; mul.rn.f64 %fd128, %fd89, %fd110; neg.f64 %fd83, %fd128; // inline asm fma.rn.f64 %fd80, %fd89, %fd110, %fd83; // inline asm // inline asm fma.rn.f64 %fd84, %fd93, %fd106, %fd80; // inline asm // inline asm fma.rn.f64 %fd88, %fd89, %fd106, %fd84; // inline asm // inline asm fma.rn.f64 %fd92, %fd93, %fd110, %fd88; // inline asm add.f64 %fd105, %fd128, %fd92; sub.f64 %fd129, %fd128, %fd105; add.f64 %fd109, %fd129, %fd92; mul.rn.f64 %fd130, %fd105, %fd110; neg.f64 %fd99, %fd130; // inline asm fma.rn.f64 %fd96, %fd105, %fd110, %fd99; // inline asm // inline asm fma.rn.f64 %fd100, %fd109, %fd106, %fd96; // inline asm // inline asm fma.rn.f64 %fd104, %fd105, %fd106, %fd100; // inline asm // inline asm fma.rn.f64 %fd108, %fd109, %fd110, %fd104; // inline asm add.f64 %fd131, %fd130, %fd108; sub.f64 %fd132, %fd130, %fd131; add.f64 %fd133, %fd132, %fd108; add.f64 %fd134, %fd110, %fd131; sub.f64 %fd135, %fd110, %fd134; add.f64 %fd136, %fd135, %fd131; add.f64 %fd137, %fd136, %fd133; add.f64 %fd138, %fd137, %fd106; add.f64 %fd139, %fd134, %fd138; sub.f64 %fd140, %fd134, %fd139; add.f64 %fd141, %fd140, %fd138; cvt.rn.f64.s32 %fd142, %r136; mov.f64 %fd143, 0d3FE62E42FEFA3000; mul.rn.f64 %fd144, %fd142, %fd143; mov.f64 %fd145, 0d3D53DE6AF278ECE6; mul.rn.f64 %fd146, %fd142, %fd145; add.f64 %fd147, %fd144, %fd139; sub.f64 %fd148, %fd144, %fd147; add.f64 %fd149, %fd148, %fd139; add.f64 %fd150, %fd149, %fd141; add.f64 %fd151, %fd150, %fd146; add.f64 %fd7, %fd147, %fd151; sub.f64 %fd152, %fd147, %fd7; add.f64 %fd8, %fd152, %fd151; // inline asm abs.f64 %fd112, %fd241; // inline asm setp.gt.f64 %p46, %fd112, 0d7F0D2A1BE4048F90; @%p46 bra BB5_59; bra.uni BB5_60; BB5_59: mov.f64 %fd154, 0d3F20000000000000; mov.f64 %fd155, 0d4000000000000000; mul.rn.f64 %fd241, %fd155, %fd154; BB5_60: mul.rn.f64 %fd165, %fd7, %fd241; neg.f64 %fd159, %fd165; // inline asm fma.rn.f64 %fd156, %fd7, %fd241, %fd159; // inline asm // inline asm fma.rn.f64 %fd160, %fd8, %fd241, %fd156; // inline asm add.f64 %fd164, %fd165, %fd160; // inline asm { .reg .b32 lo; mov.b64 {lo, %r112}, %fd164; } // inline asm setp.lt.u32 %p47, %r112, 1082535491; setp.lt.s32 %p48, %r112, -1064875759; or.pred %p49, %p47, %p48; @%p49 bra BB5_62; setp.lt.s32 %p50, %r112, 0; selp.f64 %fd167, 0d0000000000000000, %fd1, %p50; add.f64 %fd168, %fd164, %fd164; setp.nan.f64 %p51, %fd164, %fd164; selp.f64 %fd243, %fd168, %fd167, %p51; bra.uni BB5_65; BB5_62: mov.f64 %fd171, 0d3FF71547652B82FE; mul.rn.f64 %fd170, %fd164, %fd171; // inline asm abs.f64 %fd169, %fd170; // inline asm setp.gt.f64 %p52, %fd169, 0d4330000000000000; mov.f64 %fd242, %fd170; @%p52 bra BB5_64; add.f64 %fd173, %fd169, 0d3FE0000000000000; // inline asm cvt.rmi.f64.f64 %fd172, %fd173; // inline asm setp.lt.f64 %p53, %fd169, 0d3FE0000000000000; selp.f64 %fd176, 0d0000000000000000, %fd172, %p53; // inline asm { .reg .b32 lo; mov.b64 {lo, %r113}, %fd170; } // inline asm // inline asm { .reg .b32 hi; mov.b64 {%r114, hi}, %fd176; } // inline asm // inline asm { .reg .b32 lo; mov.b64 {lo, %r115}, %fd176; } // inline asm and.b32 %r118, %r113, -2147483648; and.b32 %r119, %r115, 2147483647; or.b32 %r117, %r119, %r118; // inline asm mov.b64 %fd177, {%r114, %r117}; // inline asm mov.f64 %fd242, %fd177; BB5_64: mov.f64 %fd180, 0dBFE62E42FEFA39EF; // inline asm fma.rn.f64 %fd178, %fd242, %fd180, %fd164; // inline asm mov.f64 %fd184, 0dBC7ABC9E3B39803F; // inline asm fma.rn.f64 %fd182, %fd242, %fd184, %fd178; // inline asm cvt.rzi.s32.f64 %r122, %fd242; add.s32 %r123, %r122, 55; setp.lt.s32 %p54, %r122, -1020; selp.b32 %r124, %r123, %r122, %p54; selp.f64 %fd235, 0d3C90000000000000, 0d4000000000000000, %p54; mov.f64 %fd187, 0d3E21F07FCCF58BAD; mov.f64 %fd189, 0d3E5AFD81DA6C3BAF; // inline asm fma.rn.f64 %fd186, %fd187, %fd182, %fd189; // inline asm mov.f64 %fd193, 0d3E927E55F60F80E6; // inline asm fma.rn.f64 %fd190, %fd186, %fd182, %fd193; // inline asm mov.f64 %fd197, 0d3EC71DDA8F02D666; // inline asm fma.rn.f64 %fd194, %fd190, %fd182, %fd197; // inline asm mov.f64 %fd201, 0d3EFA01A013B894E0; // inline asm fma.rn.f64 %fd198, %fd194, %fd182, %fd201; // inline asm mov.f64 %fd205, 0d3F2A01A01D3AF788; // inline asm fma.rn.f64 %fd202, %fd198, %fd182, %fd205; // inline asm mov.f64 %fd209, 0d3F56C16C16C3A1EC; // inline asm fma.rn.f64 %fd206, %fd202, %fd182, %fd209; // inline asm mov.f64 %fd213, 0d3F81111111109161; // inline asm fma.rn.f64 %fd210, %fd206, %fd182, %fd213; // inline asm mov.f64 %fd217, 0d3FA55555555554C1; // inline asm fma.rn.f64 %fd214, %fd210, %fd182, %fd217; // inline asm mov.f64 %fd221, 0d3FC555555555556F; // inline asm fma.rn.f64 %fd218, %fd214, %fd182, %fd221; // inline asm mov.f64 %fd225, 0d3FE0000000000000; // inline asm fma.rn.f64 %fd222, %fd218, %fd182, %fd225; // inline asm mul.rn.f64 %fd227, %fd222, %fd182; // inline asm fma.rn.f64 %fd226, %fd227, %fd182, %fd182; // inline asm shl.b32 %r125, %r124, 20; add.s32 %r121, %r125, 1071644672; mov.u32 %r120, 0; // inline asm mov.b64 %fd230, {%r120, %r121}; // inline asm // inline asm fma.rn.f64 %fd231, %fd226, %fd230, %fd230; // inline asm mul.rn.f64 %fd243, %fd231, %fd235; BB5_65: setp.eq.f64 %p55, %fd243, %fd2; setp.eq.f64 %p56, %fd243, %fd1; or.pred %p57, %p56, %p55; BB5_66: ret; }