900 lines · plain
1# NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 22# RUN: llc -mtriple=amdgcn-amd-amdhsa -mcpu=gfx942 -start-before=machine-scheduler -verify-misched -o - %s | FileCheck -check-prefix=GCN %s3# RUN: llc -mtriple=amdgcn-amd-amdhsa -mcpu=gfx9-4-generic --amdhsa-code-object-version=6 -start-before=machine-scheduler -verify-misched -o - %s | FileCheck -check-prefix=GCN %s4 5--- |6 define amdgpu_kernel void @smallInterleave() #0 { ret void }7 ; GCN-LABEL: smallInterleave:8 ; GCN: ; %bb.0:9 ; GCN-NEXT: ; implicit-def: $vgpr210 ; GCN-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)11 ; GCN-NEXT: v_readfirstlane_b32 s20, v212 ; GCN-NEXT: ; implicit-def: $sgpr413 ; GCN-NEXT: ; implicit-def: $vgpr314 ; GCN-NEXT: ; implicit-def: $vgpr0_vgpr115 ; GCN-NEXT: ; implicit-def: $sgpr0_sgpr1_sgpr2_sgpr316 ; GCN-NEXT: ; implicit-def: $vgpr5017 ; GCN-NEXT: ; implicit-def: $sgpr16_sgpr17_sgpr18_sgpr1918 ; GCN-NEXT: ; implicit-def: $vgpr4919 ; GCN-NEXT: ; implicit-def: $vgpr40_vgpr41_vgpr42_vgpr4320 ; GCN-NEXT: ; implicit-def: $vgpr5121 ; GCN-NEXT: ; implicit-def: $vgpr62_vgpr63_vgpr64_vgpr6522 ; GCN-NEXT: ; implicit-def: $vgpr7623 ; GCN-NEXT: ; implicit-def: $vgpr7724 ; GCN-NEXT: ; implicit-def: $vgpr7825 ; GCN-NEXT: ; implicit-def: $vgpr7926 ; GCN-NEXT: ; implicit-def: $vgpr8027 ; GCN-NEXT: ; implicit-def: $vgpr9128 ; GCN-NEXT: ; kill: killed $sgpr16_sgpr17_sgpr18_sgpr1929 ; GCN-NEXT: ; iglp_opt mask(0x00000002)30 ; GCN-NEXT: s_nop 131 ; GCN-NEXT: v_lshl_add_u32 v2, s20, 4, v332 ; GCN-NEXT: v_mad_u64_u32 v[4:5], s[4:5], s4, v2, v[0:1]33 ; GCN-NEXT: buffer_load_dwordx4 v[0:3], v4, s[0:3], 0 offen sc0 sc134 ; GCN-NEXT: s_waitcnt vmcnt(0)35 ; GCN-NEXT: buffer_inv sc0 sc136 ; GCN-NEXT: s_lshl_b32 s4, s20, 737 ; GCN-NEXT: ; implicit-def: $vgpr538 ; GCN-NEXT: v_add_lshl_u32 v48, v5, s4, 139 ; GCN-NEXT: v_add_u32_e32 v76, s20, v7640 ; GCN-NEXT: v_and_b32_e32 v76, 0x1fffffff, v7641 ; GCN-NEXT: buffer_wbl2 sc0 sc142 ; GCN-NEXT: ds_write_b128 v48, v[0:3]43 ; GCN-NEXT: s_waitcnt lgkmcnt(0)44 ; GCN-NEXT: buffer_load_dwordx4 v[32:35], v4, s[0:3], 0 offen offset:64 sc0 sc145 ; GCN-NEXT: s_waitcnt vmcnt(0)46 ; GCN-NEXT: buffer_inv sc0 sc147 ; GCN-NEXT: ; implicit-def: $vgpr048 ; GCN-NEXT: ; implicit-def: $vgpr149 ; GCN-NEXT: ; implicit-def: $sgpr0_sgpr1_sgpr2_sgpr3_sgpr4_sgpr5_sgpr6_sgpr7_sgpr8_sgpr9_sgpr10_sgpr11_sgpr12_sgpr13_sgpr14_sgpr1550 ; GCN-NEXT: ; implicit-def: $sgpr651 ; GCN-NEXT: v_add_u32_e32 v0, v0, v5052 ; GCN-NEXT: v_add_u32_e32 v1, v1, v5053 ; GCN-NEXT: buffer_load_dwordx2 v[72:73], v0, s[16:19], 0 offen sc0 sc154 ; GCN-NEXT: s_waitcnt vmcnt(0)55 ; GCN-NEXT: buffer_inv sc0 sc156 ; GCN-NEXT: buffer_load_dwordx2 v[74:75], v1, s[16:19], 0 offen sc0 sc157 ; GCN-NEXT: s_waitcnt vmcnt(0)58 ; GCN-NEXT: buffer_inv sc0 sc159 ; GCN-NEXT: ;;#ASMSTART60 ; GCN-NEXT: s_waitcnt vmcnt(8)61 ; GCN-NEXT: ;;#ASMEND62 ; GCN-NEXT: ds_read_b128 v[36:39], v4963 ; GCN-NEXT: s_waitcnt lgkmcnt(0)64 ; GCN-NEXT: buffer_inv sc0 sc165 ; GCN-NEXT: ds_read_b128 v[44:47], v49 offset:51266 ; GCN-NEXT: s_waitcnt lgkmcnt(0)67 ; GCN-NEXT: buffer_inv sc0 sc168 ; GCN-NEXT: v_mfma_f32_32x32x8_f16 v[16:31], v[36:37], v[40:41], 069 ; GCN-NEXT: ; kill: killed $vgpr170 ; GCN-NEXT: ; kill: killed $vgpr071 ; GCN-NEXT: v_mul_lo_u32 v76, v76, s672 ; GCN-NEXT: v_add_lshl_u32 v76, v77, v76, 173 ; GCN-NEXT: v_lshl_add_u32 v77, v78, 1, v7674 ; GCN-NEXT: ; implicit-def: $sgpr575 ; GCN-NEXT: v_lshl_add_u32 v78, v79, 1, v7776 ; GCN-NEXT: ; implicit-def: $sgpr277 ; GCN-NEXT: ; implicit-def: $sgpr378 ; GCN-NEXT: v_lshl_add_u32 v79, v80, 1, v7879 ; GCN-NEXT: ; implicit-def: $sgpr0_sgpr180 ; GCN-NEXT: v_mfma_f32_32x32x8_f16 v[0:15], v[44:45], v[40:41], 081 ; GCN-NEXT: v_mfma_f32_32x32x8_f16 v[16:31], v[38:39], v[42:43], v[16:31]82 ; GCN-NEXT: ds_read_b128 v[36:39], v5183 ; GCN-NEXT: s_waitcnt lgkmcnt(0)84 ; GCN-NEXT: buffer_inv sc0 sc185 ; GCN-NEXT: v_mfma_f32_32x32x8_f16 v[0:15], v[46:47], v[42:43], v[0:15]86 ; GCN-NEXT: ds_read_b128 v[44:47], v51 offset:51287 ; GCN-NEXT: s_waitcnt lgkmcnt(0)88 ; GCN-NEXT: buffer_inv sc0 sc189 ; GCN-NEXT: ; implicit-def: $vgpr40_vgpr41_vgpr42_vgpr4390 ; GCN-NEXT: ;;#ASMSTART91 ; GCN-NEXT: s_waitcnt vmcnt(8)92 ; GCN-NEXT: ;;#ASMEND93 ; GCN-NEXT: buffer_wbl2 sc0 sc194 ; GCN-NEXT: ds_write_b128 v48, v[32:35]95 ; GCN-NEXT: v_mfma_f32_32x32x8_f16 v[16:31], v[36:37], v[40:41], v[16:31]96 ; GCN-NEXT: ;;#ASMSTART97 ; GCN-NEXT: s_waitcnt vmcnt(8)98 ; GCN-NEXT: ;;#ASMEND99 ; GCN-NEXT: s_waitcnt lgkmcnt(0)100 ; GCN-NEXT: ds_read_b128 v[32:35], v49101 ; GCN-NEXT: s_waitcnt lgkmcnt(0)102 ; GCN-NEXT: buffer_inv sc0 sc1103 ; GCN-NEXT: v_mfma_f32_32x32x8_f16 v[0:15], v[44:45], v[40:41], v[0:15]104 ; GCN-NEXT: v_mfma_f32_32x32x8_f16 v[16:31], v[38:39], v[42:43], v[16:31]105 ; GCN-NEXT: ; implicit-def: $vgpr36_vgpr37_vgpr38_vgpr39106 ; GCN-NEXT: v_mfma_f32_32x32x8_f16 v[0:15], v[46:47], v[42:43], v[0:15]107 ; GCN-NEXT: ds_read_b128 v[40:43], v49 offset:512108 ; GCN-NEXT: s_waitcnt lgkmcnt(0)109 ; GCN-NEXT: buffer_inv sc0 sc1110 ; GCN-NEXT: ds_read_b128 v[68:71], v51111 ; GCN-NEXT: s_waitcnt lgkmcnt(0)112 ; GCN-NEXT: buffer_inv sc0 sc1113 ; GCN-NEXT: v_mfma_f32_32x32x8_f16 v[16:31], v[32:33], v[36:37], v[16:31]114 ; GCN-NEXT: ; implicit-def: $vgpr32115 ; GCN-NEXT: ; implicit-def: $vgpr33116 ; GCN-NEXT: v_add_u32_e32 v82, v32, v50117 ; GCN-NEXT: v_add_u32_e32 v83, v33, v50118 ; GCN-NEXT: ; kill: killed $vgpr82119 ; GCN-NEXT: ; kill: killed $vgpr83120 ; GCN-NEXT: v_mfma_f32_32x32x8_f16 v[16:31], v[34:35], v[38:39], v[16:31]121 ; GCN-NEXT: v_mfma_f32_32x32x8_f16 v[0:15], v[40:41], v[36:37], v[0:15]122 ; GCN-NEXT: v_mfma_f32_32x32x8_f16 v[16:31], v[68:69], v[62:63], v[16:31]123 ; GCN-NEXT: ds_read_b128 v[66:69], v51 offset:512124 ; GCN-NEXT: s_waitcnt lgkmcnt(0)125 ; GCN-NEXT: buffer_inv sc0 sc1126 ; GCN-NEXT: ;;#ASMSTART127 ; GCN-NEXT: s_waitcnt vmcnt(8)128 ; GCN-NEXT: ;;#ASMEND129 ; GCN-NEXT: v_mfma_f32_32x32x8_f16 v[0:15], v[42:43], v[38:39], v[0:15]130 ; GCN-NEXT: ; implicit-def: $vgpr32_vgpr33_vgpr34_vgpr35_vgpr36_vgpr37_vgpr38_vgpr39_vgpr40_vgpr41_vgpr42_vgpr43_vgpr44_vgpr45_vgpr46_vgpr47131 ; GCN-NEXT: v_mfma_f32_32x32x8_f16 v[0:15], v[66:67], v[62:63], v[0:15]132 ; GCN-NEXT: ; implicit-def: $vgpr66133 ; GCN-NEXT: ; implicit-def: $vgpr67134 ; GCN-NEXT: v_max_f32_e32 v81, v67, v67135 ; GCN-NEXT: ; implicit-def: $vgpr48_vgpr49_vgpr50_vgpr51_vgpr52_vgpr53_vgpr54_vgpr55_vgpr56_vgpr57_vgpr58_vgpr59_vgpr60_vgpr61_vgpr62_vgpr63136 ; GCN-NEXT: v_mfma_f32_32x32x8_f16 v[16:31], v[70:71], v[64:65], v[16:31]137 ; GCN-NEXT: v_perm_b32 v70, v74, v72, s2138 ; GCN-NEXT: v_perm_b32 v71, v74, v72, s3139 ; GCN-NEXT: v_perm_b32 v72, v75, v73, s2140 ; GCN-NEXT: buffer_wbl2 sc0 sc1141 ; GCN-NEXT: ds_write_b32 v76, v70142 ; GCN-NEXT: buffer_wbl2 sc0 sc1143 ; GCN-NEXT: s_waitcnt lgkmcnt(0)144 ; GCN-NEXT: ds_write_b32 v77, v71145 ; GCN-NEXT: buffer_wbl2 sc0 sc1146 ; GCN-NEXT: s_waitcnt lgkmcnt(0)147 ; GCN-NEXT: ds_write_b32 v78, v72148 ; GCN-NEXT: v_mul_f32_e32 v74, s4, v20149 ; GCN-NEXT: v_mfma_f32_32x32x8_f16 v[0:15], v[68:69], v[64:65], v[0:15]150 ; GCN-NEXT: v_mul_f32_e32 v64, s4, v16151 ; GCN-NEXT: v_mul_f32_e32 v65, s4, v17152 ; GCN-NEXT: v_mul_f32_e32 v68, s4, v18153 ; GCN-NEXT: v_mul_f32_e32 v69, s4, v19154 ; GCN-NEXT: v_max3_f32 v64, v64, s5, v65155 ; GCN-NEXT: v_mul_f32_e32 v80, s4, v21156 ; GCN-NEXT: v_max3_f32 v64, v64, v68, v69157 ; GCN-NEXT: v_mul_f32_e32 v84, s4, v22158 ; GCN-NEXT: v_mul_f32_e32 v85, s4, v23159 ; GCN-NEXT: v_max3_f32 v64, v64, v74, v80160 ; GCN-NEXT: v_mul_f32_e32 v86, s4, v24161 ; GCN-NEXT: v_mul_f32_e32 v87, s4, v25162 ; GCN-NEXT: v_max3_f32 v64, v64, v84, v85163 ; GCN-NEXT: v_mul_f32_e32 v65, s4, v26164 ; GCN-NEXT: v_mul_f32_e32 v68, s4, v27165 ; GCN-NEXT: v_max3_f32 v64, v64, v86, v87166 ; GCN-NEXT: v_mul_f32_e32 v69, s4, v28167 ; GCN-NEXT: v_mul_f32_e32 v74, s4, v29168 ; GCN-NEXT: v_max3_f32 v64, v64, v65, v68169 ; GCN-NEXT: v_mul_f32_e32 v80, s4, v30170 ; GCN-NEXT: v_mul_f32_e32 v84, s4, v31171 ; GCN-NEXT: v_max3_f32 v64, v64, v69, v74172 ; GCN-NEXT: v_mul_f32_e32 v85, s4, v0173 ; GCN-NEXT: v_mul_f32_e32 v86, s4, v1174 ; GCN-NEXT: v_max3_f32 v64, v64, v80, v84175 ; GCN-NEXT: v_mul_f32_e32 v87, s4, v2176 ; GCN-NEXT: v_mul_f32_e32 v65, s4, v3177 ; GCN-NEXT: v_max3_f32 v64, v64, v85, v86178 ; GCN-NEXT: v_mul_f32_e32 v68, s4, v4179 ; GCN-NEXT: v_mul_f32_e32 v69, s4, v5180 ; GCN-NEXT: v_max3_f32 v64, v64, v87, v65181 ; GCN-NEXT: v_mul_f32_e32 v74, s4, v6182 ; GCN-NEXT: v_mul_f32_e32 v80, s4, v7183 ; GCN-NEXT: v_max3_f32 v64, v64, v68, v69184 ; GCN-NEXT: v_mul_f32_e32 v84, s4, v8185 ; GCN-NEXT: v_mul_f32_e32 v85, s4, v9186 ; GCN-NEXT: v_max3_f32 v64, v64, v74, v80187 ; GCN-NEXT: v_mul_f32_e32 v86, s4, v10188 ; GCN-NEXT: v_mul_f32_e32 v65, s4, v11189 ; GCN-NEXT: v_max3_f32 v64, v64, v84, v85190 ; GCN-NEXT: v_mul_f32_e32 v87, s4, v12191 ; GCN-NEXT: v_mul_f32_e32 v68, s4, v13192 ; GCN-NEXT: v_max3_f32 v64, v64, v86, v65193 ; GCN-NEXT: v_mul_f32_e32 v69, s4, v14194 ; GCN-NEXT: v_mul_f32_e32 v74, s4, v15195 ; GCN-NEXT: v_max3_f32 v64, v64, v87, v68196 ; GCN-NEXT: v_max3_f32 v64, v64, v69, v74197 ; GCN-NEXT: ds_bpermute_b32 v65, v66, v64198 ; GCN-NEXT: v_perm_b32 v68, v75, v73, s3199 ; GCN-NEXT: buffer_wbl2 sc0 sc1200 ; GCN-NEXT: s_waitcnt lgkmcnt(0)201 ; GCN-NEXT: ds_write_b32 v79, v68202 ; GCN-NEXT: ; implicit-def: $vgpr84203 ; GCN-NEXT: v_max_f32_e32 v65, v65, v65204 ; GCN-NEXT: v_max_f32_e32 v70, v64, v65205 ; GCN-NEXT: s_waitcnt lgkmcnt(0)206 ; GCN-NEXT: buffer_load_dwordx2 v[64:65], v82, s[16:19], 0 offen sc0 sc1207 ; GCN-NEXT: s_waitcnt vmcnt(0)208 ; GCN-NEXT: buffer_inv sc0 sc1209 ; GCN-NEXT: buffer_load_dwordx2 v[68:69], v83, s[16:19], 0 offen sc0 sc1210 ; GCN-NEXT: s_waitcnt vmcnt(0)211 ; GCN-NEXT: buffer_inv sc0 sc1212 ; GCN-NEXT: ds_bpermute_b32 v71, v66, v70213 ; GCN-NEXT: ;;#ASMSTART214 ; GCN-NEXT: s_waitcnt vmcnt(8)215 ; GCN-NEXT: ;;#ASMEND216 ; GCN-NEXT: s_waitcnt lgkmcnt(0)217 ; GCN-NEXT: v_cndmask_b32_e64 v70, v71, v70, s[0:1]218 ; GCN-NEXT: v_max_f32_e32 v70, v70, v70219 ; GCN-NEXT: v_max_f32_e32 v72, v81, v70220 ; GCN-NEXT: v_fma_f32 v16, s4, v16, -v72221 ; GCN-NEXT: v_fma_f32 v18, s4, v18, -v72222 ; GCN-NEXT: v_fma_f32 v19, s4, v19, -v72223 ; GCN-NEXT: v_mul_f32_e32 v16, 0x3fb8aa3b, v16224 ; GCN-NEXT: v_mul_f32_e32 v18, 0x3fb8aa3b, v18225 ; GCN-NEXT: v_mul_f32_e32 v19, 0x3fb8aa3b, v19226 ; GCN-NEXT: v_fma_f32 v17, s4, v17, -v72227 ; GCN-NEXT: v_fma_f32 v20, s4, v20, -v72228 ; GCN-NEXT: v_fma_f32 v21, s4, v21, -v72229 ; GCN-NEXT: v_fma_f32 v22, s4, v22, -v72230 ; GCN-NEXT: v_fma_f32 v23, s4, v23, -v72231 ; GCN-NEXT: v_exp_f32_e32 v73, v16232 ; GCN-NEXT: v_exp_f32_e32 v74, v18233 ; GCN-NEXT: v_exp_f32_e32 v75, v19234 ; GCN-NEXT: v_mul_f32_e32 v20, 0x3fb8aa3b, v20235 ; GCN-NEXT: v_mul_f32_e32 v21, 0x3fb8aa3b, v21236 ; GCN-NEXT: v_mul_f32_e32 v22, 0x3fb8aa3b, v22237 ; GCN-NEXT: v_exp_f32_e32 v80, v20238 ; GCN-NEXT: v_cvt_f16_f32_e32 v16, v73239 ; GCN-NEXT: v_fma_f32 v18, s4, v24, -v72240 ; GCN-NEXT: v_exp_f32_e32 v81, v21241 ; GCN-NEXT: v_cvt_f16_f32_e32 v21, v74242 ; GCN-NEXT: v_fma_f32 v20, s4, v25, -v72243 ; GCN-NEXT: v_exp_f32_e32 v82, v22244 ; GCN-NEXT: v_cvt_f16_f32_e32 v22, v75245 ; GCN-NEXT: v_mul_f32_e32 v17, 0x3fb8aa3b, v17246 ; GCN-NEXT: v_mul_f32_e32 v23, 0x3fb8aa3b, v23247 ; GCN-NEXT: v_fma_f32 v26, s4, v26, -v72248 ; GCN-NEXT: v_pack_b32_f16 v71, v21, v22249 ; GCN-NEXT: v_mul_f32_e32 v22, 0x3fb8aa3b, v18250 ; GCN-NEXT: v_sub_f32_e32 v24, v67, v72251 ; GCN-NEXT: v_exp_f32_e32 v83, v23252 ; GCN-NEXT: v_fma_f32 v67, s4, v27, -v72253 ; GCN-NEXT: v_exp_f32_e32 v85, v22254 ; GCN-NEXT: v_exp_f32_e32 v17, v17255 ; GCN-NEXT: v_mul_f32_e32 v24, 0x3fb8aa3b, v24256 ; GCN-NEXT: v_mul_f32_e32 v23, 0x3fb8aa3b, v20257 ; GCN-NEXT: v_cvt_f16_f32_e32 v19, v17258 ; GCN-NEXT: v_fma_f32 v87, s4, v29, -v72259 ; GCN-NEXT: v_exp_f32_e32 v88, v23260 ; GCN-NEXT: v_fma_f32 v0, s4, v0, -v72261 ; GCN-NEXT: v_pack_b32_f16 v70, v16, v19262 ; GCN-NEXT: ds_read_b128 v[18:21], v84263 ; GCN-NEXT: s_waitcnt lgkmcnt(0)264 ; GCN-NEXT: buffer_inv sc0 sc1265 ; GCN-NEXT: v_exp_f32_e32 v16, v24266 ; GCN-NEXT: ds_read_b128 v[22:25], v84 offset:576267 ; GCN-NEXT: s_waitcnt lgkmcnt(0)268 ; GCN-NEXT: buffer_inv sc0 sc1269 ; GCN-NEXT: v_pk_mul_f32 v[48:49], v[48:49], v[16:17] op_sel_hi:[1,0]270 ; GCN-NEXT: v_pk_mul_f32 v[50:51], v[50:51], v[16:17] op_sel_hi:[1,0]271 ; GCN-NEXT: v_pk_mul_f32 v[52:53], v[52:53], v[16:17] op_sel_hi:[1,0]272 ; GCN-NEXT: v_pk_mul_f32 v[54:55], v[54:55], v[16:17] op_sel_hi:[1,0]273 ; GCN-NEXT: v_pk_mul_f32 v[56:57], v[56:57], v[16:17] op_sel_hi:[1,0]274 ; GCN-NEXT: v_pk_mul_f32 v[58:59], v[58:59], v[16:17] op_sel_hi:[1,0]275 ; GCN-NEXT: v_pk_mul_f32 v[60:61], v[60:61], v[16:17] op_sel_hi:[1,0]276 ; GCN-NEXT: v_pk_mul_f32 v[62:63], v[62:63], v[16:17] op_sel_hi:[1,0]277 ; GCN-NEXT: v_pk_mul_f32 v[32:33], v[32:33], v[16:17] op_sel_hi:[1,0]278 ; GCN-NEXT: v_pk_mul_f32 v[34:35], v[34:35], v[16:17] op_sel_hi:[1,0]279 ; GCN-NEXT: v_pk_mul_f32 v[36:37], v[36:37], v[16:17] op_sel_hi:[1,0]280 ; GCN-NEXT: v_pk_mul_f32 v[38:39], v[38:39], v[16:17] op_sel_hi:[1,0]281 ; GCN-NEXT: v_pk_mul_f32 v[40:41], v[40:41], v[16:17] op_sel_hi:[1,0]282 ; GCN-NEXT: v_pk_mul_f32 v[42:43], v[42:43], v[16:17] op_sel_hi:[1,0]283 ; GCN-NEXT: v_pk_mul_f32 v[44:45], v[44:45], v[16:17] op_sel_hi:[1,0]284 ; GCN-NEXT: v_pk_mul_f32 v[46:47], v[46:47], v[16:17] op_sel_hi:[1,0]285 ; GCN-NEXT: v_mfma_f32_32x32x8_f16 v[48:63], v[18:19], v[70:71], v[48:63]286 ; GCN-NEXT: v_add_f32_e32 v18, 0, v73287 ; GCN-NEXT: v_cvt_f16_f32_e32 v89, v83288 ; GCN-NEXT: v_fma_f32 v73, s4, v28, -v72289 ; GCN-NEXT: v_cvt_f16_f32_e32 v19, v80290 ; GCN-NEXT: v_fma_f32 v1, s4, v1, -v72291 ; GCN-NEXT: v_perm_b32 v90, v69, v65, s2292 ; GCN-NEXT: v_mfma_f32_32x32x8_f16 v[32:47], v[22:23], v[70:71], v[32:47]293 ; GCN-NEXT: v_add_f32_e32 v17, v17, v18294 ; GCN-NEXT: v_mul_f32_e32 v18, 0x3fb8aa3b, v26295 ; GCN-NEXT: v_cvt_f16_f32_e32 v86, v81296 ; GCN-NEXT: v_fma_f32 v23, s4, v30, -v72297 ; GCN-NEXT: v_exp_f32_e32 v30, v18298 ; GCN-NEXT: v_cvt_f16_f32_e32 v22, v82299 ; GCN-NEXT: v_fma_f32 v18, s4, v31, -v72300 ; GCN-NEXT: v_perm_b32 v31, v68, v64, s2301 ; GCN-NEXT: v_perm_b32 v64, v68, v64, s3302 ; GCN-NEXT: v_perm_b32 v65, v69, v65, s3303 ; GCN-NEXT: ds_read_b128 v[26:29], v91304 ; GCN-NEXT: s_waitcnt lgkmcnt(0)305 ; GCN-NEXT: buffer_inv sc0 sc1306 ; GCN-NEXT: ds_read_b128 v[68:71], v91 offset:576307 ; GCN-NEXT: s_waitcnt lgkmcnt(0)308 ; GCN-NEXT: buffer_inv sc0 sc1309 ; GCN-NEXT: ;;#ASMSTART310 ; GCN-NEXT: s_waitcnt vmcnt(8)311 ; GCN-NEXT: ;;#ASMEND312 ; GCN-NEXT: buffer_wbl2 sc0 sc1313 ; GCN-NEXT: ds_write_b32 v76, v31314 ; GCN-NEXT: v_mul_f32_e32 v31, 0x3fb8aa3b, v67315 ; GCN-NEXT: v_exp_f32_e32 v31, v31316 ; GCN-NEXT: v_mul_f32_e32 v67, 0x3fb8aa3b, v18317 ; GCN-NEXT: v_pack_b32_f16 v18, v19, v86318 ; GCN-NEXT: v_pack_b32_f16 v19, v22, v89319 ; GCN-NEXT: buffer_wbl2 sc0 sc1320 ; GCN-NEXT: s_waitcnt lgkmcnt(0)321 ; GCN-NEXT: ds_write_b32 v77, v64322 ; GCN-NEXT: buffer_wbl2 sc0 sc1323 ; GCN-NEXT: s_waitcnt lgkmcnt(0)324 ; GCN-NEXT: ds_write_b32 v78, v90325 ; GCN-NEXT: buffer_wbl2 sc0 sc1326 ; GCN-NEXT: s_waitcnt lgkmcnt(0)327 ; GCN-NEXT: ds_write_b32 v79, v65328 ; GCN-NEXT: v_mul_f32_e32 v64, 0x3fb8aa3b, v73329 ; GCN-NEXT: v_mul_f32_e32 v65, 0x3fb8aa3b, v87330 ; GCN-NEXT: v_mfma_f32_32x32x8_f16 v[48:63], v[20:21], v[18:19], v[48:63]331 ; GCN-NEXT: v_add_f32_e32 v17, v74, v17332 ; GCN-NEXT: v_cvt_f16_f32_e32 v20, v85333 ; GCN-NEXT: v_fma_f32 v2, s4, v2, -v72334 ; GCN-NEXT: v_exp_f32_e32 v22, v64335 ; GCN-NEXT: v_cvt_f16_f32_e32 v21, v88336 ; GCN-NEXT: v_exp_f32_e32 v64, v65337 ; GCN-NEXT: v_mul_f32_e32 v23, 0x3fb8aa3b, v23338 ; GCN-NEXT: v_mfma_f32_32x32x8_f16 v[32:47], v[24:25], v[18:19], v[32:47]339 ; GCN-NEXT: v_add_f32_e32 v17, v75, v17340 ; GCN-NEXT: v_cvt_f16_f32_e32 v18, v30341 ; GCN-NEXT: v_fma_f32 v24, s4, v3, -v72342 ; GCN-NEXT: v_exp_f32_e32 v23, v23343 ; GCN-NEXT: v_cvt_f16_f32_e32 v19, v31344 ; GCN-NEXT: v_mul_f32_e32 v3, 0x3fb8aa3b, v0345 ; GCN-NEXT: v_mul_f32_e32 v65, 0x3fb8aa3b, v1346 ; GCN-NEXT: v_pack_b32_f16 v0, v20, v21347 ; GCN-NEXT: v_pack_b32_f16 v1, v18, v19348 ; GCN-NEXT: v_fma_f32 v6, s4, v6, -v72349 ; GCN-NEXT: v_exp_f32_e32 v25, v67350 ; GCN-NEXT: v_mfma_f32_32x32x8_f16 v[48:63], v[26:27], v[0:1], v[48:63]351 ; GCN-NEXT: v_add_f32_e32 v17, v80, v17352 ; GCN-NEXT: v_cvt_f16_f32_e32 v18, v22353 ; GCN-NEXT: v_fma_f32 v26, s4, v4, -v72354 ; GCN-NEXT: v_exp_f32_e32 v27, v3355 ; GCN-NEXT: v_cvt_f16_f32_e32 v4, v64356 ; GCN-NEXT: v_fma_f32 v67, s4, v5, -v72357 ; GCN-NEXT: v_exp_f32_e32 v65, v65358 ; GCN-NEXT: v_mfma_f32_32x32x8_f16 v[32:47], v[68:69], v[0:1], v[32:47]359 ; GCN-NEXT: v_mul_f32_e32 v2, 0x3fb8aa3b, v2360 ; GCN-NEXT: v_add_f32_e32 v17, v81, v17361 ; GCN-NEXT: v_cvt_f16_f32_e32 v5, v23362 ; GCN-NEXT: v_fma_f32 v7, s4, v7, -v72363 ; GCN-NEXT: v_exp_f32_e32 v68, v2364 ; GCN-NEXT: v_cvt_f16_f32_e32 v19, v25365 ; GCN-NEXT: ;;#ASMSTART366 ; GCN-NEXT: s_waitcnt vmcnt(8)367 ; GCN-NEXT: ;;#ASMEND368 ; GCN-NEXT: v_mul_f32_e32 v24, 0x3fb8aa3b, v24369 ; GCN-NEXT: s_waitcnt lgkmcnt(0)370 ; GCN-NEXT: ds_read_b128 v[0:3], v84371 ; GCN-NEXT: s_waitcnt lgkmcnt(0)372 ; GCN-NEXT: buffer_inv sc0 sc1373 ; GCN-NEXT: v_pack_b32_f16 v4, v18, v4374 ; GCN-NEXT: v_pack_b32_f16 v5, v5, v19375 ; GCN-NEXT: v_exp_f32_e32 v24, v24376 ; GCN-NEXT: ds_read_b128 v[18:21], v84 offset:576377 ; GCN-NEXT: s_waitcnt lgkmcnt(0)378 ; GCN-NEXT: buffer_inv sc0 sc1379 ; GCN-NEXT: v_mul_f32_e32 v26, 0x3fb8aa3b, v26380 ; GCN-NEXT: v_mul_f32_e32 v67, 0x3fb8aa3b, v67381 ; GCN-NEXT: v_mfma_f32_32x32x8_f16 v[48:63], v[28:29], v[4:5], v[48:63]382 ; GCN-NEXT: v_add_f32_e32 v17, v82, v17383 ; GCN-NEXT: v_cvt_f16_f32_e32 v28, v27384 ; GCN-NEXT: v_exp_f32_e32 v26, v26385 ; GCN-NEXT: v_cvt_f16_f32_e32 v29, v65386 ; GCN-NEXT: v_fma_f32 v10, s4, v10, -v72387 ; GCN-NEXT: v_exp_f32_e32 v67, v67388 ; GCN-NEXT: v_mul_f32_e32 v6, 0x3fb8aa3b, v6389 ; GCN-NEXT: v_mfma_f32_32x32x8_f16 v[32:47], v[70:71], v[4:5], v[32:47]390 ; GCN-NEXT: v_add_f32_e32 v17, v83, v17391 ; GCN-NEXT: v_cvt_f16_f32_e32 v5, v68392 ; GCN-NEXT: v_exp_f32_e32 v6, v6393 ; GCN-NEXT: v_cvt_f16_f32_e32 v69, v24394 ; GCN-NEXT: v_mul_f32_e32 v7, 0x3fb8aa3b, v7395 ; GCN-NEXT: v_exp_f32_e32 v7, v7396 ; GCN-NEXT: v_pack_b32_f16 v4, v28, v29397 ; GCN-NEXT: v_pack_b32_f16 v5, v5, v69398 ; GCN-NEXT: ; implicit-def: $sgpr2399 ; GCN-NEXT: s_nop 1400 ; GCN-NEXT: v_mfma_f32_32x32x8_f16 v[48:63], v[0:1], v[4:5], v[48:63]401 ; GCN-NEXT: v_add_f32_e32 v0, v85, v17402 ; GCN-NEXT: v_cvt_f16_f32_e32 v17, v26403 ; GCN-NEXT: v_cvt_f16_f32_e32 v28, v67404 ; GCN-NEXT: v_mfma_f32_32x32x8_f16 v[32:47], v[18:19], v[4:5], v[32:47]405 ; GCN-NEXT: v_add_f32_e32 v4, v88, v0406 ; GCN-NEXT: v_mul_f32_e32 v0, 0x3fb8aa3b, v10407 ; GCN-NEXT: v_cvt_f16_f32_e32 v1, v6408 ; GCN-NEXT: v_exp_f32_e32 v10, v0409 ; GCN-NEXT: v_cvt_f16_f32_e32 v0, v7410 ; GCN-NEXT: v_pack_b32_f16 v1, v1, v0411 ; GCN-NEXT: v_pack_b32_f16 v0, v17, v28412 ; GCN-NEXT: s_nop 1413 ; GCN-NEXT: v_mfma_f32_32x32x8_f16 v[48:63], v[2:3], v[0:1], v[48:63]414 ; GCN-NEXT: v_add_f32_e32 v2, v30, v4415 ; GCN-NEXT: v_mfma_f32_32x32x8_f16 v[32:47], v[20:21], v[0:1], v[32:47]416 ; GCN-NEXT: v_add_f32_e32 v0, v31, v2417 ; GCN-NEXT: v_add_f32_e32 v0, v22, v0418 ; GCN-NEXT: v_add_f32_e32 v0, v64, v0419 ; GCN-NEXT: v_add_f32_e32 v0, v23, v0420 ; GCN-NEXT: v_add_f32_e32 v0, v25, v0421 ; GCN-NEXT: v_add_f32_e32 v0, v27, v0422 ; GCN-NEXT: v_fma_f32 v8, s4, v8, -v72423 ; GCN-NEXT: v_add_f32_e32 v0, v65, v0424 ; GCN-NEXT: v_fma_f32 v9, s4, v9, -v72425 ; GCN-NEXT: v_mul_f32_e32 v8, 0x3fb8aa3b, v8426 ; GCN-NEXT: v_add_f32_e32 v0, v68, v0427 ; GCN-NEXT: v_fma_f32 v11, s4, v11, -v72428 ; GCN-NEXT: v_mul_f32_e32 v9, 0x3fb8aa3b, v9429 ; GCN-NEXT: v_fma_f32 v12, s4, v12, -v72430 ; GCN-NEXT: v_fma_f32 v13, s4, v13, -v72431 ; GCN-NEXT: v_exp_f32_e32 v8, v8432 ; GCN-NEXT: v_add_f32_e32 v0, v24, v0433 ; GCN-NEXT: v_fma_f32 v5, s4, v14, -v72434 ; GCN-NEXT: v_exp_f32_e32 v9, v9435 ; GCN-NEXT: v_add_f32_e32 v0, v26, v0436 ; GCN-NEXT: v_add_f32_e32 v0, v67, v0437 ; GCN-NEXT: v_fma_f32 v14, s4, v15, -v72438 ; GCN-NEXT: v_mul_f32_e32 v11, 0x3fb8aa3b, v11439 ; GCN-NEXT: v_mul_f32_e32 v3, 0x3fb8aa3b, v12440 ; GCN-NEXT: v_mul_f32_e32 v1, 0x3fb8aa3b, v5441 ; GCN-NEXT: v_add_f32_e32 v0, v6, v0442 ; GCN-NEXT: v_exp_f32_e32 v11, v11443 ; GCN-NEXT: v_cvt_f16_f32_e32 v4, v8444 ; GCN-NEXT: v_exp_f32_e32 v12, v3445 ; GCN-NEXT: v_mul_f32_e32 v3, 0x3fb8aa3b, v13446 ; GCN-NEXT: v_exp_f32_e32 v17, v1447 ; GCN-NEXT: v_mul_f32_e32 v1, 0x3fb8aa3b, v14448 ; GCN-NEXT: v_add_f32_e32 v0, v7, v0449 ; GCN-NEXT: v_cvt_f16_f32_e32 v13, v9450 ; GCN-NEXT: v_exp_f32_e32 v15, v3451 ; GCN-NEXT: v_exp_f32_e32 v18, v1452 ; GCN-NEXT: v_add_f32_e32 v6, v8, v0453 ; GCN-NEXT: ds_read_b128 v[0:3], v91454 ; GCN-NEXT: s_waitcnt lgkmcnt(0)455 ; GCN-NEXT: buffer_inv sc0 sc1456 ; GCN-NEXT: v_cvt_f16_f32_e32 v5, v10457 ; GCN-NEXT: v_cvt_f16_f32_e32 v14, v11458 ; GCN-NEXT: v_add_f32_e32 v6, v9, v6459 ; GCN-NEXT: v_pack_b32_f16 v8, v4, v13460 ; GCN-NEXT: v_add_f32_e32 v6, v10, v6461 ; GCN-NEXT: v_pack_b32_f16 v9, v5, v14462 ; GCN-NEXT: v_cvt_f16_f32_e32 v7, v18463 ; GCN-NEXT: v_cvt_f16_f32_e32 v10, v15464 ; GCN-NEXT: v_mfma_f32_32x32x8_f16 v[48:63], v[0:1], v[8:9], v[48:63]465 ; GCN-NEXT: v_cvt_f16_f32_e32 v0, v17466 ; GCN-NEXT: v_cvt_f16_f32_e32 v4, v12467 ; GCN-NEXT: v_add_f32_e32 v6, v11, v6468 ; GCN-NEXT: v_add_f32_e32 v6, v12, v6469 ; GCN-NEXT: v_add_f32_e32 v1, v15, v6470 ; GCN-NEXT: v_add_f32_e32 v11, v17, v1471 ; GCN-NEXT: v_pack_b32_f16 v1, v0, v7472 ; GCN-NEXT: v_pack_b32_f16 v0, v4, v10473 ; GCN-NEXT: ds_read_b128 v[4:7], v91 offset:576474 ; GCN-NEXT: s_waitcnt lgkmcnt(0)475 ; GCN-NEXT: buffer_inv sc0 sc1476 ; GCN-NEXT: v_mfma_f32_32x32x8_f16 v[32:47], v[4:5], v[8:9], v[32:47]477 ; GCN-NEXT: ;;#ASMSTART478 ; GCN-NEXT: s_waitcnt vmcnt(8)479 ; GCN-NEXT: ;;#ASMEND480 ; GCN-NEXT: v_mov_b32_e32 v4, 0481 ; GCN-NEXT: v_mfma_f32_32x32x8_f16 v[32:47], v[6:7], v[0:1], v[32:47]482 ; GCN-NEXT: v_mfma_f32_32x32x8_f16 v[48:63], v[2:3], v[0:1], v[48:63]483 ; GCN-NEXT: v_add_f32_e32 v2, v18, v11484 ; GCN-NEXT: ds_bpermute_b32 v3, v66, v2485 ; GCN-NEXT: s_waitcnt lgkmcnt(0)486 ; GCN-NEXT: v_add_f32_e32 v2, v2, v3487 ; GCN-NEXT: ds_bpermute_b32 v3, v66, v2488 ; GCN-NEXT: s_waitcnt lgkmcnt(0)489 ; GCN-NEXT: v_cndmask_b32_e64 v2, v3, v2, s[0:1]490 ; GCN-NEXT: v_fmac_f32_e32 v2, v4, v16491 ; GCN-NEXT: s_endpgm492 attributes #0 = {"amdgpu-flat-work-group-size"="256,256"}493 494 !0 = !{i64 2862105}495...496 497---498name: smallInterleave499tracksRegLiveness: true500machineFunctionInfo:501 stackPtrOffsetReg: '$sgpr32'502body: |503 bb.0:504 liveins: $vgpr0, $sgpr0_sgpr1, $sgpr2, $sgpr3, $sgpr4505 %0:vgpr_32 = V_MOV_B32_e32 0, implicit $exec506 %1:vgpr_32 = COPY %0:vgpr_32507 %2:vgpr_32 = IMPLICIT_DEF508 %3:sreg_32 = IMPLICIT_DEF509 %4:vreg_64_align2 = IMPLICIT_DEF510 %5:sgpr_128 = IMPLICIT_DEF511 %6:vgpr_32 = IMPLICIT_DEF512 %7:vgpr_32 = IMPLICIT_DEF513 %8:sgpr_128 = IMPLICIT_DEF514 %9:vgpr_32 = IMPLICIT_DEF515 %10:sgpr_512 = IMPLICIT_DEF516 %11:sgpr_32 = IMPLICIT_DEF517 %12:sreg_64_xexec = IMPLICIT_DEF518 %13:vgpr_32 = IMPLICIT_DEF519 %14:sreg_32 = IMPLICIT_DEF520 %15:sreg_32 = IMPLICIT_DEF521 %16:vgpr_32 = IMPLICIT_DEF522 %17:sreg_32 = IMPLICIT_DEF523 %18:vgpr_32 = IMPLICIT_DEF524 %19:vgpr_32 = IMPLICIT_DEF525 %20:vgpr_32 = IMPLICIT_DEF526 %21:vgpr_32 = IMPLICIT_DEF527 %22:vgpr_32 = IMPLICIT_DEF528 %23:vgpr_32 = IMPLICIT_DEF529 %24:vgpr_32 = IMPLICIT_DEF530 %25:vgpr_32 = IMPLICIT_DEF531 %26:sreg_32 = IMPLICIT_DEF532 %42:vgpr_32 = IMPLICIT_DEF533 %44:vreg_128_align2 = IMPLICIT_DEF534 %48:vgpr_32 = IMPLICIT_DEF535 %49:vreg_128_align2 = IMPLICIT_DEF536 %52:vreg_128_align2 = IMPLICIT_DEF537 %55:vreg_128_align2 = IMPLICIT_DEF538 %106:vgpr_32 = IMPLICIT_DEF539 %29:vgpr_32 = IMPLICIT_DEF540 %37:vgpr_32 = IMPLICIT_DEF541 %259:vreg_512_align2 = IMPLICIT_DEF542 %260:vreg_512_align2 = IMPLICIT_DEF543 IGLP_OPT 2544 %27:sreg_32_xm0 = V_READFIRSTLANE_B32 %2:vgpr_32, implicit $exec545 %28:vgpr_32 = V_LSHL_ADD_U32_e64 %27:sreg_32_xm0, 4, %29:vgpr_32, implicit $exec546 %30:vreg_64_align2, dead %31:sreg_64 = V_MAD_U64_U32_e64 %3:sreg_32, %28:vgpr_32, %4:vreg_64_align2, 0, implicit $exec547 %32:vreg_128_align2 = BUFFER_LOAD_DWORDX4_OFFEN %30.sub0:vreg_64_align2, %5:sgpr_128, 0, 0, 0, 0, implicit $exec548 %33:sreg_32 = S_LSHL_B32 %27:sreg_32_xm0, 7, implicit-def dead $scc549 %34:vgpr_32 = V_ADD_LSHL_U32_e64 %6:vgpr_32, %33:sreg_32, 1, implicit $exec550 DS_WRITE_B128_gfx9 %34:vgpr_32, %32:vreg_128_align2, 0, 0, implicit $exec551 %35:vreg_128_align2 = BUFFER_LOAD_DWORDX4_OFFEN %30.sub0:vreg_64_align2, %5:sgpr_128, 0, 64, 0, 0, implicit $exec552 %36:vgpr_32 = V_ADD_U32_e32 %7:vgpr_32, %37:vgpr_32, implicit $exec553 %38:vreg_64_align2 = BUFFER_LOAD_DWORDX2_OFFEN %36:vgpr_32, %8:sgpr_128, 0, 0, 0, 0, implicit $exec554 %39:vgpr_32 = V_ADD_U32_e32 %9:vgpr_32, %37:vgpr_32, implicit $exec555 %40:vreg_64_align2 = BUFFER_LOAD_DWORDX2_OFFEN %39:vgpr_32, %8:sgpr_128, 0, 0, 0, 0, implicit $exec556 INLINEASM &"s_waitcnt vmcnt($0)", 57 /* sideeffect mayload maystore isconvergent attdialect */, 13 /* imm */, 8, !0557 %41:vreg_128_align2 = DS_READ_B128_gfx9 %42:vgpr_32, 0, 0, implicit $exec558 early-clobber %43:vreg_512_align2 = contract V_MFMA_F32_32X32X8F16_vgprcd_e64 %41.sub0_sub1:vreg_128_align2, %44.sub0_sub1:vreg_128_align2, 0, 0, 0, 0, implicit $mode, implicit $exec559 %43:vreg_512_align2 = contract V_MFMA_F32_32X32X8F16_mac_vgprcd_e64 %41.sub2_sub3:vreg_128_align2, %44.sub2_sub3:vreg_128_align2, %43:vreg_512_align2, 0, 0, 0, implicit $mode, implicit $exec560 %45:vreg_128_align2 = DS_READ_B128_gfx9 %42:vgpr_32, 512, 0, implicit $exec561 early-clobber %46:vreg_512_align2 = contract V_MFMA_F32_32X32X8F16_vgprcd_e64 %45.sub0_sub1:vreg_128_align2, %44.sub0_sub1:vreg_128_align2, 0, 0, 0, 0, implicit $mode, implicit $exec562 %46:vreg_512_align2 = contract V_MFMA_F32_32X32X8F16_mac_vgprcd_e64 %45.sub2_sub3:vreg_128_align2, %44.sub2_sub3:vreg_128_align2, %46:vreg_512_align2, 0, 0, 0, implicit $mode, implicit $exec563 %47:vreg_128_align2 = DS_READ_B128_gfx9 %48:vgpr_32, 0, 0, implicit $exec564 %43:vreg_512_align2 = contract V_MFMA_F32_32X32X8F16_mac_vgprcd_e64 %47.sub0_sub1:vreg_128_align2, %49.sub0_sub1:vreg_128_align2, %43:vreg_512_align2, 0, 0, 0, implicit $mode, implicit $exec565 %43:vreg_512_align2 = contract V_MFMA_F32_32X32X8F16_mac_vgprcd_e64 %47.sub2_sub3:vreg_128_align2, %49.sub2_sub3:vreg_128_align2, %43:vreg_512_align2, 0, 0, 0, implicit $mode, implicit $exec566 %50:vreg_128_align2 = DS_READ_B128_gfx9 %48:vgpr_32, 512, 0, implicit $exec567 %46:vreg_512_align2 = contract V_MFMA_F32_32X32X8F16_mac_vgprcd_e64 %50.sub0_sub1:vreg_128_align2, %49.sub0_sub1:vreg_128_align2, %46:vreg_512_align2, 0, 0, 0, implicit $mode, implicit $exec568 %46:vreg_512_align2 = contract V_MFMA_F32_32X32X8F16_mac_vgprcd_e64 %50.sub2_sub3:vreg_128_align2, %49.sub2_sub3:vreg_128_align2, %46:vreg_512_align2, 0, 0, 0, implicit $mode, implicit $exec569 INLINEASM &"s_waitcnt vmcnt($0)", 57 /* sideeffect mayload maystore isconvergent attdialect */, 13 /* imm */, 8, !0570 DS_WRITE_B128_gfx9 %34:vgpr_32, %35:vreg_128_align2, 0, 0, implicit $exec571 INLINEASM &"s_waitcnt vmcnt($0)", 57 /* sideeffect mayload maystore isconvergent attdialect */, 13 /* imm */, 8, !0572 %51:vreg_128_align2 = DS_READ_B128_gfx9 %42:vgpr_32, 0, 0, implicit $exec573 %43:vreg_512_align2 = contract V_MFMA_F32_32X32X8F16_mac_vgprcd_e64 %51.sub0_sub1:vreg_128_align2, %52.sub0_sub1:vreg_128_align2, %43:vreg_512_align2, 0, 0, 0, implicit $mode, implicit $exec574 %43:vreg_512_align2 = contract V_MFMA_F32_32X32X8F16_mac_vgprcd_e64 %51.sub2_sub3:vreg_128_align2, %52.sub2_sub3:vreg_128_align2, %43:vreg_512_align2, 0, 0, 0, implicit $mode, implicit $exec575 %53:vreg_128_align2 = DS_READ_B128_gfx9 %42:vgpr_32, 512, 0, implicit $exec576 %46:vreg_512_align2 = contract V_MFMA_F32_32X32X8F16_mac_vgprcd_e64 %53.sub0_sub1:vreg_128_align2, %52.sub0_sub1:vreg_128_align2, %46:vreg_512_align2, 0, 0, 0, implicit $mode, implicit $exec577 %46:vreg_512_align2 = contract V_MFMA_F32_32X32X8F16_mac_vgprcd_e64 %53.sub2_sub3:vreg_128_align2, %52.sub2_sub3:vreg_128_align2, %46:vreg_512_align2, 0, 0, 0, implicit $mode, implicit $exec578 %54:vreg_128_align2 = DS_READ_B128_gfx9 %48:vgpr_32, 0, 0, implicit $exec579 %43:vreg_512_align2 = contract V_MFMA_F32_32X32X8F16_mac_vgprcd_e64 %54.sub0_sub1:vreg_128_align2, %55.sub0_sub1:vreg_128_align2, %43:vreg_512_align2, 0, 0, 0, implicit $mode, implicit $exec580 %43:vreg_512_align2 = contract V_MFMA_F32_32X32X8F16_mac_vgprcd_e64 %54.sub2_sub3:vreg_128_align2, %55.sub2_sub3:vreg_128_align2, %43:vreg_512_align2, 0, 0, 0, implicit $mode, implicit $exec581 %56:vreg_128_align2 = DS_READ_B128_gfx9 %48:vgpr_32, 512, 0, implicit $exec582 %46:vreg_512_align2 = contract V_MFMA_F32_32X32X8F16_mac_vgprcd_e64 %56.sub0_sub1:vreg_128_align2, %55.sub0_sub1:vreg_128_align2, %46:vreg_512_align2, 0, 0, 0, implicit $mode, implicit $exec583 %46:vreg_512_align2 = contract V_MFMA_F32_32X32X8F16_mac_vgprcd_e64 %56.sub2_sub3:vreg_128_align2, %55.sub2_sub3:vreg_128_align2, %46:vreg_512_align2, 0, 0, 0, implicit $mode, implicit $exec584 %57:vgpr_32 = contract nofpexcept V_MUL_F32_e32 %10.sub4:sgpr_512, %43.sub0:vreg_512_align2, implicit $mode, implicit $exec585 %58:vgpr_32 = contract nofpexcept V_MUL_F32_e32 %10.sub4:sgpr_512, %43.sub1:vreg_512_align2, implicit $mode, implicit $exec586 %59:vgpr_32 = contract nofpexcept V_MUL_F32_e32 %10.sub4:sgpr_512, %43.sub2:vreg_512_align2, implicit $mode, implicit $exec587 %60:vgpr_32 = contract nofpexcept V_MUL_F32_e32 %10.sub4:sgpr_512, %43.sub3:vreg_512_align2, implicit $mode, implicit $exec588 %61:vgpr_32 = contract nofpexcept V_MUL_F32_e32 %10.sub4:sgpr_512, %43.sub4:vreg_512_align2, implicit $mode, implicit $exec589 %62:vgpr_32 = contract nofpexcept V_MUL_F32_e32 %10.sub4:sgpr_512, %43.sub5:vreg_512_align2, implicit $mode, implicit $exec590 %63:vgpr_32 = contract nofpexcept V_MUL_F32_e32 %10.sub4:sgpr_512, %43.sub6:vreg_512_align2, implicit $mode, implicit $exec591 %64:vgpr_32 = contract nofpexcept V_MUL_F32_e32 %10.sub4:sgpr_512, %43.sub7:vreg_512_align2, implicit $mode, implicit $exec592 %65:vgpr_32 = contract nofpexcept V_MUL_F32_e32 %10.sub4:sgpr_512, %43.sub8:vreg_512_align2, implicit $mode, implicit $exec593 %66:vgpr_32 = contract nofpexcept V_MUL_F32_e32 %10.sub4:sgpr_512, %43.sub9:vreg_512_align2, implicit $mode, implicit $exec594 %67:vgpr_32 = contract nofpexcept V_MUL_F32_e32 %10.sub4:sgpr_512, %43.sub10:vreg_512_align2, implicit $mode, implicit $exec595 %68:vgpr_32 = contract nofpexcept V_MUL_F32_e32 %10.sub4:sgpr_512, %43.sub11:vreg_512_align2, implicit $mode, implicit $exec596 %69:vgpr_32 = contract nofpexcept V_MUL_F32_e32 %10.sub4:sgpr_512, %43.sub12:vreg_512_align2, implicit $mode, implicit $exec597 %70:vgpr_32 = contract nofpexcept V_MUL_F32_e32 %10.sub4:sgpr_512, %43.sub13:vreg_512_align2, implicit $mode, implicit $exec598 %71:vgpr_32 = contract nofpexcept V_MUL_F32_e32 %10.sub4:sgpr_512, %43.sub14:vreg_512_align2, implicit $mode, implicit $exec599 %72:vgpr_32 = contract nofpexcept V_MUL_F32_e32 %10.sub4:sgpr_512, %43.sub15:vreg_512_align2, implicit $mode, implicit $exec600 %73:vgpr_32 = contract nofpexcept V_MUL_F32_e32 %10.sub4:sgpr_512, %46.sub0:vreg_512_align2, implicit $mode, implicit $exec601 %74:vgpr_32 = contract nofpexcept V_MUL_F32_e32 %10.sub4:sgpr_512, %46.sub1:vreg_512_align2, implicit $mode, implicit $exec602 %75:vgpr_32 = contract nofpexcept V_MUL_F32_e32 %10.sub4:sgpr_512, %46.sub2:vreg_512_align2, implicit $mode, implicit $exec603 %76:vgpr_32 = contract nofpexcept V_MUL_F32_e32 %10.sub4:sgpr_512, %46.sub3:vreg_512_align2, implicit $mode, implicit $exec604 %77:vgpr_32 = contract nofpexcept V_MUL_F32_e32 %10.sub4:sgpr_512, %46.sub4:vreg_512_align2, implicit $mode, implicit $exec605 %78:vgpr_32 = contract nofpexcept V_MUL_F32_e32 %10.sub4:sgpr_512, %46.sub5:vreg_512_align2, implicit $mode, implicit $exec606 %79:vgpr_32 = contract nofpexcept V_MUL_F32_e32 %10.sub4:sgpr_512, %46.sub6:vreg_512_align2, implicit $mode, implicit $exec607 %80:vgpr_32 = contract nofpexcept V_MUL_F32_e32 %10.sub4:sgpr_512, %46.sub7:vreg_512_align2, implicit $mode, implicit $exec608 %81:vgpr_32 = contract nofpexcept V_MUL_F32_e32 %10.sub4:sgpr_512, %46.sub8:vreg_512_align2, implicit $mode, implicit $exec609 %82:vgpr_32 = contract nofpexcept V_MUL_F32_e32 %10.sub4:sgpr_512, %46.sub9:vreg_512_align2, implicit $mode, implicit $exec610 %83:vgpr_32 = contract nofpexcept V_MUL_F32_e32 %10.sub4:sgpr_512, %46.sub10:vreg_512_align2, implicit $mode, implicit $exec611 %84:vgpr_32 = contract nofpexcept V_MUL_F32_e32 %10.sub4:sgpr_512, %46.sub11:vreg_512_align2, implicit $mode, implicit $exec612 %85:vgpr_32 = contract nofpexcept V_MUL_F32_e32 %10.sub4:sgpr_512, %46.sub12:vreg_512_align2, implicit $mode, implicit $exec613 %86:vgpr_32 = contract nofpexcept V_MUL_F32_e32 %10.sub4:sgpr_512, %46.sub13:vreg_512_align2, implicit $mode, implicit $exec614 %87:vgpr_32 = contract nofpexcept V_MUL_F32_e32 %10.sub4:sgpr_512, %46.sub14:vreg_512_align2, implicit $mode, implicit $exec615 %88:vgpr_32 = contract nofpexcept V_MUL_F32_e32 %10.sub4:sgpr_512, %46.sub15:vreg_512_align2, implicit $mode, implicit $exec616 %89:vgpr_32 = V_MAX3_F32_e64 0, %57:vgpr_32, 0, %11:sgpr_32, 0, %58:vgpr_32, 0, 0, implicit $mode, implicit $exec617 %90:vgpr_32 = V_MAX3_F32_e64 0, %89:vgpr_32, 0, %59:vgpr_32, 0, %60:vgpr_32, 0, 0, implicit $mode, implicit $exec618 %91:vgpr_32 = V_MAX3_F32_e64 0, %90:vgpr_32, 0, %61:vgpr_32, 0, %62:vgpr_32, 0, 0, implicit $mode, implicit $exec619 %92:vgpr_32 = V_MAX3_F32_e64 0, %91:vgpr_32, 0, %63:vgpr_32, 0, %64:vgpr_32, 0, 0, implicit $mode, implicit $exec620 %93:vgpr_32 = V_MAX3_F32_e64 0, %92:vgpr_32, 0, %65:vgpr_32, 0, %66:vgpr_32, 0, 0, implicit $mode, implicit $exec621 %94:vgpr_32 = V_MAX3_F32_e64 0, %93:vgpr_32, 0, %67:vgpr_32, 0, %68:vgpr_32, 0, 0, implicit $mode, implicit $exec622 %95:vgpr_32 = V_MAX3_F32_e64 0, %94:vgpr_32, 0, %69:vgpr_32, 0, %70:vgpr_32, 0, 0, implicit $mode, implicit $exec623 %96:vgpr_32 = V_MAX3_F32_e64 0, %95:vgpr_32, 0, %71:vgpr_32, 0, %72:vgpr_32, 0, 0, implicit $mode, implicit $exec624 %97:vgpr_32 = V_MAX3_F32_e64 0, %96:vgpr_32, 0, %73:vgpr_32, 0, %74:vgpr_32, 0, 0, implicit $mode, implicit $exec625 %98:vgpr_32 = V_MAX3_F32_e64 0, %97:vgpr_32, 0, %75:vgpr_32, 0, %76:vgpr_32, 0, 0, implicit $mode, implicit $exec626 %99:vgpr_32 = V_MAX3_F32_e64 0, %98:vgpr_32, 0, %77:vgpr_32, 0, %78:vgpr_32, 0, 0, implicit $mode, implicit $exec627 %100:vgpr_32 = V_MAX3_F32_e64 0, %99:vgpr_32, 0, %79:vgpr_32, 0, %80:vgpr_32, 0, 0, implicit $mode, implicit $exec628 %101:vgpr_32 = V_MAX3_F32_e64 0, %100:vgpr_32, 0, %81:vgpr_32, 0, %82:vgpr_32, 0, 0, implicit $mode, implicit $exec629 %102:vgpr_32 = V_MAX3_F32_e64 0, %101:vgpr_32, 0, %83:vgpr_32, 0, %84:vgpr_32, 0, 0, implicit $mode, implicit $exec630 %103:vgpr_32 = V_MAX3_F32_e64 0, %102:vgpr_32, 0, %85:vgpr_32, 0, %86:vgpr_32, 0, 0, implicit $mode, implicit $exec631 %104:vgpr_32 = V_MAX3_F32_e64 0, %103:vgpr_32, 0, %87:vgpr_32, 0, %88:vgpr_32, 0, 0, implicit $mode, implicit $exec632 %105:vgpr_32 = DS_BPERMUTE_B32 %106:vgpr_32, %104:vgpr_32, 0, implicit $exec633 %107:vgpr_32 = contract nofpexcept V_MAX_F32_e32 %105:vgpr_32, %105:vgpr_32, implicit $mode, implicit $exec634 %108:vgpr_32 = contract nofpexcept V_MAX_F32_e32 %104:vgpr_32, %107:vgpr_32, implicit $mode, implicit $exec635 %109:vgpr_32 = DS_BPERMUTE_B32 %106:vgpr_32, %108:vgpr_32, 0, implicit $exec636 %110:vgpr_32 = V_CNDMASK_B32_e64 0, %109:vgpr_32, 0, %108:vgpr_32, %12:sreg_64_xexec, implicit $exec637 %111:vgpr_32 = contract nofpexcept V_MAX_F32_e32 %110:vgpr_32, %110:vgpr_32, implicit $mode, implicit $exec638 %112:vgpr_32 = contract nofpexcept V_MAX_F32_e32 %13:vgpr_32, %13:vgpr_32, implicit $mode, implicit $exec639 %113:vgpr_32 = contract nofpexcept V_MAX_F32_e32 %112:vgpr_32, %111:vgpr_32, implicit $mode, implicit $exec640 %114:vgpr_32 = contract nofpexcept V_FMA_F32_e64 0, %10.sub4:sgpr_512, 0, %43.sub0:vreg_512_align2, 1, %113:vgpr_32, 0, 0, implicit $mode, implicit $exec641 %115:vgpr_32 = afn nofpexcept V_MUL_F32_e32 1069066811, %114:vgpr_32, implicit $mode, implicit $exec642 %116:vgpr_32 = afn nofpexcept V_EXP_F32_e32 %115:vgpr_32, implicit $mode, implicit $exec643 %117:vgpr_32 = contract nofpexcept V_FMA_F32_e64 0, %10.sub4:sgpr_512, 0, %43.sub1:vreg_512_align2, 1, %113:vgpr_32, 0, 0, implicit $mode, implicit $exec644 %118:vgpr_32 = afn nofpexcept V_MUL_F32_e32 1069066811, %117:vgpr_32, implicit $mode, implicit $exec645 %119:vgpr_32 = afn nofpexcept V_EXP_F32_e32 %118:vgpr_32, implicit $mode, implicit $exec646 %120:vgpr_32 = contract nofpexcept V_FMA_F32_e64 0, %10.sub4:sgpr_512, 0, %43.sub2:vreg_512_align2, 1, %113:vgpr_32, 0, 0, implicit $mode, implicit $exec647 %121:vgpr_32 = afn nofpexcept V_MUL_F32_e32 1069066811, %120:vgpr_32, implicit $mode, implicit $exec648 %122:vgpr_32 = afn nofpexcept V_EXP_F32_e32 %121:vgpr_32, implicit $mode, implicit $exec649 %123:vgpr_32 = contract nofpexcept V_FMA_F32_e64 0, %10.sub4:sgpr_512, 0, %43.sub3:vreg_512_align2, 1, %113:vgpr_32, 0, 0, implicit $mode, implicit $exec650 %124:vgpr_32 = afn nofpexcept V_MUL_F32_e32 1069066811, %123:vgpr_32, implicit $mode, implicit $exec651 %125:vgpr_32 = afn nofpexcept V_EXP_F32_e32 %124:vgpr_32, implicit $mode, implicit $exec652 %126:vgpr_32 = contract nofpexcept V_FMA_F32_e64 0, %10.sub4:sgpr_512, 0, %43.sub4:vreg_512_align2, 1, %113:vgpr_32, 0, 0, implicit $mode, implicit $exec653 %127:vgpr_32 = afn nofpexcept V_MUL_F32_e32 1069066811, %126:vgpr_32, implicit $mode, implicit $exec654 %128:vgpr_32 = afn nofpexcept V_EXP_F32_e32 %127:vgpr_32, implicit $mode, implicit $exec655 %129:vgpr_32 = contract nofpexcept V_FMA_F32_e64 0, %10.sub4:sgpr_512, 0, %43.sub5:vreg_512_align2, 1, %113:vgpr_32, 0, 0, implicit $mode, implicit $exec656 %130:vgpr_32 = afn nofpexcept V_MUL_F32_e32 1069066811, %129:vgpr_32, implicit $mode, implicit $exec657 %131:vgpr_32 = afn nofpexcept V_EXP_F32_e32 %130:vgpr_32, implicit $mode, implicit $exec658 %132:vgpr_32 = contract nofpexcept V_FMA_F32_e64 0, %10.sub4:sgpr_512, 0, %43.sub6:vreg_512_align2, 1, %113:vgpr_32, 0, 0, implicit $mode, implicit $exec659 %133:vgpr_32 = afn nofpexcept V_MUL_F32_e32 1069066811, %132:vgpr_32, implicit $mode, implicit $exec660 %134:vgpr_32 = afn nofpexcept V_EXP_F32_e32 %133:vgpr_32, implicit $mode, implicit $exec661 %135:vgpr_32 = contract nofpexcept V_FMA_F32_e64 0, %10.sub4:sgpr_512, 0, %43.sub7:vreg_512_align2, 1, %113:vgpr_32, 0, 0, implicit $mode, implicit $exec662 %136:vgpr_32 = afn nofpexcept V_MUL_F32_e32 1069066811, %135:vgpr_32, implicit $mode, implicit $exec663 %137:vgpr_32 = afn nofpexcept V_EXP_F32_e32 %136:vgpr_32, implicit $mode, implicit $exec664 %138:vgpr_32 = contract nofpexcept V_FMA_F32_e64 0, %10.sub4:sgpr_512, 0, %43.sub8:vreg_512_align2, 1, %113:vgpr_32, 0, 0, implicit $mode, implicit $exec665 %139:vgpr_32 = afn nofpexcept V_MUL_F32_e32 1069066811, %138:vgpr_32, implicit $mode, implicit $exec666 %140:vgpr_32 = afn nofpexcept V_EXP_F32_e32 %139:vgpr_32, implicit $mode, implicit $exec667 %141:vgpr_32 = contract nofpexcept V_FMA_F32_e64 0, %10.sub4:sgpr_512, 0, %43.sub9:vreg_512_align2, 1, %113:vgpr_32, 0, 0, implicit $mode, implicit $exec668 %142:vgpr_32 = afn nofpexcept V_MUL_F32_e32 1069066811, %141:vgpr_32, implicit $mode, implicit $exec669 %143:vgpr_32 = afn nofpexcept V_EXP_F32_e32 %142:vgpr_32, implicit $mode, implicit $exec670 %144:vgpr_32 = contract nofpexcept V_FMA_F32_e64 0, %10.sub4:sgpr_512, 0, %43.sub10:vreg_512_align2, 1, %113:vgpr_32, 0, 0, implicit $mode, implicit $exec671 %145:vgpr_32 = afn nofpexcept V_MUL_F32_e32 1069066811, %144:vgpr_32, implicit $mode, implicit $exec672 %146:vgpr_32 = afn nofpexcept V_EXP_F32_e32 %145:vgpr_32, implicit $mode, implicit $exec673 %147:vgpr_32 = contract nofpexcept V_FMA_F32_e64 0, %10.sub4:sgpr_512, 0, %43.sub11:vreg_512_align2, 1, %113:vgpr_32, 0, 0, implicit $mode, implicit $exec674 %148:vgpr_32 = afn nofpexcept V_MUL_F32_e32 1069066811, %147:vgpr_32, implicit $mode, implicit $exec675 %149:vgpr_32 = afn nofpexcept V_EXP_F32_e32 %148:vgpr_32, implicit $mode, implicit $exec676 %150:vgpr_32 = contract nofpexcept V_FMA_F32_e64 0, %10.sub4:sgpr_512, 0, %43.sub12:vreg_512_align2, 1, %113:vgpr_32, 0, 0, implicit $mode, implicit $exec677 %151:vgpr_32 = afn nofpexcept V_MUL_F32_e32 1069066811, %150:vgpr_32, implicit $mode, implicit $exec678 %152:vgpr_32 = afn nofpexcept V_EXP_F32_e32 %151:vgpr_32, implicit $mode, implicit $exec679 %153:vgpr_32 = contract nofpexcept V_FMA_F32_e64 0, %10.sub4:sgpr_512, 0, %43.sub13:vreg_512_align2, 1, %113:vgpr_32, 0, 0, implicit $mode, implicit $exec680 %154:vgpr_32 = afn nofpexcept V_MUL_F32_e32 1069066811, %153:vgpr_32, implicit $mode, implicit $exec681 %155:vgpr_32 = afn nofpexcept V_EXP_F32_e32 %154:vgpr_32, implicit $mode, implicit $exec682 %156:vgpr_32 = contract nofpexcept V_FMA_F32_e64 0, %10.sub4:sgpr_512, 0, %43.sub14:vreg_512_align2, 1, %113:vgpr_32, 0, 0, implicit $mode, implicit $exec683 %157:vgpr_32 = afn nofpexcept V_MUL_F32_e32 1069066811, %156:vgpr_32, implicit $mode, implicit $exec684 %158:vgpr_32 = afn nofpexcept V_EXP_F32_e32 %157:vgpr_32, implicit $mode, implicit $exec685 %159:vgpr_32 = contract nofpexcept V_FMA_F32_e64 0, %10.sub4:sgpr_512, 0, %43.sub15:vreg_512_align2, 1, %113:vgpr_32, 0, 0, implicit $mode, implicit $exec686 %160:vgpr_32 = afn nofpexcept V_MUL_F32_e32 1069066811, %159:vgpr_32, implicit $mode, implicit $exec687 %161:vgpr_32 = afn nofpexcept V_EXP_F32_e32 %160:vgpr_32, implicit $mode, implicit $exec688 %162:vgpr_32 = contract nofpexcept V_FMA_F32_e64 0, %10.sub4:sgpr_512, 0, %46.sub0:vreg_512_align2, 1, %113:vgpr_32, 0, 0, implicit $mode, implicit $exec689 %163:vgpr_32 = afn nofpexcept V_MUL_F32_e32 1069066811, %162:vgpr_32, implicit $mode, implicit $exec690 %164:vgpr_32 = afn nofpexcept V_EXP_F32_e32 %163:vgpr_32, implicit $mode, implicit $exec691 %165:vgpr_32 = contract nofpexcept V_FMA_F32_e64 0, %10.sub4:sgpr_512, 0, %46.sub1:vreg_512_align2, 1, %113:vgpr_32, 0, 0, implicit $mode, implicit $exec692 %166:vgpr_32 = afn nofpexcept V_MUL_F32_e32 1069066811, %165:vgpr_32, implicit $mode, implicit $exec693 %167:vgpr_32 = afn nofpexcept V_EXP_F32_e32 %166:vgpr_32, implicit $mode, implicit $exec694 %168:vgpr_32 = contract nofpexcept V_FMA_F32_e64 0, %10.sub4:sgpr_512, 0, %46.sub2:vreg_512_align2, 1, %113:vgpr_32, 0, 0, implicit $mode, implicit $exec695 %169:vgpr_32 = afn nofpexcept V_MUL_F32_e32 1069066811, %168:vgpr_32, implicit $mode, implicit $exec696 %170:vgpr_32 = afn nofpexcept V_EXP_F32_e32 %169:vgpr_32, implicit $mode, implicit $exec697 %171:vgpr_32 = contract nofpexcept V_FMA_F32_e64 0, %10.sub4:sgpr_512, 0, %46.sub3:vreg_512_align2, 1, %113:vgpr_32, 0, 0, implicit $mode, implicit $exec698 %172:vgpr_32 = afn nofpexcept V_MUL_F32_e32 1069066811, %171:vgpr_32, implicit $mode, implicit $exec699 %173:vgpr_32 = afn nofpexcept V_EXP_F32_e32 %172:vgpr_32, implicit $mode, implicit $exec700 %174:vgpr_32 = contract nofpexcept V_FMA_F32_e64 0, %10.sub4:sgpr_512, 0, %46.sub4:vreg_512_align2, 1, %113:vgpr_32, 0, 0, implicit $mode, implicit $exec701 %175:vgpr_32 = afn nofpexcept V_MUL_F32_e32 1069066811, %174:vgpr_32, implicit $mode, implicit $exec702 %176:vgpr_32 = afn nofpexcept V_EXP_F32_e32 %175:vgpr_32, implicit $mode, implicit $exec703 %177:vgpr_32 = contract nofpexcept V_FMA_F32_e64 0, %10.sub4:sgpr_512, 0, %46.sub5:vreg_512_align2, 1, %113:vgpr_32, 0, 0, implicit $mode, implicit $exec704 %178:vgpr_32 = afn nofpexcept V_MUL_F32_e32 1069066811, %177:vgpr_32, implicit $mode, implicit $exec705 %179:vgpr_32 = afn nofpexcept V_EXP_F32_e32 %178:vgpr_32, implicit $mode, implicit $exec706 %180:vgpr_32 = contract nofpexcept V_FMA_F32_e64 0, %10.sub4:sgpr_512, 0, %46.sub6:vreg_512_align2, 1, %113:vgpr_32, 0, 0, implicit $mode, implicit $exec707 %181:vgpr_32 = afn nofpexcept V_MUL_F32_e32 1069066811, %180:vgpr_32, implicit $mode, implicit $exec708 %182:vgpr_32 = afn nofpexcept V_EXP_F32_e32 %181:vgpr_32, implicit $mode, implicit $exec709 %183:vgpr_32 = contract nofpexcept V_FMA_F32_e64 0, %10.sub4:sgpr_512, 0, %46.sub7:vreg_512_align2, 1, %113:vgpr_32, 0, 0, implicit $mode, implicit $exec710 %184:vgpr_32 = afn nofpexcept V_MUL_F32_e32 1069066811, %183:vgpr_32, implicit $mode, implicit $exec711 %185:vgpr_32 = afn nofpexcept V_EXP_F32_e32 %184:vgpr_32, implicit $mode, implicit $exec712 %186:vgpr_32 = contract nofpexcept V_FMA_F32_e64 0, %10.sub4:sgpr_512, 0, %46.sub8:vreg_512_align2, 1, %113:vgpr_32, 0, 0, implicit $mode, implicit $exec713 %187:vgpr_32 = afn nofpexcept V_MUL_F32_e32 1069066811, %186:vgpr_32, implicit $mode, implicit $exec714 %188:vgpr_32 = afn nofpexcept V_EXP_F32_e32 %187:vgpr_32, implicit $mode, implicit $exec715 %189:vgpr_32 = contract nofpexcept V_FMA_F32_e64 0, %10.sub4:sgpr_512, 0, %46.sub9:vreg_512_align2, 1, %113:vgpr_32, 0, 0, implicit $mode, implicit $exec716 %190:vgpr_32 = afn nofpexcept V_MUL_F32_e32 1069066811, %189:vgpr_32, implicit $mode, implicit $exec717 %191:vgpr_32 = afn nofpexcept V_EXP_F32_e32 %190:vgpr_32, implicit $mode, implicit $exec718 %192:vgpr_32 = contract nofpexcept V_FMA_F32_e64 0, %10.sub4:sgpr_512, 0, %46.sub10:vreg_512_align2, 1, %113:vgpr_32, 0, 0, implicit $mode, implicit $exec719 %193:vgpr_32 = afn nofpexcept V_MUL_F32_e32 1069066811, %192:vgpr_32, implicit $mode, implicit $exec720 %194:vgpr_32 = afn nofpexcept V_EXP_F32_e32 %193:vgpr_32, implicit $mode, implicit $exec721 %195:vgpr_32 = contract nofpexcept V_FMA_F32_e64 0, %10.sub4:sgpr_512, 0, %46.sub11:vreg_512_align2, 1, %113:vgpr_32, 0, 0, implicit $mode, implicit $exec722 %196:vgpr_32 = afn nofpexcept V_MUL_F32_e32 1069066811, %195:vgpr_32, implicit $mode, implicit $exec723 %197:vgpr_32 = afn nofpexcept V_EXP_F32_e32 %196:vgpr_32, implicit $mode, implicit $exec724 %198:vgpr_32 = contract nofpexcept V_FMA_F32_e64 0, %10.sub4:sgpr_512, 0, %46.sub12:vreg_512_align2, 1, %113:vgpr_32, 0, 0, implicit $mode, implicit $exec725 %199:vgpr_32 = afn nofpexcept V_MUL_F32_e32 1069066811, %198:vgpr_32, implicit $mode, implicit $exec726 %200:vgpr_32 = afn nofpexcept V_EXP_F32_e32 %199:vgpr_32, implicit $mode, implicit $exec727 %201:vgpr_32 = contract nofpexcept V_FMA_F32_e64 0, %10.sub4:sgpr_512, 0, %46.sub13:vreg_512_align2, 1, %113:vgpr_32, 0, 0, implicit $mode, implicit $exec728 %202:vgpr_32 = afn nofpexcept V_MUL_F32_e32 1069066811, %201:vgpr_32, implicit $mode, implicit $exec729 %203:vgpr_32 = afn nofpexcept V_EXP_F32_e32 %202:vgpr_32, implicit $mode, implicit $exec730 %204:vgpr_32 = contract nofpexcept V_FMA_F32_e64 0, %10.sub4:sgpr_512, 0, %46.sub14:vreg_512_align2, 1, %113:vgpr_32, 0, 0, implicit $mode, implicit $exec731 %205:vgpr_32 = afn nofpexcept V_MUL_F32_e32 1069066811, %204:vgpr_32, implicit $mode, implicit $exec732 %206:vgpr_32 = afn nofpexcept V_EXP_F32_e32 %205:vgpr_32, implicit $mode, implicit $exec733 %207:vgpr_32 = contract nofpexcept V_FMA_F32_e64 0, %10.sub4:sgpr_512, 0, %46.sub15:vreg_512_align2, 1, %113:vgpr_32, 0, 0, implicit $mode, implicit $exec734 %208:vgpr_32 = afn nofpexcept V_MUL_F32_e32 1069066811, %207:vgpr_32, implicit $mode, implicit $exec735 %209:vgpr_32 = afn nofpexcept V_EXP_F32_e32 %208:vgpr_32, implicit $mode, implicit $exec736 %210:vgpr_32 = contract nofpexcept V_ADD_F32_e32 0, %116:vgpr_32, implicit $mode, implicit $exec737 %211:vgpr_32 = contract nofpexcept V_ADD_F32_e32 %119:vgpr_32, %210:vgpr_32, implicit $mode, implicit $exec738 %212:vgpr_32 = contract nofpexcept V_ADD_F32_e32 %122:vgpr_32, %211:vgpr_32, implicit $mode, implicit $exec739 %213:vgpr_32 = contract nofpexcept V_ADD_F32_e32 %125:vgpr_32, %212:vgpr_32, implicit $mode, implicit $exec740 %214:vgpr_32 = contract nofpexcept V_ADD_F32_e32 %128:vgpr_32, %213:vgpr_32, implicit $mode, implicit $exec741 %215:vgpr_32 = contract nofpexcept V_ADD_F32_e32 %131:vgpr_32, %214:vgpr_32, implicit $mode, implicit $exec742 %216:vgpr_32 = contract nofpexcept V_ADD_F32_e32 %134:vgpr_32, %215:vgpr_32, implicit $mode, implicit $exec743 %217:vgpr_32 = contract nofpexcept V_ADD_F32_e32 %137:vgpr_32, %216:vgpr_32, implicit $mode, implicit $exec744 %218:vgpr_32 = contract nofpexcept V_ADD_F32_e32 %140:vgpr_32, %217:vgpr_32, implicit $mode, implicit $exec745 %219:vgpr_32 = contract nofpexcept V_ADD_F32_e32 %143:vgpr_32, %218:vgpr_32, implicit $mode, implicit $exec746 %220:vgpr_32 = contract nofpexcept V_ADD_F32_e32 %146:vgpr_32, %219:vgpr_32, implicit $mode, implicit $exec747 %221:vgpr_32 = contract nofpexcept V_ADD_F32_e32 %149:vgpr_32, %220:vgpr_32, implicit $mode, implicit $exec748 %222:vgpr_32 = contract nofpexcept V_ADD_F32_e32 %152:vgpr_32, %221:vgpr_32, implicit $mode, implicit $exec749 %223:vgpr_32 = contract nofpexcept V_ADD_F32_e32 %155:vgpr_32, %222:vgpr_32, implicit $mode, implicit $exec750 %224:vgpr_32 = contract nofpexcept V_ADD_F32_e32 %158:vgpr_32, %223:vgpr_32, implicit $mode, implicit $exec751 %225:vgpr_32 = contract nofpexcept V_ADD_F32_e32 %161:vgpr_32, %224:vgpr_32, implicit $mode, implicit $exec752 %226:vgpr_32 = contract nofpexcept V_ADD_F32_e32 %164:vgpr_32, %225:vgpr_32, implicit $mode, implicit $exec753 %227:vgpr_32 = contract nofpexcept V_ADD_F32_e32 %167:vgpr_32, %226:vgpr_32, implicit $mode, implicit $exec754 %228:vgpr_32 = contract nofpexcept V_ADD_F32_e32 %170:vgpr_32, %227:vgpr_32, implicit $mode, implicit $exec755 %229:vgpr_32 = contract nofpexcept V_ADD_F32_e32 %173:vgpr_32, %228:vgpr_32, implicit $mode, implicit $exec756 %230:vgpr_32 = contract nofpexcept V_ADD_F32_e32 %176:vgpr_32, %229:vgpr_32, implicit $mode, implicit $exec757 %231:vgpr_32 = contract nofpexcept V_ADD_F32_e32 %179:vgpr_32, %230:vgpr_32, implicit $mode, implicit $exec758 %232:vgpr_32 = contract nofpexcept V_ADD_F32_e32 %182:vgpr_32, %231:vgpr_32, implicit $mode, implicit $exec759 %233:vgpr_32 = contract nofpexcept V_ADD_F32_e32 %185:vgpr_32, %232:vgpr_32, implicit $mode, implicit $exec760 %234:vgpr_32 = contract nofpexcept V_ADD_F32_e32 %188:vgpr_32, %233:vgpr_32, implicit $mode, implicit $exec761 %235:vgpr_32 = contract nofpexcept V_ADD_F32_e32 %191:vgpr_32, %234:vgpr_32, implicit $mode, implicit $exec762 %236:vgpr_32 = contract nofpexcept V_ADD_F32_e32 %194:vgpr_32, %235:vgpr_32, implicit $mode, implicit $exec763 %237:vgpr_32 = contract nofpexcept V_ADD_F32_e32 %197:vgpr_32, %236:vgpr_32, implicit $mode, implicit $exec764 %238:vgpr_32 = contract nofpexcept V_ADD_F32_e32 %200:vgpr_32, %237:vgpr_32, implicit $mode, implicit $exec765 %239:vgpr_32 = contract nofpexcept V_ADD_F32_e32 %203:vgpr_32, %238:vgpr_32, implicit $mode, implicit $exec766 %240:vgpr_32 = contract nofpexcept V_ADD_F32_e32 %206:vgpr_32, %239:vgpr_32, implicit $mode, implicit $exec767 %241:vgpr_32 = contract nofpexcept V_ADD_F32_e32 %209:vgpr_32, %240:vgpr_32, implicit $mode, implicit $exec768 %242:vgpr_32 = DS_BPERMUTE_B32 %106:vgpr_32, %241:vgpr_32, 0, implicit $exec769 %243:vgpr_32 = contract nofpexcept V_ADD_F32_e32 %241:vgpr_32, %242:vgpr_32, implicit $mode, implicit $exec770 %244:vgpr_32 = DS_BPERMUTE_B32 %106:vgpr_32, %243:vgpr_32, 0, implicit $exec771 %0:vgpr_32 = V_CNDMASK_B32_e64 0, %244:vgpr_32, 0, %243:vgpr_32, %12:sreg_64_xexec, implicit $exec772 %245:vgpr_32 = contract nofpexcept V_SUB_F32_e32 %13:vgpr_32, %113:vgpr_32, implicit $mode, implicit $exec773 %246:vgpr_32 = afn nofpexcept V_MUL_F32_e32 1069066811, %245:vgpr_32, implicit $mode, implicit $exec774 undef %247.sub0:vreg_64_align2 = afn nofpexcept V_EXP_F32_e32 %246:vgpr_32, implicit $mode, implicit $exec775 INLINEASM &"s_waitcnt vmcnt($0)", 57 /* sideeffect mayload maystore isconvergent attdialect */, 13 /* imm */, 8, !0776 %248:vgpr_32 = V_PERM_B32_e64 %40.sub0:vreg_64_align2, %38.sub0:vreg_64_align2, %14:sreg_32, implicit $exec777 %249:vgpr_32 = V_PERM_B32_e64 %40.sub0:vreg_64_align2, %38.sub0:vreg_64_align2, %15:sreg_32, implicit $exec778 %250:vgpr_32 = V_PERM_B32_e64 %40.sub1:vreg_64_align2, %38.sub1:vreg_64_align2, %14:sreg_32, implicit $exec779 %251:vgpr_32 = V_PERM_B32_e64 %40.sub1:vreg_64_align2, %38.sub1:vreg_64_align2, %15:sreg_32, implicit $exec780 %252:vgpr_32 = V_ADD_U32_e32 %27:sreg_32_xm0, %16:vgpr_32, implicit $exec781 %253:vgpr_32 = V_AND_B32_e32 536870911, %252:vgpr_32, implicit $exec782 %254:vgpr_32 = nsw V_MUL_LO_U32_e64 %253:vgpr_32, %17:sreg_32, implicit $exec783 %255:vgpr_32 = V_ADD_LSHL_U32_e64 %18:vgpr_32, %254:vgpr_32, 1, implicit $exec784 DS_WRITE_B32_gfx9 %255:vgpr_32, %248:vgpr_32, 0, 0, implicit $exec785 %256:vgpr_32 = V_LSHL_ADD_U32_e64 %19:vgpr_32, 1, %255:vgpr_32, implicit $exec786 DS_WRITE_B32_gfx9 %256:vgpr_32, %249:vgpr_32, 0, 0, implicit $exec787 %257:vgpr_32 = V_LSHL_ADD_U32_e64 %20:vgpr_32, 1, %256:vgpr_32, implicit $exec788 DS_WRITE_B32_gfx9 %257:vgpr_32, %250:vgpr_32, 0, 0, implicit $exec789 %258:vgpr_32 = V_LSHL_ADD_U32_e64 %21:vgpr_32, 1, %257:vgpr_32, implicit $exec790 DS_WRITE_B32_gfx9 %258:vgpr_32, %251:vgpr_32, 0, 0, implicit $exec791 %0:vgpr_32 = contract nofpexcept V_FMAC_F32_e32 %1:vgpr_32, %247.sub0:vreg_64_align2, %0:vgpr_32, implicit $mode, implicit $exec792 %259.sub0_sub1:vreg_512_align2 = contract nofpexcept V_PK_MUL_F32 8, %259.sub0_sub1:vreg_512_align2, 0, %247:vreg_64_align2, 0, 0, 0, 0, 0, implicit $mode, implicit $exec793 %259.sub2_sub3:vreg_512_align2 = contract nofpexcept V_PK_MUL_F32 8, %259.sub2_sub3:vreg_512_align2, 0, %247:vreg_64_align2, 0, 0, 0, 0, 0, implicit $mode, implicit $exec794 %259.sub4_sub5:vreg_512_align2 = contract nofpexcept V_PK_MUL_F32 8, %259.sub4_sub5:vreg_512_align2, 0, %247:vreg_64_align2, 0, 0, 0, 0, 0, implicit $mode, implicit $exec795 %259.sub6_sub7:vreg_512_align2 = contract nofpexcept V_PK_MUL_F32 8, %259.sub6_sub7:vreg_512_align2, 0, %247:vreg_64_align2, 0, 0, 0, 0, 0, implicit $mode, implicit $exec796 %259.sub8_sub9:vreg_512_align2 = contract nofpexcept V_PK_MUL_F32 8, %259.sub8_sub9:vreg_512_align2, 0, %247:vreg_64_align2, 0, 0, 0, 0, 0, implicit $mode, implicit $exec797 %259.sub10_sub11:vreg_512_align2 = contract nofpexcept V_PK_MUL_F32 8, %259.sub10_sub11:vreg_512_align2, 0, %247:vreg_64_align2, 0, 0, 0, 0, 0, implicit $mode, implicit $exec798 %259.sub12_sub13:vreg_512_align2 = contract nofpexcept V_PK_MUL_F32 8, %259.sub12_sub13:vreg_512_align2, 0, %247:vreg_64_align2, 0, 0, 0, 0, 0, implicit $mode, implicit $exec799 %259.sub14_sub15:vreg_512_align2 = contract nofpexcept V_PK_MUL_F32 8, %259.sub14_sub15:vreg_512_align2, 0, %247:vreg_64_align2, 0, 0, 0, 0, 0, implicit $mode, implicit $exec800 %260.sub0_sub1:vreg_512_align2 = contract nofpexcept V_PK_MUL_F32 8, %260.sub0_sub1:vreg_512_align2, 0, %247:vreg_64_align2, 0, 0, 0, 0, 0, implicit $mode, implicit $exec801 %260.sub2_sub3:vreg_512_align2 = contract nofpexcept V_PK_MUL_F32 8, %260.sub2_sub3:vreg_512_align2, 0, %247:vreg_64_align2, 0, 0, 0, 0, 0, implicit $mode, implicit $exec802 %260.sub4_sub5:vreg_512_align2 = contract nofpexcept V_PK_MUL_F32 8, %260.sub4_sub5:vreg_512_align2, 0, %247:vreg_64_align2, 0, 0, 0, 0, 0, implicit $mode, implicit $exec803 %260.sub6_sub7:vreg_512_align2 = contract nofpexcept V_PK_MUL_F32 8, %260.sub6_sub7:vreg_512_align2, 0, %247:vreg_64_align2, 0, 0, 0, 0, 0, implicit $mode, implicit $exec804 %260.sub8_sub9:vreg_512_align2 = contract nofpexcept V_PK_MUL_F32 8, %260.sub8_sub9:vreg_512_align2, 0, %247:vreg_64_align2, 0, 0, 0, 0, 0, implicit $mode, implicit $exec805 %260.sub10_sub11:vreg_512_align2 = contract nofpexcept V_PK_MUL_F32 8, %260.sub10_sub11:vreg_512_align2, 0, %247:vreg_64_align2, 0, 0, 0, 0, 0, implicit $mode, implicit $exec806 %260.sub12_sub13:vreg_512_align2 = contract nofpexcept V_PK_MUL_F32 8, %260.sub12_sub13:vreg_512_align2, 0, %247:vreg_64_align2, 0, 0, 0, 0, 0, implicit $mode, implicit $exec807 %260.sub14_sub15:vreg_512_align2 = contract nofpexcept V_PK_MUL_F32 8, %260.sub14_sub15:vreg_512_align2, 0, %247:vreg_64_align2, 0, 0, 0, 0, 0, implicit $mode, implicit $exec808 %261:vgpr_32 = nofpexcept V_CVT_F16_F32_e32 %116:vgpr_32, implicit $mode, implicit $exec809 %262:vgpr_32 = nofpexcept V_CVT_F16_F32_e32 %119:vgpr_32, implicit $mode, implicit $exec810 %263:vgpr_32 = nofpexcept V_CVT_F16_F32_e32 %122:vgpr_32, implicit $mode, implicit $exec811 %264:vgpr_32 = nofpexcept V_CVT_F16_F32_e32 %128:vgpr_32, implicit $mode, implicit $exec812 %265:vgpr_32 = nofpexcept V_CVT_F16_F32_e32 %131:vgpr_32, implicit $mode, implicit $exec813 %266:vgpr_32 = nofpexcept V_CVT_F16_F32_e32 %134:vgpr_32, implicit $mode, implicit $exec814 %267:vgpr_32 = nofpexcept V_CVT_F16_F32_e32 %140:vgpr_32, implicit $mode, implicit $exec815 %268:vgpr_32 = nofpexcept V_CVT_F16_F32_e32 %143:vgpr_32, implicit $mode, implicit $exec816 %269:vgpr_32 = nofpexcept V_CVT_F16_F32_e32 %146:vgpr_32, implicit $mode, implicit $exec817 %270:vgpr_32 = nofpexcept V_CVT_F16_F32_e32 %152:vgpr_32, implicit $mode, implicit $exec818 %271:vgpr_32 = nofpexcept V_CVT_F16_F32_e32 %155:vgpr_32, implicit $mode, implicit $exec819 %272:vgpr_32 = nofpexcept V_CVT_F16_F32_e32 %158:vgpr_32, implicit $mode, implicit $exec820 %273:vgpr_32 = V_ADD_U32_e32 %22:vgpr_32, %37:vgpr_32, implicit $exec821 %274:vreg_64_align2 = BUFFER_LOAD_DWORDX2_OFFEN %273:vgpr_32, %8:sgpr_128, 0, 0, 0, 0, implicit $exec822 %275:vgpr_32 = V_ADD_U32_e32 %23:vgpr_32, %37:vgpr_32, implicit $exec823 %276:vreg_64_align2 = BUFFER_LOAD_DWORDX2_OFFEN %275:vgpr_32, %8:sgpr_128, 0, 0, 0, 0, implicit $exec824 INLINEASM &"s_waitcnt vmcnt($0)", 57 /* sideeffect mayload maystore isconvergent attdialect */, 13 /* imm */, 8, !0825 %277:vreg_128_align2 = DS_READ_B128_gfx9 %24:vgpr_32, 0, 0, implicit $exec826 %278:vreg_128_align2 = DS_READ_B128_gfx9 %24:vgpr_32, 576, 0, implicit $exec827 %279:vreg_128_align2 = DS_READ_B128_gfx9 %25:vgpr_32, 0, 0, implicit $exec828 %280:vreg_128_align2 = DS_READ_B128_gfx9 %25:vgpr_32, 576, 0, implicit $exec829 INLINEASM &"s_waitcnt vmcnt($0)", 57 /* sideeffect mayload maystore isconvergent attdialect */, 13 /* imm */, 8, !0830 %281:vgpr_32 = V_PERM_B32_e64 %276.sub0:vreg_64_align2, %274.sub0:vreg_64_align2, %14:sreg_32, implicit $exec831 %282:vgpr_32 = V_PERM_B32_e64 %276.sub0:vreg_64_align2, %274.sub0:vreg_64_align2, %15:sreg_32, implicit $exec832 %283:vgpr_32 = V_PERM_B32_e64 %276.sub1:vreg_64_align2, %274.sub1:vreg_64_align2, %14:sreg_32, implicit $exec833 %284:vgpr_32 = V_PERM_B32_e64 %276.sub1:vreg_64_align2, %274.sub1:vreg_64_align2, %15:sreg_32, implicit $exec834 DS_WRITE_B32_gfx9 %255:vgpr_32, %281:vgpr_32, 0, 0, implicit $exec835 DS_WRITE_B32_gfx9 %256:vgpr_32, %282:vgpr_32, 0, 0, implicit $exec836 DS_WRITE_B32_gfx9 %257:vgpr_32, %283:vgpr_32, 0, 0, implicit $exec837 DS_WRITE_B32_gfx9 %258:vgpr_32, %284:vgpr_32, 0, 0, implicit $exec838 %285:vgpr_32 = nofpexcept V_CVT_F16_F32_e32 %125:vgpr_32, implicit $mode, implicit $exec839 %286:vgpr_32 = nofpexcept V_CVT_F16_F32_e32 %137:vgpr_32, implicit $mode, implicit $exec840 %287:vgpr_32 = nofpexcept V_CVT_F16_F32_e32 %149:vgpr_32, implicit $mode, implicit $exec841 %288:vgpr_32 = nofpexcept V_CVT_F16_F32_e32 %161:vgpr_32, implicit $mode, implicit $exec842 undef %289.sub1:vreg_64_align2 = nofpexcept V_PACK_B32_F16_e64 0, %263:vgpr_32, 0, %285:vgpr_32, 0, 0, implicit $mode, implicit $exec843 %289.sub0:vreg_64_align2 = nofpexcept V_PACK_B32_F16_e64 0, %261:vgpr_32, 0, %262:vgpr_32, 0, 0, implicit $mode, implicit $exec844 undef %290.sub1:vreg_64_align2 = nofpexcept V_PACK_B32_F16_e64 0, %266:vgpr_32, 0, %286:vgpr_32, 0, 0, implicit $mode, implicit $exec845 %290.sub0:vreg_64_align2 = nofpexcept V_PACK_B32_F16_e64 0, %264:vgpr_32, 0, %265:vgpr_32, 0, 0, implicit $mode, implicit $exec846 undef %291.sub1:vreg_64_align2 = nofpexcept V_PACK_B32_F16_e64 0, %269:vgpr_32, 0, %287:vgpr_32, 0, 0, implicit $mode, implicit $exec847 %291.sub0:vreg_64_align2 = nofpexcept V_PACK_B32_F16_e64 0, %267:vgpr_32, 0, %268:vgpr_32, 0, 0, implicit $mode, implicit $exec848 undef %292.sub1:vreg_64_align2 = nofpexcept V_PACK_B32_F16_e64 0, %272:vgpr_32, 0, %288:vgpr_32, 0, 0, implicit $mode, implicit $exec849 %292.sub0:vreg_64_align2 = nofpexcept V_PACK_B32_F16_e64 0, %270:vgpr_32, 0, %271:vgpr_32, 0, 0, implicit $mode, implicit $exec850 %259:vreg_512_align2 = contract V_MFMA_F32_32X32X8F16_mac_vgprcd_e64 %277.sub0_sub1:vreg_128_align2, %289:vreg_64_align2, %259:vreg_512_align2, 0, 0, 0, implicit $mode, implicit $exec851 %259:vreg_512_align2 = contract V_MFMA_F32_32X32X8F16_mac_vgprcd_e64 %277.sub2_sub3:vreg_128_align2, %290:vreg_64_align2, %259:vreg_512_align2, 0, 0, 0, implicit $mode, implicit $exec852 %260:vreg_512_align2 = contract V_MFMA_F32_32X32X8F16_mac_vgprcd_e64 %278.sub0_sub1:vreg_128_align2, %289:vreg_64_align2, %260:vreg_512_align2, 0, 0, 0, implicit $mode, implicit $exec853 %260:vreg_512_align2 = contract V_MFMA_F32_32X32X8F16_mac_vgprcd_e64 %278.sub2_sub3:vreg_128_align2, %290:vreg_64_align2, %260:vreg_512_align2, 0, 0, 0, implicit $mode, implicit $exec854 %259:vreg_512_align2 = contract V_MFMA_F32_32X32X8F16_mac_vgprcd_e64 %279.sub0_sub1:vreg_128_align2, %291:vreg_64_align2, %259:vreg_512_align2, 0, 0, 0, implicit $mode, implicit $exec855 %259:vreg_512_align2 = contract V_MFMA_F32_32X32X8F16_mac_vgprcd_e64 %279.sub2_sub3:vreg_128_align2, %292:vreg_64_align2, %259:vreg_512_align2, 0, 0, 0, implicit $mode, implicit $exec856 %260:vreg_512_align2 = contract V_MFMA_F32_32X32X8F16_mac_vgprcd_e64 %280.sub0_sub1:vreg_128_align2, %291:vreg_64_align2, %260:vreg_512_align2, 0, 0, 0, implicit $mode, implicit $exec857 %260:vreg_512_align2 = contract V_MFMA_F32_32X32X8F16_mac_vgprcd_e64 %280.sub2_sub3:vreg_128_align2, %292:vreg_64_align2, %260:vreg_512_align2, 0, 0, 0, implicit $mode, implicit $exec858 %293:vgpr_32 = nofpexcept V_CVT_F16_F32_e32 %209:vgpr_32, implicit $mode, implicit $exec859 %294:vgpr_32 = nofpexcept V_CVT_F16_F32_e32 %203:vgpr_32, implicit $mode, implicit $exec860 %295:vgpr_32 = nofpexcept V_CVT_F16_F32_e32 %197:vgpr_32, implicit $mode, implicit $exec861 %296:vgpr_32 = nofpexcept V_CVT_F16_F32_e32 %191:vgpr_32, implicit $mode, implicit $exec862 %297:vgpr_32 = nofpexcept V_CVT_F16_F32_e32 %185:vgpr_32, implicit $mode, implicit $exec863 %298:vgpr_32 = nofpexcept V_CVT_F16_F32_e32 %179:vgpr_32, implicit $mode, implicit $exec864 %299:vgpr_32 = nofpexcept V_CVT_F16_F32_e32 %173:vgpr_32, implicit $mode, implicit $exec865 %300:vgpr_32 = nofpexcept V_CVT_F16_F32_e32 %167:vgpr_32, implicit $mode, implicit $exec866 %301:vgpr_32 = nofpexcept V_CVT_F16_F32_e32 %206:vgpr_32, implicit $mode, implicit $exec867 %302:vgpr_32 = nofpexcept V_CVT_F16_F32_e32 %200:vgpr_32, implicit $mode, implicit $exec868 %303:vgpr_32 = nofpexcept V_CVT_F16_F32_e32 %194:vgpr_32, implicit $mode, implicit $exec869 %304:vgpr_32 = nofpexcept V_CVT_F16_F32_e32 %188:vgpr_32, implicit $mode, implicit $exec870 %305:vgpr_32 = nofpexcept V_CVT_F16_F32_e32 %182:vgpr_32, implicit $mode, implicit $exec871 %306:vgpr_32 = nofpexcept V_CVT_F16_F32_e32 %176:vgpr_32, implicit $mode, implicit $exec872 %307:vgpr_32 = nofpexcept V_CVT_F16_F32_e32 %170:vgpr_32, implicit $mode, implicit $exec873 %308:vgpr_32 = nofpexcept V_CVT_F16_F32_e32 %164:vgpr_32, implicit $mode, implicit $exec874 INLINEASM &"s_waitcnt vmcnt($0)", 57 /* sideeffect mayload maystore isconvergent attdialect */, 13 /* imm */, 8, !0875 undef %309.sub1:vreg_64_align2 = nofpexcept V_PACK_B32_F16_e64 0, %307:vgpr_32, 0, %299:vgpr_32, 0, 0, implicit $mode, implicit $exec876 %309.sub0:vreg_64_align2 = nofpexcept V_PACK_B32_F16_e64 0, %308:vgpr_32, 0, %300:vgpr_32, 0, 0, implicit $mode, implicit $exec877 undef %310.sub1:vreg_64_align2 = nofpexcept V_PACK_B32_F16_e64 0, %305:vgpr_32, 0, %297:vgpr_32, 0, 0, implicit $mode, implicit $exec878 %310.sub0:vreg_64_align2 = nofpexcept V_PACK_B32_F16_e64 0, %306:vgpr_32, 0, %298:vgpr_32, 0, 0, implicit $mode, implicit $exec879 undef %311.sub1:vreg_64_align2 = nofpexcept V_PACK_B32_F16_e64 0, %303:vgpr_32, 0, %295:vgpr_32, 0, 0, implicit $mode, implicit $exec880 %311.sub0:vreg_64_align2 = nofpexcept V_PACK_B32_F16_e64 0, %304:vgpr_32, 0, %296:vgpr_32, 0, 0, implicit $mode, implicit $exec881 undef %312.sub1:vreg_64_align2 = nofpexcept V_PACK_B32_F16_e64 0, %301:vgpr_32, 0, %293:vgpr_32, 0, 0, implicit $mode, implicit $exec882 %312.sub0:vreg_64_align2 = nofpexcept V_PACK_B32_F16_e64 0, %302:vgpr_32, 0, %294:vgpr_32, 0, 0, implicit $mode, implicit $exec883 %313:vreg_128_align2 = DS_READ_B128_gfx9 %24:vgpr_32, 0, 0, implicit $exec884 %259:vreg_512_align2 = contract V_MFMA_F32_32X32X8F16_mac_vgprcd_e64 %313.sub0_sub1:vreg_128_align2, %309:vreg_64_align2, %259:vreg_512_align2, 0, 0, 0, implicit $mode, implicit $exec885 %259:vreg_512_align2 = contract V_MFMA_F32_32X32X8F16_mac_vgprcd_e64 %313.sub2_sub3:vreg_128_align2, %310:vreg_64_align2, %259:vreg_512_align2, 0, 0, 0, implicit $mode, implicit $exec886 %314:vreg_128_align2 = DS_READ_B128_gfx9 %24:vgpr_32, 576, 0, implicit $exec887 %260:vreg_512_align2 = contract V_MFMA_F32_32X32X8F16_mac_vgprcd_e64 %314.sub0_sub1:vreg_128_align2, %309:vreg_64_align2, %260:vreg_512_align2, 0, 0, 0, implicit $mode, implicit $exec888 %260:vreg_512_align2 = contract V_MFMA_F32_32X32X8F16_mac_vgprcd_e64 %314.sub2_sub3:vreg_128_align2, %310:vreg_64_align2, %260:vreg_512_align2, 0, 0, 0, implicit $mode, implicit $exec889 %315:vreg_128_align2 = DS_READ_B128_gfx9 %25:vgpr_32, 0, 0, implicit $exec890 %259:vreg_512_align2 = contract V_MFMA_F32_32X32X8F16_mac_vgprcd_e64 %315.sub0_sub1:vreg_128_align2, %311:vreg_64_align2, %259:vreg_512_align2, 0, 0, 0, implicit $mode, implicit $exec891 %259:vreg_512_align2 = contract V_MFMA_F32_32X32X8F16_mac_vgprcd_e64 %315.sub2_sub3:vreg_128_align2, %312:vreg_64_align2, %259:vreg_512_align2, 0, 0, 0, implicit $mode, implicit $exec892 %316:vreg_128_align2 = DS_READ_B128_gfx9 %25:vgpr_32, 576, 0, implicit $exec893 %260:vreg_512_align2 = contract V_MFMA_F32_32X32X8F16_mac_vgprcd_e64 %316.sub0_sub1:vreg_128_align2, %311:vreg_64_align2, %260:vreg_512_align2, 0, 0, 0, implicit $mode, implicit $exec894 %260:vreg_512_align2 = contract V_MFMA_F32_32X32X8F16_mac_vgprcd_e64 %316.sub2_sub3:vreg_128_align2, %312:vreg_64_align2, %260:vreg_512_align2, 0, 0, 0, implicit $mode, implicit $exec895 INLINEASM &"s_waitcnt vmcnt($0)", 57 /* sideeffect mayload maystore isconvergent attdialect */, 13 /* imm */, 8, !0896 %37:vgpr_32 = V_ADD_U32_e32 %26:sreg_32, %37:vgpr_32, implicit $exec897 %29:vgpr_32 = nuw V_ADD_U32_e32 64, %29:vgpr_32, implicit $exec898 S_ENDPGM 0899...900