.func (.param .align 4 .b8 func_retval0[32]) activation_8( .param .align 4 .b8 activation_8_param_0[32] ) { .pragma "abi_param_reg all"; .reg .pred %p<17>; .reg .b16 %rs<33>; .reg .f32 %f<58>; .reg .b32 %r<17>; ld.param.u32 %r1, [activation_8_param_0+28]; ld.param.u32 %r2, [activation_8_param_0+24]; ld.param.u32 %r3, [activation_8_param_0+20]; ld.param.u32 %r4, [activation_8_param_0+16]; ld.param.u32 %r5, [activation_8_param_0+12]; ld.param.u32 %r6, [activation_8_param_0+8]; ld.param.u32 %r7, [activation_8_param_0+4]; ld.param.u32 %r8, [activation_8_param_0]; ld.const.f32 %f33, [params]; mov.b32 {%rs1, %rs3}, %r8; // begin inline asm { cvt.f32.f16 %f1, %rs1;} // end inline asm setp.lt.ftz.f32 %p1, %f1, 0f00000000; mul.ftz.f32 %f34, %f1, %f33; selp.f32 %f2, %f34, %f1, %p1; // begin inline asm { cvt.f32.f16 %f3, %rs3;} // end inline asm setp.lt.ftz.f32 %p2, %f3, 0f00000000; mul.ftz.f32 %f35, %f3, %f33; selp.f32 %f4, %f35, %f3, %p2; // begin inline asm { cvt.rn.f16.f32 %rs4, %f4;} // end inline asm // begin inline asm { cvt.rn.f16.f32 %rs2, %f2;} // end inline asm mov.b32 %r9, {%rs2, %rs4}; mov.b32 {%rs5, %rs7}, %r7; // begin inline asm { cvt.f32.f16 %f5, %rs5;} // end inline asm setp.lt.ftz.f32 %p3, %f5, 0f00000000; mul.ftz.f32 %f36, %f5, %f33; selp.f32 %f6, %f36, %f5, %p3; // begin inline asm { cvt.f32.f16 %f7, %rs7;} // end inline asm setp.lt.ftz.f32 %p4, %f7, 0f00000000; mul.ftz.f32 %f37, %f7, %f33; selp.f32 %f8, %f37, %f7, %p4; // begin inline asm { cvt.rn.f16.f32 %rs8, %f8;} // end inline asm // begin inline asm { cvt.rn.f16.f32 %rs6, %f6;} // end inline asm mov.b32 %r10, {%rs6, %rs8}; mov.b32 {%rs9, %rs11}, %r6; // begin inline asm { cvt.f32.f16 %f9, %rs9;} // end inline asm setp.lt.ftz.f32 %p5, %f9, 0f00000000; mul.ftz.f32 %f38, %f9, %f33; selp.f32 %f10, %f38, %f9, %p5; // begin inline asm { cvt.f32.f16 %f11, %rs11;} // end inline asm setp.lt.ftz.f32 %p6, %f11, 0f00000000; mul.ftz.f32 %f39, %f11, %f33; selp.f32 %f12, %f39, %f11, %p6; // begin inline asm { cvt.rn.f16.f32 %rs12, %f12;} // end inline asm // begin inline asm { cvt.rn.f16.f32 %rs10, %f10;} // end inline asm mov.b32 %r11, {%rs10, %rs12}; mov.b32 {%rs13, %rs15}, %r5; // begin inline asm { cvt.f32.f16 %f13, %rs13;} // end inline asm setp.lt.ftz.f32 %p7, %f13, 0f00000000; mul.ftz.f32 %f40, %f13, %f33; selp.f32 %f14, %f40, %f13, %p7; // begin inline asm { cvt.f32.f16 %f15, %rs15;} // end inline asm setp.lt.ftz.f32 %p8, %f15, 0f00000000; mul.ftz.f32 %f41, %f15, %f33; selp.f32 %f16, %f41, %f15, %p8; // begin inline asm { cvt.rn.f16.f32 %rs16, %f16;} // end inline asm // begin inline asm { cvt.rn.f16.f32 %rs14, %f14;} // end inline asm mov.b32 %r12, {%rs14, %rs16}; mov.b32 {%rs17, %rs19}, %r4; // begin inline asm { cvt.f32.f16 %f17, %rs17;} // end inline asm setp.lt.ftz.f32 %p9, %f17, 0f00000000; mul.ftz.f32 %f42, %f17, %f33; selp.f32 %f18, %f42, %f17, %p9; // begin inline asm { cvt.f32.f16 %f19, %rs19;} // end inline asm setp.lt.ftz.f32 %p10, %f19, 0f00000000; mul.ftz.f32 %f43, %f19, %f33; selp.f32 %f20, %f43, %f19, %p10; // begin inline asm { cvt.rn.f16.f32 %rs20, %f20;} // end inline asm // begin inline asm { cvt.rn.f16.f32 %rs18, %f18;} // end inline asm mov.b32 %r13, {%rs18, %rs20}; mov.b32 {%rs21, %rs23}, %r3; // begin inline asm { cvt.f32.f16 %f21, %rs21;} // end inline asm setp.lt.ftz.f32 %p11, %f21, 0f00000000; mul.ftz.f32 %f44, %f21, %f33; selp.f32 %f22, %f44, %f21, %p11; // begin inline asm { cvt.f32.f16 %f23, %rs23;} // end inline asm setp.lt.ftz.f32 %p12, %f23, 0f00000000; mul.ftz.f32 %f45, %f23, %f33; selp.f32 %f24, %f45, %f23, %p12; // begin inline asm { cvt.rn.f16.f32 %rs24, %f24;} // end inline asm // begin inline asm { cvt.rn.f16.f32 %rs22, %f22;} // end inline asm mov.b32 %r14, {%rs22, %rs24}; mov.b32 {%rs25, %rs27}, %r2; // begin inline asm { cvt.f32.f16 %f25, %rs25;} // end inline asm setp.lt.ftz.f32 %p13, %f25, 0f00000000; mul.ftz.f32 %f46, %f25, %f33; selp.f32 %f26, %f46, %f25, %p13; // begin inline asm { cvt.f32.f16 %f27, %rs27;} // end inline asm setp.lt.ftz.f32 %p14, %f27, 0f00000000; mul.ftz.f32 %f47, %f27, %f33; selp.f32 %f28, %f47, %f27, %p14; // begin inline asm { cvt.rn.f16.f32 %rs28, %f28;} // end inline asm // begin inline asm { cvt.rn.f16.f32 %rs26, %f26;} // end inline asm mov.b32 %r15, {%rs26, %rs28}; mov.b32 {%rs29, %rs31}, %r1; // begin inline asm { cvt.f32.f16 %f29, %rs29;} // end inline asm setp.lt.ftz.f32 %p15, %f29, 0f00000000; mul.ftz.f32 %f48, %f29, %f33; selp.f32 %f30, %f48, %f29, %p15; // begin inline asm { cvt.f32.f16 %f31, %rs31;} // end inline asm setp.lt.ftz.f32 %p16, %f31, 0f00000000; mul.ftz.f32 %f49, %f31, %f33; selp.f32 %f32, %f49, %f31, %p16; // begin inline asm { cvt.rn.f16.f32 %rs32, %f32;} // end inline asm // begin inline asm { cvt.rn.f16.f32 %rs30, %f30;} // end inline asm mov.b32 %r16, {%rs30, %rs32}; mov.b32 %f50, %r16; mov.b32 %f51, %r15; mov.b32 %f52, %r14; mov.b32 %f53, %r13; mov.b32 %f54, %r12; mov.b32 %f55, %r11; mov.b32 %f56, %r10; mov.b32 %f57, %r9; st.param.f32 [func_retval0+0], %f57; st.param.f32 [func_retval0+4], %f56; st.param.f32 [func_retval0+8], %f55; st.param.f32 [func_retval0+12], %f54; st.param.f32 [func_retval0+16], %f53; st.param.f32 [func_retval0+20], %f52; st.param.f32 [func_retval0+24], %f51; st.param.f32 [func_retval0+28], %f50; ret; }