ase 11.8, V11.8.85 // Based on NVVM 7.0.1 // .version 7.8 .target sm_80 .address_size 64 // .globl activation_16 .visible .const .align 4 .b8 params[8]; .visible .func (.param .align 4 .b8 func_retval0[64]) activation_16( .param .align 4 .b8 activation_16_param_0[64] ) { .pragma "abi_param_reg all"; .reg .pred %p<33>; .reg .b16 %rs<65>; .reg .f32 %f<114>; .reg .b32 %r<33>; ld.param.u32 %r1, [activation_16_param_0+60]; ld.param.u32 %r2, [activation_16_param_0+56]; ld.param.u32 %r3, [activation_16_param_0+52]; ld.param.u32 %r4, [activation_16_param_0+48]; ld.param.u32 %r5, [activation_16_param_0+44]; ld.param.u32 %r6, [activation_16_param_0+40]; ld.param.u32 %r7, [activation_16_param_0+36]; ld.param.u32 %r8, [activation_16_param_0+32]; ld.param.u32 %r9, [activation_16_param_0+28]; ld.param.u32 %r10, [activation_16_param_0+24]; ld.param.u32 %r11, [activation_16_param_0+20]; ld.param.u32 %r12, [activation_16_param_0+16]; ld.param.u32 %r13, [activation_16_param_0+12]; ld.param.u32 %r14, [activation_16_param_0+8]; ld.param.u32 %r15, [activation_16_param_0+4]; ld.param.u32 %r16, [activation_16_param_0]; ld.const.f32 %f65, [params]; mov.b32 {%rs1, %rs3}, %r16; // begin inline asm { cvt.f32.f16 %f1, %rs1;} // end inline asm setp.lt.ftz.f32 %p1, %f1, 0f00000000; mul.ftz.f32 %f66, %f1, %f65; selp.f32 %f2, %f66, %f1, %p1; // begin inline asm { cvt.f32.f16 %f3, %rs3;} // end inline asm setp.lt.ftz.f32 %p2, %f3, 0f00000000; mul.ftz.f32 %f67, %f3, %f65; selp.f32 %f4, %f67, %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 %r17, {%rs2, %rs4}; mov.b32 {%rs5, %rs7}, %r15; // begin inline asm { cvt.f32.f16 %f5, %rs5;} // end inline asm setp.lt.ftz.f32 %p3, %f5, 0f00000000; mul.ftz.f32 %f68, %f5, %f65; selp.f32 %f6, %f68, %f5, %p3; // begin inline asm { cvt.f32.f16 %f7, %rs7;} // end inline asm setp.lt.ftz.f32 %p4, %f7, 0f00000000; mul.ftz.f32 %f69, %f7, %f65; selp.f32 %f8, %f69, %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 %r18, {%rs6, %rs8}; mov.b32 {%rs9, %rs11}, %r14; // begin inline asm { cvt.f32.f16 %f9, %rs9;} // end inline asm setp.lt.ftz.f32 %p5, %f9, 0f00000000; mul.ftz.f32 %f70, %f9, %f65; selp.f32 %f10, %f70, %f9, %p5; // begin inline asm { cvt.f32.f16 %f11, %rs11;} // end inline asm setp.lt.ftz.f32 %p6, %f11, 0f00000000; mul.ftz.f32 %f71, %f11, %f65; selp.f32 %f12, %f71, %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 %r19, {%rs10, %rs12}; mov.b32 {%rs13, %rs15}, %r13; // begin inline asm { cvt.f32.f16 %f13, %rs13;} // end inline asm setp.lt.ftz.f32 %p7, %f13, 0f00000000; mul.ftz.f32 %f72, %f13, %f65; selp.f32 %f14, %f72, %f13, %p7; // begin inline asm { cvt.f32.f16 %f15, %rs15;} // end inline asm setp.lt.ftz.f32 %p8, %f15, 0f00000000; mul.ftz.f32 %f73, %f15, %f65; selp.f32 %f16, %f73, %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 %r20, {%rs14, %rs16}; mov.b32 {%rs17, %rs19}, %r12; // begin inline asm { cvt.f32.f16 %f17, %rs17;} // end inline asm setp.lt.ftz.f32 %p9, %f17, 0f00000000; mul.ftz.f32 %f74, %f17, %f65; selp.f32 %f18, %f74, %f17, %p9; // begin inline asm { cvt.f32.f16 %f19, %rs19;} // end inline asm setp.lt.ftz.f32 %p10, %f19, 0f00000000; mul.ftz.f32 %f75, %f19, %f65; selp.f32 %f20, %f75, %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 %r21, {%rs18, %rs20}; mov.b32 {%rs21, %rs23}, %r11; // begin inline asm { cvt.f32.f16 %f21, %rs21;} // end inline asm setp.lt.ftz.f32 %p11, %f21, 0f00000000; mul.ftz.f32 %f76, %f21, %f65; selp.f32 %f22, %f76, %f21, %p11; // begin inline asm { cvt.f32.f16 %f23, %rs23;} // end inline asm setp.lt.ftz.f32 %p12, %f23, 0f00000000; mul.ftz.f32 %f77, %f23, %f65; selp.f32 %f24, %f77, %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 %r22, {%rs22, %rs24}; mov.b32 {%rs25, %rs27}, %r10; // begin inline asm { cvt.f32.f16 %f25, %rs25;} // end inline asm setp.lt.ftz.f32 %p13, %f25, 0f00000000; mul.ftz.f32 %f78, %f25, %f65; selp.f32 %f26, %f78, %f25, %p13; // begin inline asm { cvt.f32.f16 %f27, %rs27;} // end inline asm setp.lt.ftz.f32 %p14, %f27, 0f00000000; mul.ftz.f32 %f79, %f27, %f65; selp.f32 %f28, %f79, %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 %r23, {%rs26, %rs28}; mov.b32 {%rs29, %rs31}, %r9; // begin inline asm { cvt.f32.f16 %f29, %rs29;} // end inline asm setp.lt.ftz.f32 %p15, %f29, 0f00000000; mul.ftz.f32 %f80, %f29, %f65; selp.f32 %f30, %f80, %f29, %p15; // begin inline asm { cvt.f32.f16 %f31, %rs31;} // end inline asm setp.lt.ftz.f32 %p16, %f31, 0f00000000; mul.ftz.f32 %f81, %f31, %f65; selp.f32 %f32, %f81, %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 %r24, {%rs30, %rs32}; mov.b32 {%rs33, %rs35}, %r8; // begin inline asm { cvt.f32.f16 %f33, %rs33;} // end inline asm setp.lt.ftz.f32 %p17, %f33, 0f00000000; mul.ftz.f32 %f82, %f33, %f65; selp.f32 %f34, %f82, %f33, %p17; // begin inline asm { cvt.f32.f16 %f35, %rs35;} // end inline asm setp.lt.ftz.f32 %p18, %f35, 0f00000000; mul.ftz.f32 %f83, %f35, %f65; selp.f32 %f36, %f83, %f35, %p18; // begin inline asm { cvt.rn.f16.f32 %rs36, %f36;} // end inline asm // begin inline asm { cvt.rn.f16.f32 %rs34, %f34;} // end inline asm mov.b32 %r25, {%rs34, %rs36}; mov.b32 {%rs37, %rs39}, %r7; // begin inline asm { cvt.f32.f16 %f37, %rs37;} // end inline asm setp.lt.ftz.f32 %p19, %f37, 0f00000000; mul.ftz.f32 %f84, %f37, %f65; selp.f32 %f38, %f84, %f37, %p19; // begin inline asm { cvt.f32.f16 %f39, %rs39;} // end inline asm setp.lt.ftz.f32 %p20, %f39, 0f00000000; mul.ftz.f32 %f85, %f39, %f65; selp.f32 %f40, %f85, %f39, %p20; // begin inline asm { cvt.rn.f16.f32 %rs40, %f40;} // end inline asm // begin inline asm { cvt.rn.f16.f32 %rs38, %f38;} // end inline asm mov.b32 %r26, {%rs38, %rs40}; mov.b32 {%rs41, %rs43}, %r6; // begin inline asm { cvt.f32.f16 %f41, %rs41;} // end inline asm setp.lt.ftz.f32 %p21, %f41, 0f00000000; mul.ftz.f32 %f86, %f41, %f65; selp.f32 %f42, %f86, %f41, %p21; // begin inline asm { cvt.f32.f16 %f43, %rs43;} // end inline asm setp.lt.ftz.f32 %p22, %f43, 0f00000000; mul.ftz.f32 %f87, %f43, %f65; selp.f32 %f44, %f87, %f43, %p22; // begin inline asm { cvt.rn.f16.f32 %rs44, %f44;} // end inline asm // begin inline asm { cvt.rn.f16.f32 %rs42, %f42;} // end inline asm mov.b32 %r27, {%rs42, %rs44}; mov.b32 {%rs45, %rs47}, %r5; // begin inline asm { cvt.f32.f16 %f45, %rs45;} // end inline asm setp.lt.ftz.f32 %p23, %f45, 0f00000000; mul.ftz.f32 %f88, %f45, %f65; selp.f32 %f46, %f88, %f45, %p23; // begin inline asm { cvt.f32.f16 %f47, %rs47;} // end inline asm setp.lt.ftz.f32 %p24, %f47, 0f00000000; mul.ftz.f32 %f89, %f47, %f65; selp.f32 %f48, %f89, %f47, %p24; // begin inline asm { cvt.rn.f16.f32 %rs48, %f48;} // end inline asm // begin inline asm { cvt.rn.f16.f32 %rs46, %f46;} // end inline asm mov.b32 %r28, {%rs46, %rs48}; mov.b32 {%rs49, %rs51}, %r4; // begin inline asm { cvt.f32.f16 %f49, %rs49;} // end inline asm setp.lt.ftz.f32 %p25, %f49, 0f00000000; mul.ftz.f32 %f90, %f49, %f65; selp.f32 %f50, %f90, %f49, %p25; // begin inline asm { cvt.f32.f16 %f51, %rs51;} // end inline asm setp.lt.ftz.f32 %p26, %f51, 0f00000000; mul.ftz.f32 %f91, %f51, %f65; selp.f32 %f52, %f91, %f51, %p26; // begin inline asm { cvt.rn.f16.f32 %rs52, %f52;} // end inline asm // begin inline asm { cvt.rn.f16.f32 %rs50, %f50;} // end inline asm mov.b32 %r29, {%rs50, %rs52}; mov.b32 {%rs53, %rs55}, %r3; // begin inline asm { cvt.f32.f16 %f53, %rs53;} // end inline asm setp.lt.ftz.f32 %p27, %f53, 0f00000000; mul.ftz.f32 %f92, %f53, %f65; selp.f32 %f54, %f92, %f53, %p27; // begin inline asm { cvt.f32.f16 %f55, %rs55;} // end inline asm setp.lt.ftz.f32 %p28, %f55, 0f00000000; mul.ftz.f32 %f93, %f55, %f65; selp.f32 %f56, %f93, %f55, %p28; // begin inline asm { cvt.rn.f16.f32 %rs56, %f56;} // end inline asm // begin inline asm { cvt.rn.f16.f32 %rs54, %f54;} // end inline asm mov.b32 %r30, {%rs54, %rs56}; mov.b32 {%rs57, %rs59}, %r2; // begin inline asm { cvt.f32.f16 %f57, %rs57;} // end inline asm setp.lt.ftz.f32 %p29, %f57, 0f00000000; mul.ftz.f32 %f94, %f57, %f65; selp.f32 %f58, %f94, %f57, %p29; // begin inline asm { cvt.f32.f16 %f59, %rs59;} // end inline asm setp.lt.ftz.f32 %p30, %f59, 0f00000000; mul.ftz.f32 %f95, %f59, %f65; selp.f32 %f60, %f95, %f59, %p30; // begin inline asm { cvt.rn.f16.f32 %rs60, %f60;} // end inline asm // begin inline asm { cvt.rn.f16.f32 %rs58, %f58;} // end inline asm mov.b32 %r31, {%rs58, %rs60}; mov.b32 {%rs61, %rs63}, %r1; // begin inline asm { cvt.f32.f16 %f61, %rs61;} // end inline asm setp.lt.ftz.f32 %p31, %f61, 0f00000000; mul.ftz.f32 %f96, %f61, %f65; selp.f32 %f62, %f96, %f61, %p31; // begin inline asm { cvt.f32.f16 %f63, %rs63;} // end inline asm setp.lt.ftz.f32 %p32, %f63, 0f00000000; mul.ftz.f32 %f97, %f63, %f65; selp.f32 %f64, %f97, %f63, %p32; // begin inline asm { cvt.rn.f16.f32 %rs64, %f64;} // end inline asm // begin inline asm { cvt.rn.f16.f32 %rs62, %f62;} // end inline asm mov.b32 %r32, {%rs62, %rs64}; mov.b32 %f98, %r32; mov.b32 %f99, %r31; mov.b32 %f100, %r30; mov.b32 %f101, %r29; mov.b32 %f102, %r28; mov.b32 %f103, %r27; mov.b32 %f104, %r26; mov.b32 %f105, %r25; mov.b32 %f106, %r24; mov.b32 %f107, %r23; mov.b32 %f108, %r22; mov.b32 %f109, %r21; mov.b32 %f110, %r20; mov.b32 %f111, %r19; mov.b32 %f112, %r18; mov.b32 %f113, %r17; st.param.f32 [func_retval0+0], %f113; st.param.f32 [func_retval0+4], %f112; st.param.f32 [func_retval0+8], %f111; st.param.f32 [func_retval0+12], %f110; st.param.f32 [func_retval0+16], %f109; st.param.f32 [func_retval0+20], %f108; st.param.f32 [func_retval0+24], %f107; st.param.f32 [func_retval0+28], %f106; st.param.f32 [func_retval0+32], %f105; st.param.f32 [func_retval0+36], %f104; st.param.f32 [func_retval0+40], %f103; st.param.f32 [func_retval0+44], %f102; st.param.f32 [func_retval0+48], %f101; st.param.f32 [func_retval0+52], %f100; st.param.f32 [func_retval0+56], %f99; st.param.f32 [func_retval0+60], %f98; ret; }