.version 7.8 .target sm_80 .address_size 64 // .globl activation_4 .visible .const .align 4 .b8 params[8]; .visible .func (.param .align 4 .b8 func_retval0[16]) activation_4( .param .align 4 .b8 activation_4_param_0[16] ) { .pragma "abi_param_reg all"; .reg .b16 %rs<17>; .reg .f32 %f<41>; .reg .b32 %r<9>; ld.param.u32 %r1, [activation_4_param_0+12]; ld.param.u32 %r2, [activation_4_param_0+8]; ld.param.u32 %r3, [activation_4_param_0+4]; ld.param.u32 %r4, [activation_4_param_0]; ld.const.f32 %f17, [params]; ld.const.f32 %f18, [params+4]; mov.b32 {%rs1, %rs3}, %r4; // begin inline asm { cvt.f32.f16 %f1, %rs1;} // end inline asm fma.rn.ftz.f32 %f19, %f1, %f17, %f18; mov.f32 %f20, 0f3F800000; min.ftz.f32 %f21, %f20, %f19; mov.f32 %f22, 0f00000000; max.ftz.f32 %f2, %f22, %f21; // begin inline asm { cvt.f32.f16 %f3, %rs3;} // end inline asm fma.rn.ftz.f32 %f23, %f3, %f17, %f18; min.ftz.f32 %f24, %f20, %f23; max.ftz.f32 %f4, %f22, %f24; // 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 %r5, {%rs2, %rs4}; mov.b32 {%rs5, %rs7}, %r3; // begin inline asm { cvt.f32.f16 %f5, %rs5;} // end inline asm fma.rn.ftz.f32 %f25, %f5, %f17, %f18; min.ftz.f32 %f26, %f20, %f25; max.ftz.f32 %f6, %f22, %f26; // begin inline asm { cvt.f32.f16 %f7, %rs7;} // end inline asm fma.rn.ftz.f32 %f27, %f7, %f17, %f18; min.ftz.f32 %f28, %f20, %f27; max.ftz.f32 %f8, %f22, %f28; // 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 %r6, {%rs6, %rs8}; mov.b32 {%rs9, %rs11}, %r2; // begin inline asm { cvt.f32.f16 %f9, %rs9;} // end inline asm fma.rn.ftz.f32 %f29, %f9, %f17, %f18; min.ftz.f32 %f30, %f20, %f29; max.ftz.f32 %f10, %f22, %f30; // begin inline asm { cvt.f32.f16 %f11, %rs11;} // end inline asm fma.rn.ftz.f32 %f31, %f11, %f17, %f18; min.ftz.f32 %f32, %f20, %f31; max.ftz.f32 %f12, %f22, %f32; // 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 %r7, {%rs10, %rs12}; mov.b32 {%rs13, %rs15}, %r1; // begin inline asm { cvt.f32.f16 %f13, %rs13;} // end inline asm fma.rn.ftz.f32 %f33, %f13, %f17, %f18; min.ftz.f32 %f34, %f20, %f33; max.ftz.f32 %f14, %f22, %f34; // begin inline asm { cvt.f32.f16 %f15, %rs15;} // end inline asm fma.rn.ftz.f32 %f35, %f15, %f17, %f18; min.ftz.f32 %f36, %f20, %f35; max.ftz.f32 %f16, %f22, %f36; // 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 %r8, {%rs14, %rs16}; mov.b32 %f37, %r8; mov.b32 %f38, %r7; mov.b32 %f39, %r6; mov.b32 %f40, %r5; st.param.f32 [func_retval0+0], %f40; st.param.f32 [func_retval0+4], %f39; st.param.f32 [func_retval0+8], %f38; st.param.f32 [func_retval0+12], %f37; ret; }