.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 .pred %p<9>; .reg .b16 %rs<17>; .reg .f32 %f<22>; .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]; mov.b32 {%rs1, %rs3}, %r4; // begin inline asm { cvt.f32.f16 %f1, %rs1;} // end inline asm setp.ge.ftz.f32 %p1, %f17, %f1; selp.f32 %f2, 0f00000000, %f1, %p1; // begin inline asm { cvt.f32.f16 %f3, %rs3;} // end inline asm setp.ge.ftz.f32 %p2, %f17, %f3; selp.f32 %f4, 0f00000000, %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 %r5, {%rs2, %rs4}; mov.b32 {%rs5, %rs7}, %r3; // begin inline asm { cvt.f32.f16 %f5, %rs5;} // end inline asm setp.ge.ftz.f32 %p3, %f17, %f5; selp.f32 %f6, 0f00000000, %f5, %p3; // begin inline asm { cvt.f32.f16 %f7, %rs7;} // end inline asm setp.ge.ftz.f32 %p4, %f17, %f7; selp.f32 %f8, 0f00000000, %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 %r6, {%rs6, %rs8}; mov.b32 {%rs9, %rs11}, %r2; // begin inline asm { cvt.f32.f16 %f9, %rs9;} // end inline asm setp.ge.ftz.f32 %p5, %f17, %f9; selp.f32 %f10, 0f00000000, %f9, %p5; // begin inline asm { cvt.f32.f16 %f11, %rs11;} // end inline asm setp.ge.ftz.f32 %p6, %f17, %f11; selp.f32 %f12, 0f00000000, %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 %r7, {%rs10, %rs12}; mov.b32 {%rs13, %rs15}, %r1; // begin inline asm { cvt.f32.f16 %f13, %rs13;} // end inline asm setp.ge.ftz.f32 %p7, %f17, %f13; selp.f32 %f14, 0f00000000, %f13, %p7; // begin inline asm { cvt.f32.f16 %f15, %rs15;} // end inline asm setp.ge.ftz.f32 %p8, %f17, %f15; selp.f32 %f16, 0f00000000, %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 %r8, {%rs14, %rs16}; mov.b32 %f18, %r8; mov.b32 %f19, %r7; mov.b32 %f20, %r6; mov.b32 %f21, %r5; st.param.f32 [func_retval0+0], %f21; st.param.f32 [func_retval0+4], %f20; st.param.f32 [func_retval0+8], %f19; st.param.f32 [func_retval0+12], %f18; ret; }