2060 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 @largeInterleave() #0 { ret void }7 ; GCN-LABEL: largeInterleave:8 ; GCN: ; %bb.0:9 ; GCN-NEXT: ; implicit-def: $vgpr1610 ; GCN-NEXT: ; implicit-def: $vgpr2511 ; GCN-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)12 ; GCN-NEXT: v_readfirstlane_b32 s17, v1613 ; GCN-NEXT: ; implicit-def: $sgpr0_sgpr1_sgpr2_sgpr3_sgpr4_sgpr5_sgpr6_sgpr7_sgpr8_sgpr9_sgpr10_sgpr11_sgpr12_sgpr13_sgpr14_sgpr1514 ; GCN-NEXT: ; implicit-def: $vgpr1715 ; GCN-NEXT: ; implicit-def: $sgpr1516 ; GCN-NEXT: ; implicit-def: $sgpr8_sgpr9_sgpr10_sgpr1117 ; GCN-NEXT: s_lshl_b32 s18, s17, 718 ; GCN-NEXT: ; implicit-def: $vgpr1819 ; GCN-NEXT: v_add_lshl_u32 v230, v18, s18, 120 ; GCN-NEXT: v_lshl_add_u32 v25, s17, 4, v2521 ; GCN-NEXT: v_mul_lo_u32 v25, v25, s622 ; GCN-NEXT: v_add_lshl_u32 v226, v25, v17, 123 ; GCN-NEXT: v_add_u32_e32 v17, s15, v22624 ; GCN-NEXT: buffer_load_dwordx4 v[64:67], v226, s[8:11], 0 offen sc0 sc125 ; GCN-NEXT: s_waitcnt vmcnt(0)26 ; GCN-NEXT: buffer_inv sc0 sc127 ; GCN-NEXT: buffer_load_dwordx4 v[68:71], v17, s[8:11], 0 offen sc0 sc128 ; GCN-NEXT: s_waitcnt vmcnt(0)29 ; GCN-NEXT: buffer_inv sc0 sc130 ; GCN-NEXT: v_add_u32_e32 v72, 64, v1731 ; GCN-NEXT: ; implicit-def: $vgpr21332 ; GCN-NEXT: ; implicit-def: $vgpr152_vgpr153_vgpr154_vgpr15533 ; GCN-NEXT: ; implicit-def: $vgpr24634 ; GCN-NEXT: v_add_u32_e32 v188, 0x80, v1735 ; GCN-NEXT: ; implicit-def: $vgpr156_vgpr157_vgpr158_vgpr15936 ; GCN-NEXT: ; implicit-def: $vgpr144_vgpr145_vgpr146_vgpr14737 ; GCN-NEXT: ; implicit-def: $vgpr1938 ; GCN-NEXT: ; implicit-def: $vgpr2639 ; GCN-NEXT: ; implicit-def: $vgpr2740 ; GCN-NEXT: v_add_u32_e32 v227, 0xc0, v1741 ; GCN-NEXT: v_add_u32_e32 v231, v19, v2642 ; GCN-NEXT: v_add_u32_e32 v232, v19, v2743 ; GCN-NEXT: ; implicit-def: $sgpr0_sgpr1_sgpr2_sgpr344 ; GCN-NEXT: ; implicit-def: $vgpr2845 ; GCN-NEXT: ; implicit-def: $vgpr2946 ; GCN-NEXT: v_add_u32_e32 v233, v19, v2847 ; GCN-NEXT: v_add_u32_e32 v234, v19, v2948 ; GCN-NEXT: ; implicit-def: $vgpr140_vgpr141_vgpr142_vgpr14349 ; GCN-NEXT: ; implicit-def: $sgpr550 ; GCN-NEXT: ; implicit-def: $sgpr751 ; GCN-NEXT: ; implicit-def: $vgpr148_vgpr149_vgpr150_vgpr15152 ; GCN-NEXT: ; implicit-def: $vgpr136_vgpr137_vgpr138_vgpr13953 ; GCN-NEXT: ; implicit-def: $vgpr132_vgpr133_vgpr134_vgpr13554 ; GCN-NEXT: ; implicit-def: $vgpr2055 ; GCN-NEXT: v_add_u32_e32 v18, s17, v2056 ; GCN-NEXT: v_and_b32_e32 v18, 0x1fffffff, v1857 ; GCN-NEXT: ; implicit-def: $sgpr1658 ; GCN-NEXT: v_mul_lo_u32 v18, v18, s1659 ; GCN-NEXT: ; implicit-def: $vgpr2160 ; GCN-NEXT: v_add_lshl_u32 v199, v21, v18, 161 ; GCN-NEXT: ; implicit-def: $vgpr2262 ; GCN-NEXT: v_lshl_add_u32 v200, v22, 1, v19963 ; GCN-NEXT: ; implicit-def: $vgpr2364 ; GCN-NEXT: v_lshl_add_u32 v201, v23, 1, v20065 ; GCN-NEXT: ; implicit-def: $vgpr2466 ; GCN-NEXT: v_lshl_add_u32 v202, v24, 1, v20167 ; GCN-NEXT: ; implicit-def: $vgpr1668 ; GCN-NEXT: ; implicit-def: $vgpr1869 ; GCN-NEXT: ; implicit-def: $vgpr2070 ; GCN-NEXT: ; implicit-def: $vgpr2471 ; GCN-NEXT: v_add_u32_e32 v247, v19, v2472 ; GCN-NEXT: v_add_u32_e32 v248, v19, v1673 ; GCN-NEXT: v_add_u32_e32 v249, v19, v1874 ; GCN-NEXT: v_add_u32_e32 v250, v19, v2075 ; GCN-NEXT: ; implicit-def: $vgpr128_vgpr129_vgpr130_vgpr13176 ; GCN-NEXT: ; implicit-def: $sgpr1477 ; GCN-NEXT: ; implicit-def: $vgpr19678 ; GCN-NEXT: ; implicit-def: $sgpr12_sgpr1379 ; GCN-NEXT: ; implicit-def: $vgpr21180 ; GCN-NEXT: v_max_f32_e32 v212, v211, v21181 ; GCN-NEXT: ; implicit-def: $vgpr19882 ; GCN-NEXT: ; implicit-def: $vgpr0_vgpr1_vgpr2_vgpr3_vgpr4_vgpr5_vgpr6_vgpr7_vgpr8_vgpr9_vgpr10_vgpr11_vgpr12_vgpr13_vgpr14_vgpr1583 ; GCN-NEXT: ; implicit-def: $vgpr3284 ; GCN-NEXT: ; implicit-def: $vgpr3385 ; GCN-NEXT: ; implicit-def: $vgpr3486 ; GCN-NEXT: v_add_u32_e32 v210, v19, v3487 ; GCN-NEXT: v_add_u32_e32 v206, v19, v3388 ; GCN-NEXT: v_add_u32_e32 v205, v19, v3289 ; GCN-NEXT: ; implicit-def: $vgpr32_vgpr33_vgpr34_vgpr35_vgpr36_vgpr37_vgpr38_vgpr39_vgpr40_vgpr41_vgpr42_vgpr43_vgpr44_vgpr45_vgpr46_vgpr4790 ; GCN-NEXT: ; implicit-def: $vgpr2191 ; GCN-NEXT: ; implicit-def: $vgpr2292 ; GCN-NEXT: ; implicit-def: $vgpr2393 ; GCN-NEXT: ; implicit-def: $vgpr3094 ; GCN-NEXT: ; implicit-def: $vgpr3195 ; GCN-NEXT: v_add_u32_e32 v207, v19, v2196 ; GCN-NEXT: v_add_u32_e32 v208, v19, v2297 ; GCN-NEXT: v_add_u32_e32 v209, v19, v2398 ; GCN-NEXT: v_add_u32_e32 v203, v19, v3099 ; GCN-NEXT: v_add_u32_e32 v204, v19, v31100 ; GCN-NEXT: ; kill: killed $vgpr17101 ; GCN-NEXT: ; implicit-def: $vgpr16_vgpr17_vgpr18_vgpr19_vgpr20_vgpr21_vgpr22_vgpr23_vgpr24_vgpr25_vgpr26_vgpr27_vgpr28_vgpr29_vgpr30_vgpr31102 ; GCN-NEXT: ; implicit-def: $vgpr48_vgpr49_vgpr50_vgpr51_vgpr52_vgpr53_vgpr54_vgpr55_vgpr56_vgpr57_vgpr58_vgpr59_vgpr60_vgpr61_vgpr62_vgpr63103 ; GCN-NEXT: ; implicit-def: $vgpr197104 ; GCN-NEXT: ; iglp_opt mask(0x00000002)105 ; GCN-NEXT: buffer_wbl2 sc0 sc1106 ; GCN-NEXT: ds_write_b128 v230, v[64:67]107 ; GCN-NEXT: buffer_wbl2 sc0 sc1108 ; GCN-NEXT: s_waitcnt lgkmcnt(0)109 ; GCN-NEXT: ds_write_b128 v230, v[68:71] offset:1024110 ; GCN-NEXT: s_waitcnt lgkmcnt(0)111 ; GCN-NEXT: buffer_load_dwordx4 v[160:163], v226, s[8:11], 0 offen offset:64 sc0 sc1112 ; GCN-NEXT: s_waitcnt vmcnt(0)113 ; GCN-NEXT: buffer_inv sc0 sc1114 ; GCN-NEXT: buffer_load_dwordx4 v[164:167], v72, s[8:11], 0 offen sc0 sc1115 ; GCN-NEXT: s_waitcnt vmcnt(0)116 ; GCN-NEXT: buffer_inv sc0 sc1117 ; GCN-NEXT: ;;#ASMSTART118 ; GCN-NEXT: s_waitcnt vmcnt(8)119 ; GCN-NEXT: ;;#ASMEND120 ; GCN-NEXT: ds_read_b128 v[64:67], v213121 ; GCN-NEXT: s_waitcnt lgkmcnt(0)122 ; GCN-NEXT: buffer_inv sc0 sc1123 ; GCN-NEXT: v_mfma_f32_32x32x8_f16 v[112:127], v[64:65], v[152:153], 0124 ; GCN-NEXT: v_mfma_f32_32x32x8_f16 v[112:127], v[66:67], v[154:155], v[112:127]125 ; GCN-NEXT: ds_read_b128 v[64:67], v213 offset:512126 ; GCN-NEXT: s_waitcnt lgkmcnt(0)127 ; GCN-NEXT: buffer_inv sc0 sc1128 ; GCN-NEXT: v_mfma_f32_32x32x8_f16 v[96:111], v[64:65], v[152:153], 0129 ; GCN-NEXT: v_mfma_f32_32x32x8_f16 v[96:111], v[66:67], v[154:155], v[96:111]130 ; GCN-NEXT: ds_read_b128 v[64:67], v213 offset:1024131 ; GCN-NEXT: s_waitcnt lgkmcnt(0)132 ; GCN-NEXT: buffer_inv sc0 sc1133 ; GCN-NEXT: ds_read_b128 v[168:171], v213 offset:1536134 ; GCN-NEXT: s_waitcnt lgkmcnt(0)135 ; GCN-NEXT: buffer_inv sc0 sc1136 ; GCN-NEXT: ds_read_b128 v[172:175], v246137 ; GCN-NEXT: s_waitcnt lgkmcnt(0)138 ; GCN-NEXT: buffer_inv sc0 sc1139 ; GCN-NEXT: ds_read_b128 v[176:179], v246 offset:512140 ; GCN-NEXT: s_waitcnt lgkmcnt(0)141 ; GCN-NEXT: buffer_inv sc0 sc1142 ; GCN-NEXT: ds_read_b128 v[180:183], v246 offset:1024143 ; GCN-NEXT: s_waitcnt lgkmcnt(0)144 ; GCN-NEXT: buffer_inv sc0 sc1145 ; GCN-NEXT: ds_read_b128 v[184:187], v246 offset:1536146 ; GCN-NEXT: s_waitcnt lgkmcnt(0)147 ; GCN-NEXT: buffer_inv sc0 sc1148 ; GCN-NEXT: ;;#ASMSTART149 ; GCN-NEXT: s_waitcnt vmcnt(8)150 ; GCN-NEXT: ;;#ASMEND151 ; GCN-NEXT: v_mfma_f32_32x32x8_f16 v[80:95], v[64:65], v[152:153], 0152 ; GCN-NEXT: buffer_wbl2 sc0 sc1153 ; GCN-NEXT: ds_write_b128 v230, v[160:163]154 ; GCN-NEXT: v_mfma_f32_32x32x8_f16 v[80:95], v[66:67], v[154:155], v[80:95]155 ; GCN-NEXT: buffer_wbl2 sc0 sc1156 ; GCN-NEXT: s_waitcnt lgkmcnt(0)157 ; GCN-NEXT: ds_write_b128 v230, v[164:167] offset:1024158 ; GCN-NEXT: v_mfma_f32_32x32x8_f16 v[64:79], v[168:169], v[152:153], 0159 ; GCN-NEXT: v_mfma_f32_32x32x8_f16 v[64:79], v[170:171], v[154:155], v[64:79]160 ; GCN-NEXT: s_waitcnt lgkmcnt(0)161 ; GCN-NEXT: buffer_load_dwordx4 v[152:155], v226, s[8:11], 0 offen offset:128 sc0 sc1162 ; GCN-NEXT: s_waitcnt vmcnt(0)163 ; GCN-NEXT: buffer_inv sc0 sc1164 ; GCN-NEXT: buffer_load_dwordx4 v[160:163], v188, s[8:11], 0 offen sc0 sc1165 ; GCN-NEXT: s_waitcnt vmcnt(0)166 ; GCN-NEXT: buffer_inv sc0 sc1167 ; GCN-NEXT: ;;#ASMSTART168 ; GCN-NEXT: s_waitcnt vmcnt(8)169 ; GCN-NEXT: ;;#ASMEND170 ; GCN-NEXT: ds_read_b128 v[188:191], v213171 ; GCN-NEXT: s_waitcnt lgkmcnt(0)172 ; GCN-NEXT: buffer_inv sc0 sc1173 ; GCN-NEXT: ds_read_b128 v[192:195], v213 offset:512174 ; GCN-NEXT: s_waitcnt lgkmcnt(0)175 ; GCN-NEXT: buffer_inv sc0 sc1176 ; GCN-NEXT: ds_read_b128 v[164:167], v213 offset:1024177 ; GCN-NEXT: s_waitcnt lgkmcnt(0)178 ; GCN-NEXT: buffer_inv sc0 sc1179 ; GCN-NEXT: ds_read_b128 v[214:217], v213 offset:1536180 ; GCN-NEXT: s_waitcnt lgkmcnt(0)181 ; GCN-NEXT: buffer_inv sc0 sc1182 ; GCN-NEXT: v_mfma_f32_32x32x8_f16 v[112:127], v[172:173], v[156:157], v[112:127]183 ; GCN-NEXT: ds_read_b128 v[218:221], v246184 ; GCN-NEXT: s_waitcnt lgkmcnt(0)185 ; GCN-NEXT: buffer_inv sc0 sc1186 ; GCN-NEXT: ds_read_b128 v[222:225], v246 offset:512187 ; GCN-NEXT: s_waitcnt lgkmcnt(0)188 ; GCN-NEXT: buffer_inv sc0 sc1189 ; GCN-NEXT: ds_read_b128 v[168:171], v246 offset:1024190 ; GCN-NEXT: s_waitcnt lgkmcnt(0)191 ; GCN-NEXT: buffer_inv sc0 sc1192 ; GCN-NEXT: v_mfma_f32_32x32x8_f16 v[112:127], v[174:175], v[158:159], v[112:127]193 ; GCN-NEXT: v_mfma_f32_32x32x8_f16 v[112:127], v[188:189], v[144:145], v[112:127]194 ; GCN-NEXT: v_mfma_f32_32x32x8_f16 v[112:127], v[190:191], v[146:147], v[112:127]195 ; GCN-NEXT: ds_read_b128 v[188:191], v246 offset:1536196 ; GCN-NEXT: s_waitcnt lgkmcnt(0)197 ; GCN-NEXT: buffer_inv sc0 sc1198 ; GCN-NEXT: ;;#ASMSTART199 ; GCN-NEXT: s_waitcnt vmcnt(8)200 ; GCN-NEXT: ;;#ASMEND201 ; GCN-NEXT: buffer_wbl2 sc0 sc1202 ; GCN-NEXT: ds_write_b128 v230, v[152:155]203 ; GCN-NEXT: buffer_wbl2 sc0 sc1204 ; GCN-NEXT: s_waitcnt lgkmcnt(0)205 ; GCN-NEXT: ds_write_b128 v230, v[160:163] offset:1024206 ; GCN-NEXT: s_waitcnt lgkmcnt(0)207 ; GCN-NEXT: buffer_load_dwordx4 v[152:155], v226, s[8:11], 0 offen offset:192 sc0 sc1208 ; GCN-NEXT: s_waitcnt vmcnt(0)209 ; GCN-NEXT: buffer_inv sc0 sc1210 ; GCN-NEXT: v_mfma_f32_32x32x8_f16 v[64:79], v[184:185], v[156:157], v[64:79]211 ; GCN-NEXT: buffer_load_dwordx4 v[226:229], v227, s[8:11], 0 offen sc0 sc1212 ; GCN-NEXT: s_waitcnt vmcnt(0)213 ; GCN-NEXT: buffer_inv sc0 sc1214 ; GCN-NEXT: buffer_load_dwordx2 v[160:161], v231, s[0:3], 0 offen sc0 sc1215 ; GCN-NEXT: s_waitcnt vmcnt(0)216 ; GCN-NEXT: buffer_inv sc0 sc1217 ; GCN-NEXT: buffer_load_dwordx2 v[162:163], v232, s[0:3], 0 offen sc0 sc1218 ; GCN-NEXT: s_waitcnt vmcnt(0)219 ; GCN-NEXT: buffer_inv sc0 sc1220 ; GCN-NEXT: buffer_load_dwordx2 v[172:173], v233, s[0:3], 0 offen sc0 sc1221 ; GCN-NEXT: s_waitcnt vmcnt(0)222 ; GCN-NEXT: buffer_inv sc0 sc1223 ; GCN-NEXT: buffer_load_dwordx2 v[174:175], v234, s[0:3], 0 offen sc0 sc1224 ; GCN-NEXT: s_waitcnt vmcnt(0)225 ; GCN-NEXT: buffer_inv sc0 sc1226 ; GCN-NEXT: ;;#ASMSTART227 ; GCN-NEXT: s_waitcnt vmcnt(8)228 ; GCN-NEXT: ;;#ASMEND229 ; GCN-NEXT: v_mfma_f32_32x32x8_f16 v[64:79], v[186:187], v[158:159], v[64:79]230 ; GCN-NEXT: v_perm_b32 v238, v162, v160, s5231 ; GCN-NEXT: v_mfma_f32_32x32x8_f16 v[112:127], v[218:219], v[140:141], v[112:127]232 ; GCN-NEXT: v_perm_b32 v240, v162, v160, s7233 ; GCN-NEXT: v_perm_b32 v242, v163, v161, s5234 ; GCN-NEXT: v_perm_b32 v244, v163, v161, s7235 ; GCN-NEXT: ds_read_b128 v[160:163], v213236 ; GCN-NEXT: s_waitcnt lgkmcnt(0)237 ; GCN-NEXT: buffer_inv sc0 sc1238 ; GCN-NEXT: v_perm_b32 v239, v174, v172, s5239 ; GCN-NEXT: v_perm_b32 v241, v174, v172, s7240 ; GCN-NEXT: v_perm_b32 v243, v175, v173, s5241 ; GCN-NEXT: v_mfma_f32_32x32x8_f16 v[64:79], v[214:215], v[144:145], v[64:79]242 ; GCN-NEXT: v_perm_b32 v245, v175, v173, s7243 ; GCN-NEXT: v_mfma_f32_32x32x8_f16 v[96:111], v[176:177], v[156:157], v[96:111]244 ; GCN-NEXT: v_mfma_f32_32x32x8_f16 v[112:127], v[220:221], v[142:143], v[112:127]245 ; GCN-NEXT: ds_read_b128 v[218:221], v213 offset:512246 ; GCN-NEXT: s_waitcnt lgkmcnt(0)247 ; GCN-NEXT: buffer_inv sc0 sc1248 ; GCN-NEXT: ds_read_b128 v[172:175], v213 offset:1024249 ; GCN-NEXT: s_waitcnt lgkmcnt(0)250 ; GCN-NEXT: buffer_inv sc0 sc1251 ; GCN-NEXT: v_mfma_f32_32x32x8_f16 v[64:79], v[216:217], v[146:147], v[64:79]252 ; GCN-NEXT: v_mfma_f32_32x32x8_f16 v[96:111], v[178:179], v[158:159], v[96:111]253 ; GCN-NEXT: v_mfma_f32_32x32x8_f16 v[112:127], v[160:161], v[148:149], v[112:127]254 ; GCN-NEXT: v_mfma_f32_32x32x8_f16 v[64:79], v[188:189], v[140:141], v[64:79]255 ; GCN-NEXT: v_mfma_f32_32x32x8_f16 v[96:111], v[192:193], v[144:145], v[96:111]256 ; GCN-NEXT: v_mfma_f32_32x32x8_f16 v[112:127], v[162:163], v[150:151], v[112:127]257 ; GCN-NEXT: ds_read_b128 v[160:163], v213 offset:1536258 ; GCN-NEXT: s_waitcnt lgkmcnt(0)259 ; GCN-NEXT: buffer_inv sc0 sc1260 ; GCN-NEXT: ds_read_b128 v[184:187], v246261 ; GCN-NEXT: s_waitcnt lgkmcnt(0)262 ; GCN-NEXT: buffer_inv sc0 sc1263 ; GCN-NEXT: ds_read_b128 v[214:217], v246 offset:512264 ; GCN-NEXT: s_waitcnt lgkmcnt(0)265 ; GCN-NEXT: buffer_inv sc0 sc1266 ; GCN-NEXT: ds_read_b128 v[176:179], v246 offset:1024267 ; GCN-NEXT: s_waitcnt lgkmcnt(0)268 ; GCN-NEXT: buffer_inv sc0 sc1269 ; GCN-NEXT: v_mfma_f32_32x32x8_f16 v[64:79], v[190:191], v[142:143], v[64:79]270 ; GCN-NEXT: v_mfma_f32_32x32x8_f16 v[96:111], v[194:195], v[146:147], v[96:111]271 ; GCN-NEXT: v_mfma_f32_32x32x8_f16 v[64:79], v[160:161], v[148:149], v[64:79]272 ; GCN-NEXT: v_mfma_f32_32x32x8_f16 v[80:95], v[180:181], v[156:157], v[80:95]273 ; GCN-NEXT: v_mfma_f32_32x32x8_f16 v[112:127], v[184:185], v[136:137], v[112:127]274 ; GCN-NEXT: v_mfma_f32_32x32x8_f16 v[96:111], v[222:223], v[140:141], v[96:111]275 ; GCN-NEXT: v_mfma_f32_32x32x8_f16 v[64:79], v[162:163], v[150:151], v[64:79]276 ; GCN-NEXT: ds_read_b128 v[160:163], v246 offset:1536277 ; GCN-NEXT: s_waitcnt lgkmcnt(0)278 ; GCN-NEXT: buffer_inv sc0 sc1279 ; GCN-NEXT: ;;#ASMSTART280 ; GCN-NEXT: s_waitcnt vmcnt(8)281 ; GCN-NEXT: ;;#ASMEND282 ; GCN-NEXT: buffer_wbl2 sc0 sc1283 ; GCN-NEXT: ds_write_b128 v230, v[152:155]284 ; GCN-NEXT: buffer_wbl2 sc0 sc1285 ; GCN-NEXT: s_waitcnt lgkmcnt(0)286 ; GCN-NEXT: ds_write_b128 v230, v[226:229] offset:1024287 ; GCN-NEXT: ;;#ASMSTART288 ; GCN-NEXT: s_waitcnt vmcnt(8)289 ; GCN-NEXT: ;;#ASMEND290 ; GCN-NEXT: v_mfma_f32_32x32x8_f16 v[80:95], v[182:183], v[158:159], v[80:95]291 ; GCN-NEXT: s_waitcnt lgkmcnt(0)292 ; GCN-NEXT: ds_read_b128 v[156:159], v213293 ; GCN-NEXT: s_waitcnt lgkmcnt(0)294 ; GCN-NEXT: buffer_inv sc0 sc1295 ; GCN-NEXT: ds_read_b128 v[226:229], v213 offset:512296 ; GCN-NEXT: s_waitcnt lgkmcnt(0)297 ; GCN-NEXT: buffer_inv sc0 sc1298 ; GCN-NEXT: ds_read_b128 v[180:183], v213 offset:1024299 ; GCN-NEXT: s_waitcnt lgkmcnt(0)300 ; GCN-NEXT: buffer_inv sc0 sc1301 ; GCN-NEXT: ds_read_b128 v[152:155], v213 offset:1536302 ; GCN-NEXT: s_waitcnt lgkmcnt(0)303 ; GCN-NEXT: buffer_inv sc0 sc1304 ; GCN-NEXT: ds_read_b128 v[230:233], v246305 ; GCN-NEXT: s_waitcnt lgkmcnt(0)306 ; GCN-NEXT: buffer_inv sc0 sc1307 ; GCN-NEXT: ds_read_b128 v[234:237], v246 offset:512308 ; GCN-NEXT: s_waitcnt lgkmcnt(0)309 ; GCN-NEXT: buffer_inv sc0 sc1310 ; GCN-NEXT: v_mfma_f32_32x32x8_f16 v[112:127], v[186:187], v[138:139], v[112:127]311 ; GCN-NEXT: ds_read_b128 v[184:187], v246 offset:1024312 ; GCN-NEXT: s_waitcnt lgkmcnt(0)313 ; GCN-NEXT: buffer_inv sc0 sc1314 ; GCN-NEXT: v_mfma_f32_32x32x8_f16 v[96:111], v[224:225], v[142:143], v[96:111]315 ; GCN-NEXT: v_mfma_f32_32x32x8_f16 v[112:127], v[156:157], v[132:133], v[112:127]316 ; GCN-NEXT: v_mfma_f32_32x32x8_f16 v[96:111], v[218:219], v[148:149], v[96:111]317 ; GCN-NEXT: v_mfma_f32_32x32x8_f16 v[112:127], v[158:159], v[134:135], v[112:127]318 ; GCN-NEXT: ds_read_b128 v[156:159], v246 offset:1536319 ; GCN-NEXT: s_waitcnt lgkmcnt(0)320 ; GCN-NEXT: buffer_inv sc0 sc1321 ; GCN-NEXT: ;;#ASMSTART322 ; GCN-NEXT: s_waitcnt vmcnt(8)323 ; GCN-NEXT: ;;#ASMEND324 ; GCN-NEXT: buffer_wbl2 sc0 sc1325 ; GCN-NEXT: ds_write_b64 v199, v[238:239]326 ; GCN-NEXT: buffer_wbl2 sc0 sc1327 ; GCN-NEXT: s_waitcnt lgkmcnt(0)328 ; GCN-NEXT: ds_write_b64 v200, v[240:241]329 ; GCN-NEXT: buffer_wbl2 sc0 sc1330 ; GCN-NEXT: s_waitcnt lgkmcnt(0)331 ; GCN-NEXT: ds_write_b64 v201, v[242:243]332 ; GCN-NEXT: buffer_wbl2 sc0 sc1333 ; GCN-NEXT: s_waitcnt lgkmcnt(0)334 ; GCN-NEXT: ds_write_b64 v202, v[244:245]335 ; GCN-NEXT: s_waitcnt lgkmcnt(0)336 ; GCN-NEXT: buffer_load_dwordx2 v[192:193], v247, s[0:3], 0 offen sc0 sc1337 ; GCN-NEXT: s_waitcnt vmcnt(0)338 ; GCN-NEXT: buffer_inv sc0 sc1339 ; GCN-NEXT: v_mfma_f32_32x32x8_f16 v[96:111], v[220:221], v[150:151], v[96:111]340 ; GCN-NEXT: buffer_load_dwordx2 v[194:195], v248, s[0:3], 0 offen sc0 sc1341 ; GCN-NEXT: s_waitcnt vmcnt(0)342 ; GCN-NEXT: buffer_inv sc0 sc1343 ; GCN-NEXT: buffer_load_dwordx2 v[218:219], v249, s[0:3], 0 offen sc0 sc1344 ; GCN-NEXT: s_waitcnt vmcnt(0)345 ; GCN-NEXT: buffer_inv sc0 sc1346 ; GCN-NEXT: buffer_load_dwordx2 v[220:221], v250, s[0:3], 0 offen sc0 sc1347 ; GCN-NEXT: s_waitcnt vmcnt(0)348 ; GCN-NEXT: buffer_inv sc0 sc1349 ; GCN-NEXT: ;;#ASMSTART350 ; GCN-NEXT: s_waitcnt vmcnt(8)351 ; GCN-NEXT: ;;#ASMEND352 ; GCN-NEXT: v_perm_b32 v188, v194, v192, s5353 ; GCN-NEXT: v_mfma_f32_32x32x8_f16 v[80:95], v[164:165], v[144:145], v[80:95]354 ; GCN-NEXT: v_perm_b32 v189, v220, v218, s5355 ; GCN-NEXT: v_perm_b32 v191, v220, v218, s7356 ; GCN-NEXT: v_perm_b32 v190, v194, v192, s7357 ; GCN-NEXT: v_perm_b32 v192, v195, v193, s5358 ; GCN-NEXT: v_perm_b32 v194, v195, v193, s7359 ; GCN-NEXT: v_perm_b32 v193, v221, v219, s5360 ; GCN-NEXT: v_perm_b32 v195, v221, v219, s7361 ; GCN-NEXT: v_mfma_f32_32x32x8_f16 v[80:95], v[166:167], v[146:147], v[80:95]362 ; GCN-NEXT: v_mfma_f32_32x32x8_f16 v[80:95], v[168:169], v[140:141], v[80:95]363 ; GCN-NEXT: v_mfma_f32_32x32x8_f16 v[80:95], v[170:171], v[142:143], v[80:95]364 ; GCN-NEXT: v_mfma_f32_32x32x8_f16 v[80:95], v[172:173], v[148:149], v[80:95]365 ; GCN-NEXT: v_mfma_f32_32x32x8_f16 v[96:111], v[214:215], v[136:137], v[96:111]366 ; GCN-NEXT: v_mfma_f32_32x32x8_f16 v[80:95], v[174:175], v[150:151], v[80:95]367 ; GCN-NEXT: v_mfma_f32_32x32x8_f16 v[96:111], v[216:217], v[138:139], v[96:111]368 ; GCN-NEXT: v_mfma_f32_32x32x8_f16 v[80:95], v[176:177], v[136:137], v[80:95]369 ; GCN-NEXT: v_mfma_f32_32x32x8_f16 v[96:111], v[226:227], v[132:133], v[96:111]370 ; GCN-NEXT: v_mfma_f32_32x32x8_f16 v[80:95], v[178:179], v[138:139], v[80:95]371 ; GCN-NEXT: v_mfma_f32_32x32x8_f16 v[64:79], v[160:161], v[136:137], v[64:79]372 ; GCN-NEXT: v_mfma_f32_32x32x8_f16 v[112:127], v[230:231], v[128:129], v[112:127]373 ; GCN-NEXT: v_mfma_f32_32x32x8_f16 v[96:111], v[228:229], v[134:135], v[96:111]374 ; GCN-NEXT: v_mfma_f32_32x32x8_f16 v[80:95], v[180:181], v[132:133], v[80:95]375 ; GCN-NEXT: v_mfma_f32_32x32x8_f16 v[64:79], v[162:163], v[138:139], v[64:79]376 ; GCN-NEXT: v_mfma_f32_32x32x8_f16 v[112:127], v[232:233], v[130:131], v[112:127]377 ; GCN-NEXT: v_mfma_f32_32x32x8_f16 v[96:111], v[234:235], v[128:129], v[96:111]378 ; GCN-NEXT: s_nop 9379 ; GCN-NEXT: v_mul_f32_e32 v213, s4, v112380 ; GCN-NEXT: v_mul_f32_e32 v218, s4, v113381 ; GCN-NEXT: v_max3_f32 v213, v213, s14, v218382 ; GCN-NEXT: v_mul_f32_e32 v218, s4, v114383 ; GCN-NEXT: v_mul_f32_e32 v219, s4, v115384 ; GCN-NEXT: v_max3_f32 v213, v213, v218, v219385 ; GCN-NEXT: v_mul_f32_e32 v218, s4, v116386 ; GCN-NEXT: v_mfma_f32_32x32x8_f16 v[80:95], v[182:183], v[134:135], v[80:95]387 ; GCN-NEXT: v_mul_f32_e32 v219, s4, v117388 ; GCN-NEXT: v_max3_f32 v213, v213, v218, v219389 ; GCN-NEXT: v_mul_f32_e32 v218, s4, v118390 ; GCN-NEXT: v_mul_f32_e32 v219, s4, v119391 ; GCN-NEXT: v_max3_f32 v213, v213, v218, v219392 ; GCN-NEXT: v_mul_f32_e32 v218, s4, v120393 ; GCN-NEXT: v_mul_f32_e32 v219, s4, v121394 ; GCN-NEXT: v_mfma_f32_32x32x8_f16 v[64:79], v[152:153], v[132:133], v[64:79]395 ; GCN-NEXT: v_max3_f32 v213, v213, v218, v219396 ; GCN-NEXT: v_mul_f32_e32 v218, s4, v122397 ; GCN-NEXT: v_mul_f32_e32 v219, s4, v123398 ; GCN-NEXT: v_max3_f32 v213, v213, v218, v219399 ; GCN-NEXT: v_mul_f32_e32 v218, s4, v124400 ; GCN-NEXT: v_mul_f32_e32 v219, s4, v125401 ; GCN-NEXT: v_max3_f32 v213, v213, v218, v219402 ; GCN-NEXT: v_mfma_f32_32x32x8_f16 v[96:111], v[236:237], v[130:131], v[96:111]403 ; GCN-NEXT: v_mul_f32_e32 v218, s4, v126404 ; GCN-NEXT: v_mul_f32_e32 v219, s4, v127405 ; GCN-NEXT: v_max3_f32 v213, v213, v218, v219406 ; GCN-NEXT: v_mfma_f32_32x32x8_f16 v[80:95], v[184:185], v[128:129], v[80:95]407 ; GCN-NEXT: s_nop 6408 ; GCN-NEXT: v_mul_f32_e32 v214, s4, v96409 ; GCN-NEXT: v_mul_f32_e32 v215, s4, v97410 ; GCN-NEXT: v_max3_f32 v213, v213, v214, v215411 ; GCN-NEXT: v_mul_f32_e32 v214, s4, v98412 ; GCN-NEXT: v_mul_f32_e32 v215, s4, v99413 ; GCN-NEXT: v_max3_f32 v213, v213, v214, v215414 ; GCN-NEXT: v_mul_f32_e32 v214, s4, v100415 ; GCN-NEXT: v_mfma_f32_32x32x8_f16 v[64:79], v[154:155], v[134:135], v[64:79]416 ; GCN-NEXT: v_mul_f32_e32 v215, s4, v101417 ; GCN-NEXT: v_max3_f32 v213, v213, v214, v215418 ; GCN-NEXT: v_mul_f32_e32 v214, s4, v102419 ; GCN-NEXT: v_mul_f32_e32 v215, s4, v103420 ; GCN-NEXT: v_max3_f32 v213, v213, v214, v215421 ; GCN-NEXT: v_mul_f32_e32 v214, s4, v104422 ; GCN-NEXT: v_mul_f32_e32 v215, s4, v105423 ; GCN-NEXT: v_mfma_f32_32x32x8_f16 v[80:95], v[186:187], v[130:131], v[80:95]424 ; GCN-NEXT: v_max3_f32 v213, v213, v214, v215425 ; GCN-NEXT: v_mul_f32_e32 v214, s4, v106426 ; GCN-NEXT: v_mul_f32_e32 v215, s4, v107427 ; GCN-NEXT: v_max3_f32 v213, v213, v214, v215428 ; GCN-NEXT: v_mul_f32_e32 v214, s4, v108429 ; GCN-NEXT: v_mul_f32_e32 v215, s4, v109430 ; GCN-NEXT: v_max3_f32 v213, v213, v214, v215431 ; GCN-NEXT: v_mfma_f32_32x32x8_f16 v[64:79], v[156:157], v[128:129], v[64:79]432 ; GCN-NEXT: v_mul_f32_e32 v214, s4, v110433 ; GCN-NEXT: v_mul_f32_e32 v215, s4, v111434 ; GCN-NEXT: v_max3_f32 v213, v213, v214, v215435 ; GCN-NEXT: v_mul_f32_e32 v140, s4, v80436 ; GCN-NEXT: v_mul_f32_e32 v141, s4, v81437 ; GCN-NEXT: v_max3_f32 v140, v213, v140, v141438 ; GCN-NEXT: v_mul_f32_e32 v141, s4, v82439 ; GCN-NEXT: v_mfma_f32_32x32x8_f16 v[64:79], v[158:159], v[130:131], v[64:79]440 ; GCN-NEXT: v_mul_f32_e32 v142, s4, v83441 ; GCN-NEXT: v_max3_f32 v140, v140, v141, v142442 ; GCN-NEXT: v_mul_f32_e32 v141, s4, v84443 ; GCN-NEXT: v_mul_f32_e32 v142, s4, v85444 ; GCN-NEXT: v_max3_f32 v140, v140, v141, v142445 ; GCN-NEXT: v_mul_f32_e32 v141, s4, v86446 ; GCN-NEXT: v_mul_f32_e32 v142, s4, v87447 ; GCN-NEXT: v_max3_f32 v140, v140, v141, v142448 ; GCN-NEXT: v_mul_f32_e32 v141, s4, v88449 ; GCN-NEXT: v_mul_f32_e32 v142, s4, v89450 ; GCN-NEXT: v_max3_f32 v140, v140, v141, v142451 ; GCN-NEXT: v_mul_f32_e32 v141, s4, v90452 ; GCN-NEXT: v_mul_f32_e32 v142, s4, v91453 ; GCN-NEXT: v_max3_f32 v140, v140, v141, v142454 ; GCN-NEXT: v_mul_f32_e32 v141, s4, v92455 ; GCN-NEXT: v_mul_f32_e32 v142, s4, v93456 ; GCN-NEXT: v_max3_f32 v140, v140, v141, v142457 ; GCN-NEXT: v_mul_f32_e32 v141, s4, v94458 ; GCN-NEXT: v_mul_f32_e32 v142, s4, v95459 ; GCN-NEXT: v_max3_f32 v140, v140, v141, v142460 ; GCN-NEXT: v_mul_f32_e32 v128, s4, v64461 ; GCN-NEXT: v_mul_f32_e32 v129, s4, v65462 ; GCN-NEXT: v_max3_f32 v128, v140, v128, v129463 ; GCN-NEXT: v_mul_f32_e32 v129, s4, v66464 ; GCN-NEXT: v_mul_f32_e32 v130, s4, v67465 ; GCN-NEXT: v_max3_f32 v128, v128, v129, v130466 ; GCN-NEXT: v_mul_f32_e32 v129, s4, v68467 ; GCN-NEXT: v_mul_f32_e32 v130, s4, v69468 ; GCN-NEXT: v_max3_f32 v128, v128, v129, v130469 ; GCN-NEXT: v_mul_f32_e32 v129, s4, v70470 ; GCN-NEXT: v_mul_f32_e32 v130, s4, v71471 ; GCN-NEXT: v_max3_f32 v128, v128, v129, v130472 ; GCN-NEXT: v_mul_f32_e32 v129, s4, v72473 ; GCN-NEXT: v_mul_f32_e32 v130, s4, v73474 ; GCN-NEXT: v_max3_f32 v128, v128, v129, v130475 ; GCN-NEXT: v_mul_f32_e32 v129, s4, v74476 ; GCN-NEXT: v_mul_f32_e32 v130, s4, v75477 ; GCN-NEXT: v_max3_f32 v128, v128, v129, v130478 ; GCN-NEXT: v_mul_f32_e32 v129, s4, v76479 ; GCN-NEXT: v_mul_f32_e32 v130, s4, v77480 ; GCN-NEXT: v_max3_f32 v128, v128, v129, v130481 ; GCN-NEXT: v_mul_f32_e32 v129, s4, v78482 ; GCN-NEXT: v_mul_f32_e32 v130, s4, v79483 ; GCN-NEXT: v_max3_f32 v128, v128, v129, v130484 ; GCN-NEXT: ds_bpermute_b32 v129, v196, v128485 ; GCN-NEXT: s_waitcnt lgkmcnt(0)486 ; GCN-NEXT: ds_read_b128 v[130:133], v198487 ; GCN-NEXT: s_waitcnt lgkmcnt(0)488 ; GCN-NEXT: buffer_inv sc0 sc1489 ; GCN-NEXT: ds_read_b128 v[134:137], v198 offset:576490 ; GCN-NEXT: s_waitcnt lgkmcnt(0)491 ; GCN-NEXT: buffer_inv sc0 sc1492 ; GCN-NEXT: v_max_f32_e32 v129, v129, v129493 ; GCN-NEXT: v_max_f32_e32 v128, v128, v129494 ; GCN-NEXT: ds_bpermute_b32 v129, v196, v128495 ; GCN-NEXT: s_waitcnt lgkmcnt(0)496 ; GCN-NEXT: v_cndmask_b32_e64 v128, v129, v128, s[12:13]497 ; GCN-NEXT: v_max_f32_e32 v128, v128, v128498 ; GCN-NEXT: v_max_f32_e32 v128, v212, v128499 ; GCN-NEXT: v_fma_f32 v113, s4, v113, -v128500 ; GCN-NEXT: v_mul_f32_e32 v138, 0x3fb8aa3b, v113501 ; GCN-NEXT: v_fma_f32 v113, s4, v114, -v128502 ; GCN-NEXT: v_mul_f32_e32 v139, 0x3fb8aa3b, v113503 ; GCN-NEXT: v_fma_f32 v113, s4, v115, -v128504 ; GCN-NEXT: v_mul_f32_e32 v140, 0x3fb8aa3b, v113505 ; GCN-NEXT: v_fma_f32 v113, s4, v116, -v128506 ; GCN-NEXT: v_mul_f32_e32 v141, 0x3fb8aa3b, v113507 ; GCN-NEXT: v_fma_f32 v113, s4, v117, -v128508 ; GCN-NEXT: v_mul_f32_e32 v142, 0x3fb8aa3b, v113509 ; GCN-NEXT: v_fma_f32 v113, s4, v118, -v128510 ; GCN-NEXT: v_fma_f32 v112, s4, v112, -v128511 ; GCN-NEXT: v_mul_f32_e32 v143, 0x3fb8aa3b, v113512 ; GCN-NEXT: v_fma_f32 v113, s4, v119, -v128513 ; GCN-NEXT: v_fma_f32 v118, s4, v120, -v128514 ; GCN-NEXT: v_fma_f32 v120, s4, v121, -v128515 ; GCN-NEXT: v_mul_f32_e32 v112, 0x3fb8aa3b, v112516 ; GCN-NEXT: v_mul_f32_e32 v144, 0x3fb8aa3b, v113517 ; GCN-NEXT: v_mul_f32_e32 v149, 0x3fb8aa3b, v120518 ; GCN-NEXT: v_fma_f32 v120, s4, v122, -v128519 ; GCN-NEXT: v_exp_f32_e32 v114, v138520 ; GCN-NEXT: v_exp_f32_e32 v115, v139521 ; GCN-NEXT: v_exp_f32_e32 v116, v140522 ; GCN-NEXT: v_exp_f32_e32 v117, v141523 ; GCN-NEXT: v_mul_f32_e32 v148, 0x3fb8aa3b, v118524 ; GCN-NEXT: v_exp_f32_e32 v118, v142525 ; GCN-NEXT: v_mul_f32_e32 v150, 0x3fb8aa3b, v120526 ; GCN-NEXT: v_exp_f32_e32 v120, v144527 ; GCN-NEXT: v_exp_f32_e32 v113, v112528 ; GCN-NEXT: v_cvt_f16_f32_e32 v119, v114529 ; GCN-NEXT: v_cvt_f16_f32_e32 v121, v116530 ; GCN-NEXT: v_sub_f32_e32 v129, v211, v128531 ; GCN-NEXT: v_cvt_f16_f32_e32 v112, v113532 ; GCN-NEXT: v_mul_f32_e32 v129, 0x3fb8aa3b, v129533 ; GCN-NEXT: ds_read_b128 v[138:141], v198 offset:1152534 ; GCN-NEXT: s_waitcnt lgkmcnt(0)535 ; GCN-NEXT: buffer_inv sc0 sc1536 ; GCN-NEXT: v_fma_f32 v122, s4, v123, -v128537 ; GCN-NEXT: v_pack_b32_f16 v146, v112, v119538 ; GCN-NEXT: v_cvt_f16_f32_e32 v112, v115539 ; GCN-NEXT: v_mul_f32_e32 v151, 0x3fb8aa3b, v122540 ; GCN-NEXT: v_cvt_f16_f32_e32 v123, v117541 ; GCN-NEXT: v_fma_f32 v122, s4, v124, -v128542 ; GCN-NEXT: v_pack_b32_f16 v147, v112, v121543 ; GCN-NEXT: v_exp_f32_e32 v112, v129544 ; GCN-NEXT: v_cvt_f16_f32_e32 v124, v118545 ; GCN-NEXT: v_mul_f32_e32 v129, 0x3fb8aa3b, v122546 ; GCN-NEXT: v_fma_f32 v125, s4, v125, -v128547 ; GCN-NEXT: v_pk_mul_f32 v[0:1], v[0:1], v[112:113] op_sel_hi:[1,0]548 ; GCN-NEXT: v_pk_mul_f32 v[2:3], v[2:3], v[112:113] op_sel_hi:[1,0]549 ; GCN-NEXT: v_pk_mul_f32 v[4:5], v[4:5], v[112:113] op_sel_hi:[1,0]550 ; GCN-NEXT: v_pk_mul_f32 v[6:7], v[6:7], v[112:113] op_sel_hi:[1,0]551 ; GCN-NEXT: v_pk_mul_f32 v[8:9], v[8:9], v[112:113] op_sel_hi:[1,0]552 ; GCN-NEXT: v_pk_mul_f32 v[10:11], v[10:11], v[112:113] op_sel_hi:[1,0]553 ; GCN-NEXT: v_pk_mul_f32 v[12:13], v[12:13], v[112:113] op_sel_hi:[1,0]554 ; GCN-NEXT: v_pk_mul_f32 v[14:15], v[14:15], v[112:113] op_sel_hi:[1,0]555 ; GCN-NEXT: v_pk_mul_f32 v[32:33], v[32:33], v[112:113] op_sel_hi:[1,0]556 ; GCN-NEXT: v_pk_mul_f32 v[34:35], v[34:35], v[112:113] op_sel_hi:[1,0]557 ; GCN-NEXT: v_mfma_f32_32x32x8_f16 v[0:15], v[130:131], v[146:147], v[0:15]558 ; GCN-NEXT: v_exp_f32_e32 v119, v143559 ; GCN-NEXT: ds_read_b128 v[142:145], v198 offset:1728560 ; GCN-NEXT: s_waitcnt lgkmcnt(0)561 ; GCN-NEXT: buffer_inv sc0 sc1562 ; GCN-NEXT: v_pk_mul_f32 v[36:37], v[36:37], v[112:113] op_sel_hi:[1,0]563 ; GCN-NEXT: v_pk_mul_f32 v[38:39], v[38:39], v[112:113] op_sel_hi:[1,0]564 ; GCN-NEXT: v_pk_mul_f32 v[40:41], v[40:41], v[112:113] op_sel_hi:[1,0]565 ; GCN-NEXT: v_pk_mul_f32 v[42:43], v[42:43], v[112:113] op_sel_hi:[1,0]566 ; GCN-NEXT: v_pk_mul_f32 v[44:45], v[44:45], v[112:113] op_sel_hi:[1,0]567 ; GCN-NEXT: v_pk_mul_f32 v[46:47], v[46:47], v[112:113] op_sel_hi:[1,0]568 ; GCN-NEXT: v_pk_mul_f32 v[16:17], v[16:17], v[112:113] op_sel_hi:[1,0]569 ; GCN-NEXT: v_pk_mul_f32 v[18:19], v[18:19], v[112:113] op_sel_hi:[1,0]570 ; GCN-NEXT: v_mfma_f32_32x32x8_f16 v[32:47], v[134:135], v[146:147], v[32:47]571 ; GCN-NEXT: v_mul_f32_e64 v20, v20, v112572 ; GCN-NEXT: v_mul_f32_e64 v21, v21, v112573 ; GCN-NEXT: v_mul_f32_e64 v22, v22, v112574 ; GCN-NEXT: v_mul_f32_e64 v23, v23, v112575 ; GCN-NEXT: v_mul_f32_e64 v24, v24, v112576 ; GCN-NEXT: v_mul_f32_e64 v25, v25, v112577 ; GCN-NEXT: v_pk_mul_f32 v[26:27], v[26:27], v[112:113] op_sel_hi:[1,0]578 ; GCN-NEXT: v_pk_mul_f32 v[28:29], v[28:29], v[112:113] op_sel_hi:[1,0]579 ; GCN-NEXT: v_pk_mul_f32 v[30:31], v[30:31], v[112:113] op_sel_hi:[1,0]580 ; GCN-NEXT: v_pk_mul_f32 v[48:49], v[48:49], v[112:113] op_sel_hi:[1,0]581 ; GCN-NEXT: v_pk_mul_f32 v[50:51], v[50:51], v[112:113] op_sel_hi:[1,0]582 ; GCN-NEXT: v_pk_mul_f32 v[52:53], v[52:53], v[112:113] op_sel_hi:[1,0]583 ; GCN-NEXT: v_pk_mul_f32 v[54:55], v[54:55], v[112:113] op_sel_hi:[1,0]584 ; GCN-NEXT: v_pk_mul_f32 v[56:57], v[56:57], v[112:113] op_sel_hi:[1,0]585 ; GCN-NEXT: v_pk_mul_f32 v[58:59], v[58:59], v[112:113] op_sel_hi:[1,0]586 ; GCN-NEXT: v_pk_mul_f32 v[60:61], v[60:61], v[112:113] op_sel_hi:[1,0]587 ; GCN-NEXT: v_pk_mul_f32 v[62:63], v[62:63], v[112:113] op_sel_hi:[1,0]588 ; GCN-NEXT: v_pack_b32_f16 v134, v123, v124589 ; GCN-NEXT: v_cvt_f16_f32_e32 v130, v119590 ; GCN-NEXT: v_fma_f32 v124, s4, v126, -v128591 ; GCN-NEXT: v_cvt_f16_f32_e32 v126, v120592 ; GCN-NEXT: v_exp_f32_e32 v121, v148593 ; GCN-NEXT: v_mfma_f32_32x32x8_f16 v[16:31], v[138:139], v[146:147], v[16:31]594 ; GCN-NEXT: v_exp_f32_e32 v122, v149595 ; GCN-NEXT: v_pack_b32_f16 v135, v130, v126596 ; GCN-NEXT: v_mul_f32_e32 v138, 0x3fb8aa3b, v124597 ; GCN-NEXT: v_cvt_f16_f32_e32 v126, v121598 ; GCN-NEXT: v_mul_f32_e32 v125, 0x3fb8aa3b, v125599 ; GCN-NEXT: v_fma_f32 v139, s4, v96, -v128600 ; GCN-NEXT: v_fma_f32 v127, s4, v127, -v128601 ; GCN-NEXT: v_mfma_f32_32x32x8_f16 v[48:63], v[142:143], v[146:147], v[48:63]602 ; GCN-NEXT: v_exp_f32_e32 v123, v150603 ; GCN-NEXT: v_mul_f32_e32 v127, 0x3fb8aa3b, v127604 ; GCN-NEXT: v_fma_f32 v143, s4, v101, -v128605 ; GCN-NEXT: v_fma_f32 v64, s4, v64, -v128606 ; GCN-NEXT: v_fma_f32 v65, s4, v65, -v128607 ; GCN-NEXT: v_fma_f32 v68, s4, v68, -v128608 ; GCN-NEXT: v_fma_f32 v69, s4, v69, -v128609 ; GCN-NEXT: v_mfma_f32_32x32x8_f16 v[0:15], v[132:133], v[134:135], v[0:15]610 ; GCN-NEXT: v_exp_f32_e32 v124, v151611 ; GCN-NEXT: ds_read_b128 v[130:133], v197612 ; GCN-NEXT: s_waitcnt lgkmcnt(0)613 ; GCN-NEXT: buffer_inv sc0 sc1614 ; GCN-NEXT: ds_read_b128 v[146:149], v197 offset:576615 ; GCN-NEXT: s_waitcnt lgkmcnt(0)616 ; GCN-NEXT: buffer_inv sc0 sc1617 ; GCN-NEXT: v_mfma_f32_32x32x8_f16 v[32:47], v[136:137], v[134:135], v[32:47]618 ; GCN-NEXT: v_cvt_f16_f32_e32 v136, v122619 ; GCN-NEXT: v_exp_f32_e32 v96, v129620 ; GCN-NEXT: v_fma_f32 v137, s4, v97, -v128621 ; GCN-NEXT: v_mul_f32_e32 v129, 0x3fb8aa3b, v139622 ; GCN-NEXT: v_pack_b32_f16 v126, v126, v136623 ; GCN-NEXT: v_cvt_f16_f32_e32 v136, v123624 ; GCN-NEXT: v_mfma_f32_32x32x8_f16 v[16:31], v[140:141], v[134:135], v[16:31]625 ; GCN-NEXT: v_exp_f32_e32 v97, v125626 ; GCN-NEXT: v_mul_f32_e32 v125, 0x3fb8aa3b, v137627 ; GCN-NEXT: v_fma_f32 v137, s4, v98, -v128628 ; GCN-NEXT: v_mul_f32_e32 v142, 0x3fb8aa3b, v137629 ; GCN-NEXT: v_mfma_f32_32x32x8_f16 v[48:63], v[144:145], v[134:135], v[48:63]630 ; GCN-NEXT: v_cvt_f16_f32_e32 v134, v124631 ; GCN-NEXT: v_fma_f32 v135, s4, v99, -v128632 ; GCN-NEXT: v_exp_f32_e32 v98, v138633 ; GCN-NEXT: v_exp_f32_e32 v99, v127634 ; GCN-NEXT: v_mul_f32_e32 v150, 0x3fb8aa3b, v135635 ; GCN-NEXT: v_pack_b32_f16 v127, v136, v134636 ; GCN-NEXT: ds_read_b128 v[134:137], v197 offset:1152637 ; GCN-NEXT: s_waitcnt lgkmcnt(0)638 ; GCN-NEXT: buffer_inv sc0 sc1639 ; GCN-NEXT: ds_read_b128 v[138:141], v197 offset:1728640 ; GCN-NEXT: s_waitcnt lgkmcnt(0)641 ; GCN-NEXT: buffer_inv sc0 sc1642 ; GCN-NEXT: v_mfma_f32_32x32x8_f16 v[0:15], v[130:131], v[126:127], v[0:15]643 ; GCN-NEXT: v_fma_f32 v131, s4, v100, -v128644 ; GCN-NEXT: v_cvt_f16_f32_e32 v130, v96645 ; GCN-NEXT: v_exp_f32_e32 v100, v129646 ; GCN-NEXT: v_mul_f32_e32 v129, 0x3fb8aa3b, v131647 ; GCN-NEXT: v_cvt_f16_f32_e32 v131, v97648 ; GCN-NEXT: ;;#ASMSTART649 ; GCN-NEXT: s_waitcnt vmcnt(8)650 ; GCN-NEXT: ;;#ASMEND651 ; GCN-NEXT: buffer_wbl2 sc0 sc1652 ; GCN-NEXT: ds_write_b64 v199, v[188:189]653 ; GCN-NEXT: buffer_wbl2 sc0 sc1654 ; GCN-NEXT: s_waitcnt lgkmcnt(0)655 ; GCN-NEXT: ds_write_b64 v200, v[190:191]656 ; GCN-NEXT: buffer_wbl2 sc0 sc1657 ; GCN-NEXT: s_waitcnt lgkmcnt(0)658 ; GCN-NEXT: ds_write_b64 v201, v[192:193]659 ; GCN-NEXT: buffer_wbl2 sc0 sc1660 ; GCN-NEXT: s_waitcnt lgkmcnt(0)661 ; GCN-NEXT: ds_write_b64 v202, v[194:195]662 ; GCN-NEXT: v_mfma_f32_32x32x8_f16 v[32:47], v[146:147], v[126:127], v[32:47]663 ; GCN-NEXT: v_exp_f32_e32 v101, v125664 ; GCN-NEXT: v_pack_b32_f16 v146, v130, v131665 ; GCN-NEXT: s_waitcnt lgkmcnt(0)666 ; GCN-NEXT: buffer_load_dwordx2 v[130:131], v210, s[0:3], 0 offen sc0 sc1667 ; GCN-NEXT: s_waitcnt vmcnt(0)668 ; GCN-NEXT: buffer_inv sc0 sc1669 ; GCN-NEXT: v_mul_f32_e32 v125, 0x3fb8aa3b, v143670 ; GCN-NEXT: v_cvt_f16_f32_e32 v147, v98671 ; GCN-NEXT: v_mfma_f32_32x32x8_f16 v[16:31], v[134:135], v[126:127], v[16:31]672 ; GCN-NEXT: v_fma_f32 v134, s4, v102, -v128673 ; GCN-NEXT: v_mul_f32_e32 v156, 0x3fb8aa3b, v134674 ; GCN-NEXT: buffer_load_dwordx2 v[134:135], v207, s[0:3], 0 offen sc0 sc1675 ; GCN-NEXT: s_waitcnt vmcnt(0)676 ; GCN-NEXT: buffer_inv sc0 sc1677 ; GCN-NEXT: v_exp_f32_e32 v102, v142678 ; GCN-NEXT: buffer_load_dwordx2 v[142:143], v208, s[0:3], 0 offen sc0 sc1679 ; GCN-NEXT: s_waitcnt vmcnt(0)680 ; GCN-NEXT: buffer_inv sc0 sc1681 ; GCN-NEXT: buffer_load_dwordx2 v[144:145], v209, s[0:3], 0 offen sc0 sc1682 ; GCN-NEXT: s_waitcnt vmcnt(0)683 ; GCN-NEXT: buffer_inv sc0 sc1684 ; GCN-NEXT: ;;#ASMSTART685 ; GCN-NEXT: s_waitcnt vmcnt(8)686 ; GCN-NEXT: ;;#ASMEND687 ; GCN-NEXT: v_mfma_f32_32x32x8_f16 v[48:63], v[138:139], v[126:127], v[48:63]688 ; GCN-NEXT: v_cvt_f16_f32_e32 v126, v99689 ; GCN-NEXT: v_fma_f32 v127, s4, v103, -v128690 ; GCN-NEXT: v_exp_f32_e32 v103, v150691 ; GCN-NEXT: v_fma_f32 v139, s4, v105, -v128692 ; GCN-NEXT: v_pack_b32_f16 v147, v147, v126693 ; GCN-NEXT: v_mul_f32_e32 v138, 0x3fb8aa3b, v127694 ; GCN-NEXT: v_perm_b32 v152, v135, v131, s5695 ; GCN-NEXT: v_perm_b32 v154, v135, v131, s7696 ; GCN-NEXT: v_fma_f32 v135, s4, v104, -v128697 ; GCN-NEXT: v_perm_b32 v126, v134, v130, s5698 ; GCN-NEXT: v_mfma_f32_32x32x8_f16 v[0:15], v[132:133], v[146:147], v[0:15]699 ; GCN-NEXT: v_perm_b32 v150, v134, v130, s7700 ; GCN-NEXT: v_cvt_f16_f32_e32 v134, v100701 ; GCN-NEXT: v_exp_f32_e32 v104, v129702 ; GCN-NEXT: v_mul_f32_e32 v129, 0x3fb8aa3b, v135703 ; GCN-NEXT: v_cvt_f16_f32_e32 v135, v101704 ; GCN-NEXT: ds_read_b128 v[130:133], v198705 ; GCN-NEXT: s_waitcnt lgkmcnt(0)706 ; GCN-NEXT: buffer_inv sc0 sc1707 ; GCN-NEXT: v_perm_b32 v127, v144, v142, s5708 ; GCN-NEXT: v_mfma_f32_32x32x8_f16 v[32:47], v[148:149], v[146:147], v[32:47]709 ; GCN-NEXT: v_pack_b32_f16 v148, v134, v135710 ; GCN-NEXT: v_fma_f32 v135, s4, v106, -v128711 ; GCN-NEXT: v_exp_f32_e32 v105, v125712 ; GCN-NEXT: v_cvt_f16_f32_e32 v134, v102713 ; GCN-NEXT: v_perm_b32 v151, v144, v142, s7714 ; GCN-NEXT: v_perm_b32 v153, v145, v143, s5715 ; GCN-NEXT: v_perm_b32 v155, v145, v143, s7716 ; GCN-NEXT: v_mfma_f32_32x32x8_f16 v[16:31], v[136:137], v[146:147], v[16:31]717 ; GCN-NEXT: v_exp_f32_e32 v106, v156718 ; GCN-NEXT: v_mul_f32_e32 v156, 0x3fb8aa3b, v135719 ; GCN-NEXT: v_cvt_f16_f32_e32 v135, v103720 ; GCN-NEXT: v_fma_f32 v136, s4, v107, -v128721 ; GCN-NEXT: ds_read_b128 v[142:145], v198 offset:576722 ; GCN-NEXT: s_waitcnt lgkmcnt(0)723 ; GCN-NEXT: buffer_inv sc0 sc1724 ; GCN-NEXT: v_mul_f32_e32 v125, 0x3fb8aa3b, v139725 ; GCN-NEXT: v_pack_b32_f16 v149, v134, v135726 ; GCN-NEXT: v_mfma_f32_32x32x8_f16 v[48:63], v[140:141], v[146:147], v[48:63]727 ; GCN-NEXT: v_mul_f32_e32 v146, 0x3fb8aa3b, v136728 ; GCN-NEXT: ds_read_b128 v[134:137], v198 offset:1152729 ; GCN-NEXT: s_waitcnt lgkmcnt(0)730 ; GCN-NEXT: buffer_inv sc0 sc1731 ; GCN-NEXT: v_exp_f32_e32 v107, v138732 ; GCN-NEXT: ds_read_b128 v[138:141], v198 offset:1728733 ; GCN-NEXT: s_waitcnt lgkmcnt(0)734 ; GCN-NEXT: buffer_inv sc0 sc1735 ; GCN-NEXT: v_mfma_f32_32x32x8_f16 v[0:15], v[130:131], v[148:149], v[0:15]736 ; GCN-NEXT: v_fma_f32 v131, s4, v108, -v128737 ; GCN-NEXT: v_cvt_f16_f32_e32 v130, v104738 ; GCN-NEXT: v_exp_f32_e32 v108, v129739 ; GCN-NEXT: v_mul_f32_e32 v129, 0x3fb8aa3b, v131740 ; GCN-NEXT: v_cvt_f16_f32_e32 v131, v105741 ; GCN-NEXT: v_mfma_f32_32x32x8_f16 v[32:47], v[142:143], v[148:149], v[32:47]742 ; GCN-NEXT: v_fma_f32 v142, s4, v109, -v128743 ; GCN-NEXT: v_exp_f32_e32 v109, v125744 ; GCN-NEXT: v_mul_f32_e32 v125, 0x3fb8aa3b, v142745 ; GCN-NEXT: v_pack_b32_f16 v142, v130, v131746 ; GCN-NEXT: v_fma_f32 v131, s4, v110, -v128747 ; GCN-NEXT: v_cvt_f16_f32_e32 v130, v106748 ; GCN-NEXT: v_mfma_f32_32x32x8_f16 v[16:31], v[134:135], v[148:149], v[16:31]749 ; GCN-NEXT: v_mul_f32_e32 v134, 0x3fb8aa3b, v131750 ; GCN-NEXT: v_cvt_f16_f32_e32 v131, v107751 ; GCN-NEXT: v_exp_f32_e32 v110, v156752 ; GCN-NEXT: v_fma_f32 v135, s4, v111, -v128753 ; GCN-NEXT: v_mul_f32_e32 v135, 0x3fb8aa3b, v135754 ; GCN-NEXT: v_pack_b32_f16 v143, v130, v131755 ; GCN-NEXT: v_mfma_f32_32x32x8_f16 v[48:63], v[138:139], v[148:149], v[48:63]756 ; GCN-NEXT: v_exp_f32_e32 v111, v146757 ; GCN-NEXT: v_fma_f32 v139, s4, v80, -v128758 ; GCN-NEXT: v_cvt_f16_f32_e32 v138, v108759 ; GCN-NEXT: v_mfma_f32_32x32x8_f16 v[0:15], v[132:133], v[142:143], v[0:15]760 ; GCN-NEXT: v_exp_f32_e32 v80, v129761 ; GCN-NEXT: ds_read_b128 v[130:133], v197762 ; GCN-NEXT: s_waitcnt lgkmcnt(0)763 ; GCN-NEXT: buffer_inv sc0 sc1764 ; GCN-NEXT: ds_read_b128 v[146:149], v197 offset:576765 ; GCN-NEXT: s_waitcnt lgkmcnt(0)766 ; GCN-NEXT: buffer_inv sc0 sc1767 ; GCN-NEXT: v_mul_f32_e32 v129, 0x3fb8aa3b, v139768 ; GCN-NEXT: v_cvt_f16_f32_e32 v139, v109769 ; GCN-NEXT: v_mfma_f32_32x32x8_f16 v[32:47], v[144:145], v[142:143], v[32:47]770 ; GCN-NEXT: v_fma_f32 v144, s4, v81, -v128771 ; GCN-NEXT: v_exp_f32_e32 v81, v125772 ; GCN-NEXT: v_mul_f32_e32 v125, 0x3fb8aa3b, v144773 ; GCN-NEXT: v_pack_b32_f16 v144, v138, v139774 ; GCN-NEXT: v_mfma_f32_32x32x8_f16 v[16:31], v[136:137], v[142:143], v[16:31]775 ; GCN-NEXT: v_cvt_f16_f32_e32 v136, v110776 ; GCN-NEXT: v_fma_f32 v137, s4, v82, -v128777 ; GCN-NEXT: v_exp_f32_e32 v82, v134778 ; GCN-NEXT: v_cvt_f16_f32_e32 v134, v111779 ; GCN-NEXT: v_mul_f32_e32 v156, 0x3fb8aa3b, v137780 ; GCN-NEXT: v_fma_f32 v137, s4, v83, -v128781 ; GCN-NEXT: v_mul_f32_e32 v157, 0x3fb8aa3b, v137782 ; GCN-NEXT: v_mfma_f32_32x32x8_f16 v[48:63], v[140:141], v[142:143], v[48:63]783 ; GCN-NEXT: v_exp_f32_e32 v83, v135784 ; GCN-NEXT: v_pack_b32_f16 v145, v136, v134785 ; GCN-NEXT: ds_read_b128 v[134:137], v197 offset:1152786 ; GCN-NEXT: s_waitcnt lgkmcnt(0)787 ; GCN-NEXT: buffer_inv sc0 sc1788 ; GCN-NEXT: ds_read_b128 v[138:141], v197 offset:1728789 ; GCN-NEXT: s_waitcnt lgkmcnt(0)790 ; GCN-NEXT: buffer_inv sc0 sc1791 ; GCN-NEXT: ;;#ASMSTART792 ; GCN-NEXT: s_waitcnt vmcnt(8)793 ; GCN-NEXT: ;;#ASMEND794 ; GCN-NEXT: buffer_wbl2 sc0 sc1795 ; GCN-NEXT: ds_write_b64 v199, v[126:127]796 ; GCN-NEXT: buffer_wbl2 sc0 sc1797 ; GCN-NEXT: s_waitcnt lgkmcnt(0)798 ; GCN-NEXT: ds_write_b64 v200, v[150:151]799 ; GCN-NEXT: v_mfma_f32_32x32x8_f16 v[0:15], v[130:131], v[144:145], v[0:15]800 ; GCN-NEXT: buffer_wbl2 sc0 sc1801 ; GCN-NEXT: s_waitcnt lgkmcnt(0)802 ; GCN-NEXT: ds_write_b64 v201, v[152:153]803 ; GCN-NEXT: buffer_wbl2 sc0 sc1804 ; GCN-NEXT: s_waitcnt lgkmcnt(0)805 ; GCN-NEXT: ds_write_b64 v202, v[154:155]806 ; GCN-NEXT: v_fma_f32 v127, s4, v84, -v128807 ; GCN-NEXT: v_exp_f32_e32 v84, v129808 ; GCN-NEXT: v_fma_f32 v130, s4, v85, -v128809 ; GCN-NEXT: v_cvt_f16_f32_e32 v126, v80810 ; GCN-NEXT: v_mul_f32_e32 v129, 0x3fb8aa3b, v127811 ; GCN-NEXT: v_mfma_f32_32x32x8_f16 v[32:47], v[146:147], v[144:145], v[32:47]812 ; GCN-NEXT: v_exp_f32_e32 v85, v125813 ; GCN-NEXT: v_mul_f32_e32 v125, 0x3fb8aa3b, v130814 ; GCN-NEXT: s_waitcnt lgkmcnt(0)815 ; GCN-NEXT: buffer_load_dwordx2 v[130:131], v206, s[0:3], 0 offen sc0 sc1816 ; GCN-NEXT: s_waitcnt vmcnt(0)817 ; GCN-NEXT: buffer_inv sc0 sc1818 ; GCN-NEXT: v_cvt_f16_f32_e32 v127, v81819 ; GCN-NEXT: v_pack_b32_f16 v126, v126, v127820 ; GCN-NEXT: v_mfma_f32_32x32x8_f16 v[16:31], v[134:135], v[144:145], v[16:31]821 ; GCN-NEXT: v_fma_f32 v134, s4, v86, -v128822 ; GCN-NEXT: v_mul_f32_e32 v158, 0x3fb8aa3b, v134823 ; GCN-NEXT: buffer_load_dwordx2 v[134:135], v203, s[0:3], 0 offen sc0 sc1824 ; GCN-NEXT: s_waitcnt vmcnt(0)825 ; GCN-NEXT: buffer_inv sc0 sc1826 ; GCN-NEXT: buffer_load_dwordx2 v[142:143], v204, s[0:3], 0 offen sc0 sc1827 ; GCN-NEXT: s_waitcnt vmcnt(0)828 ; GCN-NEXT: buffer_inv sc0 sc1829 ; GCN-NEXT: buffer_load_dwordx2 v[146:147], v205, s[0:3], 0 offen sc0 sc1830 ; GCN-NEXT: s_waitcnt vmcnt(0)831 ; GCN-NEXT: buffer_inv sc0 sc1832 ; GCN-NEXT: v_cvt_f16_f32_e32 v127, v82833 ; GCN-NEXT: v_exp_f32_e32 v86, v156834 ; GCN-NEXT: v_mfma_f32_32x32x8_f16 v[48:63], v[138:139], v[144:145], v[48:63]835 ; GCN-NEXT: v_cvt_f16_f32_e32 v138, v83836 ; GCN-NEXT: ;;#ASMSTART837 ; GCN-NEXT: s_waitcnt vmcnt(8)838 ; GCN-NEXT: ;;#ASMEND839 ; GCN-NEXT: v_fma_f32 v139, s4, v87, -v128840 ; GCN-NEXT: v_exp_f32_e32 v87, v157841 ; GCN-NEXT: v_pack_b32_f16 v127, v127, v138842 ; GCN-NEXT: v_fma_f32 v138, s4, v89, -v128843 ; GCN-NEXT: v_mul_f32_e32 v139, 0x3fb8aa3b, v139844 ; GCN-NEXT: v_mfma_f32_32x32x8_f16 v[0:15], v[132:133], v[126:127], v[0:15]845 ; GCN-NEXT: ; implicit-def: $sgpr0846 ; GCN-NEXT: v_perm_b32 v154, v135, v131, s5847 ; GCN-NEXT: v_perm_b32 v156, v135, v131, s7848 ; GCN-NEXT: v_fma_f32 v135, s4, v88, -v128849 ; GCN-NEXT: v_perm_b32 v150, v134, v130, s5850 ; GCN-NEXT: v_perm_b32 v152, v134, v130, s7851 ; GCN-NEXT: ds_read_b128 v[130:133], v198852 ; GCN-NEXT: s_waitcnt lgkmcnt(0)853 ; GCN-NEXT: buffer_inv sc0 sc1854 ; GCN-NEXT: v_cvt_f16_f32_e32 v134, v84855 ; GCN-NEXT: v_exp_f32_e32 v88, v129856 ; GCN-NEXT: v_mul_f32_e32 v129, 0x3fb8aa3b, v135857 ; GCN-NEXT: v_cvt_f16_f32_e32 v135, v85858 ; GCN-NEXT: v_perm_b32 v151, v146, v142, s5859 ; GCN-NEXT: v_perm_b32 v153, v146, v142, s7860 ; GCN-NEXT: v_perm_b32 v155, v147, v143, s5861 ; GCN-NEXT: v_perm_b32 v157, v147, v143, s7862 ; GCN-NEXT: ds_read_b128 v[142:145], v198 offset:576863 ; GCN-NEXT: s_waitcnt lgkmcnt(0)864 ; GCN-NEXT: buffer_inv sc0 sc1865 ; GCN-NEXT: v_mfma_f32_32x32x8_f16 v[32:47], v[148:149], v[126:127], v[32:47]866 ; GCN-NEXT: v_exp_f32_e32 v89, v125867 ; GCN-NEXT: v_pack_b32_f16 v146, v134, v135868 ; GCN-NEXT: v_cvt_f16_f32_e32 v134, v86869 ; GCN-NEXT: v_fma_f32 v135, s4, v90, -v128870 ; GCN-NEXT: v_mul_f32_e32 v125, 0x3fb8aa3b, v138871 ; GCN-NEXT: v_mul_f32_e32 v148, 0x3fb8aa3b, v135872 ; GCN-NEXT: v_mfma_f32_32x32x8_f16 v[16:31], v[136:137], v[126:127], v[16:31]873 ; GCN-NEXT: v_exp_f32_e32 v90, v158874 ; GCN-NEXT: v_mul_f32_e32 v158, 0x3fb8aa3b, v64875 ; GCN-NEXT: v_mfma_f32_32x32x8_f16 v[48:63], v[140:141], v[126:127], v[48:63]876 ; GCN-NEXT: v_cvt_f16_f32_e32 v126, v87877 ; GCN-NEXT: v_fma_f32 v127, s4, v91, -v128878 ; GCN-NEXT: v_exp_f32_e32 v91, v139879 ; GCN-NEXT: v_mul_f32_e32 v127, 0x3fb8aa3b, v127880 ; GCN-NEXT: v_pack_b32_f16 v147, v134, v126881 ; GCN-NEXT: ds_read_b128 v[134:137], v198 offset:1152882 ; GCN-NEXT: s_waitcnt lgkmcnt(0)883 ; GCN-NEXT: buffer_inv sc0 sc1884 ; GCN-NEXT: ds_read_b128 v[138:141], v198 offset:1728885 ; GCN-NEXT: s_waitcnt lgkmcnt(0)886 ; GCN-NEXT: buffer_inv sc0 sc1887 ; GCN-NEXT: v_mfma_f32_32x32x8_f16 v[0:15], v[130:131], v[146:147], v[0:15]888 ; GCN-NEXT: v_fma_f32 v130, s4, v92, -v128889 ; GCN-NEXT: v_cvt_f16_f32_e32 v126, v88890 ; GCN-NEXT: v_exp_f32_e32 v92, v129891 ; GCN-NEXT: v_mul_f32_e32 v129, 0x3fb8aa3b, v130892 ; GCN-NEXT: v_cvt_f16_f32_e32 v130, v89893 ; GCN-NEXT: v_fma_f32 v131, s4, v93, -v128894 ; GCN-NEXT: v_pack_b32_f16 v130, v126, v130895 ; GCN-NEXT: v_mfma_f32_32x32x8_f16 v[32:47], v[142:143], v[146:147], v[32:47]896 ; GCN-NEXT: v_exp_f32_e32 v93, v125897 ; GCN-NEXT: v_fma_f32 v126, s4, v94, -v128898 ; GCN-NEXT: v_cvt_f16_f32_e32 v125, v90899 ; GCN-NEXT: v_mul_f32_e32 v143, 0x3fb8aa3b, v126900 ; GCN-NEXT: v_cvt_f16_f32_e32 v126, v91901 ; GCN-NEXT: v_mul_f32_e32 v142, 0x3fb8aa3b, v131902 ; GCN-NEXT: v_fma_f32 v131, s4, v95, -v128903 ; GCN-NEXT: v_mfma_f32_32x32x8_f16 v[16:31], v[134:135], v[146:147], v[16:31]904 ; GCN-NEXT: v_exp_f32_e32 v94, v148905 ; GCN-NEXT: v_cvt_f16_f32_e32 v64, v93906 ; GCN-NEXT: v_mfma_f32_32x32x8_f16 v[48:63], v[138:139], v[146:147], v[48:63]907 ; GCN-NEXT: v_exp_f32_e32 v95, v127908 ; GCN-NEXT: v_cvt_f16_f32_e32 v127, v92909 ; GCN-NEXT: v_mul_f32_e32 v138, 0x3fb8aa3b, v131910 ; GCN-NEXT: v_pack_b32_f16 v131, v125, v126911 ; GCN-NEXT: s_nop 1912 ; GCN-NEXT: v_mfma_f32_32x32x8_f16 v[0:15], v[132:133], v[130:131], v[0:15]913 ; GCN-NEXT: v_exp_f32_e32 v125, v129914 ; GCN-NEXT: ds_read_b128 v[132:135], v197915 ; GCN-NEXT: s_waitcnt lgkmcnt(0)916 ; GCN-NEXT: buffer_inv sc0 sc1917 ; GCN-NEXT: ds_read_b128 v[146:149], v197 offset:576918 ; GCN-NEXT: s_waitcnt lgkmcnt(0)919 ; GCN-NEXT: buffer_inv sc0 sc1920 ; GCN-NEXT: v_mfma_f32_32x32x8_f16 v[32:47], v[144:145], v[130:131], v[32:47]921 ; GCN-NEXT: v_mul_f32_e32 v144, 0x3fb8aa3b, v65922 ; GCN-NEXT: v_fma_f32 v65, s4, v66, -v128923 ; GCN-NEXT: v_exp_f32_e32 v126, v142924 ; GCN-NEXT: v_pack_b32_f16 v142, v127, v64925 ; GCN-NEXT: v_cvt_f16_f32_e32 v64, v94926 ; GCN-NEXT: v_mul_f32_e32 v145, 0x3fb8aa3b, v65927 ; GCN-NEXT: v_cvt_f16_f32_e32 v65, v95928 ; GCN-NEXT: v_fma_f32 v66, s4, v67, -v128929 ; GCN-NEXT: v_mfma_f32_32x32x8_f16 v[16:31], v[136:137], v[130:131], v[16:31]930 ; GCN-NEXT: v_exp_f32_e32 v127, v143931 ; GCN-NEXT: v_pack_b32_f16 v143, v64, v65932 ; GCN-NEXT: v_mfma_f32_32x32x8_f16 v[48:63], v[140:141], v[130:131], v[48:63]933 ; GCN-NEXT: v_exp_f32_e32 v129, v138934 ; GCN-NEXT: v_mul_f32_e32 v141, 0x3fb8aa3b, v66935 ; GCN-NEXT: ds_read_b128 v[64:67], v197 offset:1152936 ; GCN-NEXT: s_waitcnt lgkmcnt(0)937 ; GCN-NEXT: buffer_inv sc0 sc1938 ; GCN-NEXT: ds_read_b128 v[136:139], v197 offset:1728939 ; GCN-NEXT: s_waitcnt lgkmcnt(0)940 ; GCN-NEXT: buffer_inv sc0 sc1941 ; GCN-NEXT: ;;#ASMSTART942 ; GCN-NEXT: s_waitcnt vmcnt(8)943 ; GCN-NEXT: ;;#ASMEND944 ; GCN-NEXT: buffer_wbl2 sc0 sc1945 ; GCN-NEXT: ds_write_b64 v199, v[150:151]946 ; GCN-NEXT: buffer_wbl2 sc0 sc1947 ; GCN-NEXT: s_waitcnt lgkmcnt(0)948 ; GCN-NEXT: ds_write_b64 v200, v[152:153]949 ; GCN-NEXT: v_mfma_f32_32x32x8_f16 v[0:15], v[132:133], v[142:143], v[0:15]950 ; GCN-NEXT: v_cvt_f16_f32_e32 v132, v125951 ; GCN-NEXT: v_exp_f32_e32 v130, v158952 ; GCN-NEXT: buffer_wbl2 sc0 sc1953 ; GCN-NEXT: s_waitcnt lgkmcnt(0)954 ; GCN-NEXT: ds_write_b64 v201, v[154:155]955 ; GCN-NEXT: buffer_wbl2 sc0 sc1956 ; GCN-NEXT: s_waitcnt lgkmcnt(0)957 ; GCN-NEXT: ds_write_b64 v202, v[156:157]958 ; GCN-NEXT: ;;#ASMSTART959 ; GCN-NEXT: s_waitcnt vmcnt(8)960 ; GCN-NEXT: ;;#ASMEND961 ; GCN-NEXT: v_mfma_f32_32x32x8_f16 v[32:47], v[146:147], v[142:143], v[32:47]962 ; GCN-NEXT: v_mul_f32_e32 v146, 0x3fb8aa3b, v68963 ; GCN-NEXT: v_cvt_f16_f32_e32 v68, v126964 ; GCN-NEXT: v_exp_f32_e32 v131, v144965 ; GCN-NEXT: v_mul_f32_e32 v144, 0x3fb8aa3b, v69966 ; GCN-NEXT: v_fma_f32 v69, s4, v71, -v128967 ; GCN-NEXT: v_pack_b32_f16 v140, v132, v68968 ; GCN-NEXT: v_cvt_f16_f32_e32 v68, v129969 ; GCN-NEXT: v_mfma_f32_32x32x8_f16 v[16:31], v[64:65], v[142:143], v[16:31]970 ; GCN-NEXT: v_cvt_f16_f32_e32 v64, v127971 ; GCN-NEXT: v_exp_f32_e32 v132, v145972 ; GCN-NEXT: v_fma_f32 v65, s4, v70, -v128973 ; GCN-NEXT: v_mul_f32_e32 v65, 0x3fb8aa3b, v65974 ; GCN-NEXT: v_fma_f32 v145, s4, v73, -v128975 ; GCN-NEXT: v_mul_f32_e32 v147, 0x3fb8aa3b, v145976 ; GCN-NEXT: v_mfma_f32_32x32x8_f16 v[48:63], v[136:137], v[142:143], v[48:63]977 ; GCN-NEXT: v_exp_f32_e32 v133, v141978 ; GCN-NEXT: v_mul_f32_e32 v142, 0x3fb8aa3b, v69979 ; GCN-NEXT: v_pack_b32_f16 v141, v64, v68980 ; GCN-NEXT: s_waitcnt lgkmcnt(0)981 ; GCN-NEXT: ds_read_b128 v[68:71], v198982 ; GCN-NEXT: s_waitcnt lgkmcnt(0)983 ; GCN-NEXT: buffer_inv sc0 sc1984 ; GCN-NEXT: v_fma_f32 v143, s4, v72, -v128985 ; GCN-NEXT: v_cvt_f16_f32_e32 v64, v130986 ; GCN-NEXT: v_mfma_f32_32x32x8_f16 v[0:15], v[134:135], v[140:141], v[0:15]987 ; GCN-NEXT: v_exp_f32_e32 v72, v146988 ; GCN-NEXT: v_mul_f32_e32 v146, 0x3fb8aa3b, v143989 ; GCN-NEXT: v_cvt_f16_f32_e32 v143, v131990 ; GCN-NEXT: ds_read_b128 v[134:137], v198 offset:576991 ; GCN-NEXT: s_waitcnt lgkmcnt(0)992 ; GCN-NEXT: buffer_inv sc0 sc1993 ; GCN-NEXT: v_pack_b32_f16 v64, v64, v143994 ; GCN-NEXT: v_mfma_f32_32x32x8_f16 v[32:47], v[148:149], v[140:141], v[32:47]995 ; GCN-NEXT: v_exp_f32_e32 v73, v144996 ; GCN-NEXT: v_mfma_f32_32x32x8_f16 v[16:31], v[66:67], v[140:141], v[16:31]997 ; GCN-NEXT: v_cvt_f16_f32_e32 v66, v132998 ; GCN-NEXT: v_fma_f32 v67, s4, v74, -v128999 ; GCN-NEXT: v_exp_f32_e32 v74, v651000 ; GCN-NEXT: v_cvt_f16_f32_e32 v65, v1331001 ; GCN-NEXT: v_mul_f32_e32 v67, 0x3fb8aa3b, v671002 ; GCN-NEXT: v_pack_b32_f16 v65, v66, v651003 ; GCN-NEXT: v_mfma_f32_32x32x8_f16 v[48:63], v[138:139], v[140:141], v[48:63]1004 ; GCN-NEXT: v_fma_f32 v138, s4, v75, -v1281005 ; GCN-NEXT: v_exp_f32_e32 v75, v1421006 ; GCN-NEXT: v_mul_f32_e32 v148, 0x3fb8aa3b, v1381007 ; GCN-NEXT: ds_read_b128 v[138:141], v198 offset:11521008 ; GCN-NEXT: s_waitcnt lgkmcnt(0)1009 ; GCN-NEXT: buffer_inv sc0 sc11010 ; GCN-NEXT: ds_read_b128 v[142:145], v198 offset:17281011 ; GCN-NEXT: s_waitcnt lgkmcnt(0)1012 ; GCN-NEXT: buffer_inv sc0 sc11013 ; GCN-NEXT: v_cvt_f16_f32_e32 v66, v721014 ; GCN-NEXT: v_mfma_f32_32x32x8_f16 v[0:15], v[68:69], v[64:65], v[0:15]1015 ; GCN-NEXT: v_fma_f32 v68, s4, v76, -v1281016 ; GCN-NEXT: v_exp_f32_e32 v76, v1461017 ; GCN-NEXT: v_mul_f32_e32 v146, 0x3fb8aa3b, v681018 ; GCN-NEXT: v_cvt_f16_f32_e32 v68, v731019 ; GCN-NEXT: v_fma_f32 v69, s4, v77, -v1281020 ; GCN-NEXT: v_mfma_f32_32x32x8_f16 v[32:47], v[134:135], v[64:65], v[32:47]1021 ; GCN-NEXT: v_exp_f32_e32 v77, v1471022 ; GCN-NEXT: v_pack_b32_f16 v134, v66, v681023 ; GCN-NEXT: v_fma_f32 v68, s4, v78, -v1281024 ; GCN-NEXT: v_cvt_f16_f32_e32 v66, v741025 ; GCN-NEXT: v_mul_f32_e32 v147, 0x3fb8aa3b, v691026 ; GCN-NEXT: v_mfma_f32_32x32x8_f16 v[16:31], v[138:139], v[64:65], v[16:31]1027 ; GCN-NEXT: v_exp_f32_e32 v78, v671028 ; GCN-NEXT: v_mul_f32_e32 v138, 0x3fb8aa3b, v681029 ; GCN-NEXT: v_cvt_f16_f32_e32 v139, v761030 ; GCN-NEXT: v_mfma_f32_32x32x8_f16 v[48:63], v[142:143], v[64:65], v[48:63]1031 ; GCN-NEXT: v_cvt_f16_f32_e32 v64, v751032 ; GCN-NEXT: v_fma_f32 v65, s4, v79, -v1281033 ; GCN-NEXT: v_exp_f32_e32 v79, v1481034 ; GCN-NEXT: v_mul_f32_e32 v128, 0x3fb8aa3b, v651035 ; GCN-NEXT: v_pack_b32_f16 v135, v66, v641036 ; GCN-NEXT: s_nop 11037 ; GCN-NEXT: v_mfma_f32_32x32x8_f16 v[0:15], v[70:71], v[134:135], v[0:15]1038 ; GCN-NEXT: v_exp_f32_e32 v142, v1461039 ; GCN-NEXT: ds_read_b128 v[68:71], v1971040 ; GCN-NEXT: s_waitcnt lgkmcnt(0)1041 ; GCN-NEXT: buffer_inv sc0 sc11042 ; GCN-NEXT: ds_read_b128 v[64:67], v197 offset:5761043 ; GCN-NEXT: s_waitcnt lgkmcnt(0)1044 ; GCN-NEXT: buffer_inv sc0 sc11045 ; GCN-NEXT: v_mfma_f32_32x32x8_f16 v[32:47], v[136:137], v[134:135], v[32:47]1046 ; GCN-NEXT: v_exp_f32_e32 v137, v1471047 ; GCN-NEXT: v_cvt_f16_f32_e32 v136, v771048 ; GCN-NEXT: v_mfma_f32_32x32x8_f16 v[16:31], v[140:141], v[134:135], v[16:31]1049 ; GCN-NEXT: v_exp_f32_e32 v138, v1381050 ; GCN-NEXT: v_cvt_f16_f32_e32 v140, v781051 ; GCN-NEXT: v_mfma_f32_32x32x8_f16 v[48:63], v[144:145], v[134:135], v[48:63]1052 ; GCN-NEXT: s_nop 101053 ; GCN-NEXT: v_exp_f32_e32 v52, v1281054 ; GCN-NEXT: v_cvt_f16_f32_e32 v50, v1371055 ; GCN-NEXT: v_cvt_f16_f32_e32 v51, v1421056 ; GCN-NEXT: v_cvt_f16_f32_e32 v54, v1381057 ; GCN-NEXT: v_cvt_f16_f32_e32 v53, v521058 ; GCN-NEXT: v_cvt_f16_f32_e32 v49, v791059 ; GCN-NEXT: v_pack_b32_f16 v50, v51, v501060 ; GCN-NEXT: v_pack_b32_f16 v48, v139, v1361061 ; GCN-NEXT: v_pack_b32_f16 v51, v54, v531062 ; GCN-NEXT: v_add_f32_e32 v53, 0, v1131063 ; GCN-NEXT: v_add_f32_e32 v53, v114, v531064 ; GCN-NEXT: v_add_f32_e32 v53, v115, v531065 ; GCN-NEXT: v_add_f32_e32 v53, v116, v531066 ; GCN-NEXT: v_add_f32_e32 v53, v117, v531067 ; GCN-NEXT: v_add_f32_e32 v53, v118, v531068 ; GCN-NEXT: v_add_f32_e32 v53, v119, v531069 ; GCN-NEXT: v_add_f32_e32 v53, v120, v531070 ; GCN-NEXT: v_add_f32_e32 v53, v121, v531071 ; GCN-NEXT: v_add_f32_e32 v53, v122, v531072 ; GCN-NEXT: v_add_f32_e32 v53, v123, v531073 ; GCN-NEXT: v_add_f32_e32 v53, v124, v531074 ; GCN-NEXT: v_add_f32_e32 v53, v96, v531075 ; GCN-NEXT: v_add_f32_e32 v53, v97, v531076 ; GCN-NEXT: v_add_f32_e32 v53, v98, v531077 ; GCN-NEXT: v_add_f32_e32 v53, v99, v531078 ; GCN-NEXT: v_add_f32_e32 v53, v100, v531079 ; GCN-NEXT: v_add_f32_e32 v53, v101, v531080 ; GCN-NEXT: v_add_f32_e32 v53, v102, v531081 ; GCN-NEXT: v_add_f32_e32 v53, v103, v531082 ; GCN-NEXT: v_add_f32_e32 v53, v104, v531083 ; GCN-NEXT: v_add_f32_e32 v53, v105, v531084 ; GCN-NEXT: v_add_f32_e32 v53, v106, v531085 ; GCN-NEXT: v_add_f32_e32 v53, v107, v531086 ; GCN-NEXT: v_add_f32_e32 v53, v108, v531087 ; GCN-NEXT: v_add_f32_e32 v53, v109, v531088 ; GCN-NEXT: v_add_f32_e32 v53, v110, v531089 ; GCN-NEXT: v_add_f32_e32 v53, v111, v531090 ; GCN-NEXT: v_add_f32_e32 v53, v80, v531091 ; GCN-NEXT: v_add_f32_e32 v53, v81, v531092 ; GCN-NEXT: v_add_f32_e32 v53, v82, v531093 ; GCN-NEXT: v_add_f32_e32 v53, v83, v531094 ; GCN-NEXT: v_add_f32_e32 v53, v84, v531095 ; GCN-NEXT: v_add_f32_e32 v53, v85, v531096 ; GCN-NEXT: v_add_f32_e32 v53, v86, v531097 ; GCN-NEXT: v_add_f32_e32 v53, v87, v531098 ; GCN-NEXT: v_add_f32_e32 v53, v88, v531099 ; GCN-NEXT: v_add_f32_e32 v53, v89, v531100 ; GCN-NEXT: v_pack_b32_f16 v49, v140, v491101 ; GCN-NEXT: v_add_f32_e32 v53, v90, v531102 ; GCN-NEXT: v_add_f32_e32 v53, v91, v531103 ; GCN-NEXT: v_mfma_f32_32x32x8_f16 v[0:15], v[68:69], v[48:49], v[0:15]1104 ; GCN-NEXT: v_add_f32_e32 v53, v92, v531105 ; GCN-NEXT: v_add_f32_e32 v53, v93, v531106 ; GCN-NEXT: v_add_f32_e32 v53, v94, v531107 ; GCN-NEXT: v_add_f32_e32 v53, v95, v531108 ; GCN-NEXT: v_add_f32_e32 v53, v125, v531109 ; GCN-NEXT: v_add_f32_e32 v53, v126, v531110 ; GCN-NEXT: v_add_f32_e32 v53, v127, v531111 ; GCN-NEXT: v_add_f32_e32 v53, v129, v531112 ; GCN-NEXT: v_mfma_f32_32x32x8_f16 v[0:15], v[70:71], v[50:51], v[0:15]1113 ; GCN-NEXT: v_mfma_f32_32x32x8_f16 v[32:47], v[64:65], v[48:49], v[32:47]1114 ; GCN-NEXT: s_nop 91115 ; GCN-NEXT: v_add_f32_e32 v0, v130, v531116 ; GCN-NEXT: v_add_f32_e32 v0, v131, v01117 ; GCN-NEXT: v_add_f32_e32 v0, v132, v01118 ; GCN-NEXT: v_add_f32_e32 v0, v133, v01119 ; GCN-NEXT: v_add_f32_e32 v0, v72, v01120 ; GCN-NEXT: v_add_f32_e32 v0, v73, v01121 ; GCN-NEXT: v_add_f32_e32 v0, v74, v01122 ; GCN-NEXT: v_add_f32_e32 v0, v75, v01123 ; GCN-NEXT: v_add_f32_e32 v0, v76, v01124 ; GCN-NEXT: v_add_f32_e32 v0, v77, v01125 ; GCN-NEXT: v_add_f32_e32 v0, v78, v01126 ; GCN-NEXT: v_add_f32_e32 v0, v79, v01127 ; GCN-NEXT: v_add_f32_e32 v0, v142, v01128 ; GCN-NEXT: v_add_f32_e32 v0, v137, v01129 ; GCN-NEXT: v_add_f32_e32 v0, v138, v01130 ; GCN-NEXT: v_add_f32_e32 v4, v52, v01131 ; GCN-NEXT: ds_bpermute_b32 v5, v196, v41132 ; GCN-NEXT: s_waitcnt lgkmcnt(0)1133 ; GCN-NEXT: ds_read_b128 v[0:3], v197 offset:11521134 ; GCN-NEXT: s_waitcnt lgkmcnt(0)1135 ; GCN-NEXT: buffer_inv sc0 sc11136 ; GCN-NEXT: v_mfma_f32_32x32x8_f16 v[16:31], v[0:1], v[48:49], v[16:31]1137 ; GCN-NEXT: v_add_f32_e32 v2, v4, v51138 ; GCN-NEXT: ds_bpermute_b32 v3, v196, v21139 ; GCN-NEXT: ; implicit-def: $vgpr41140 ; GCN-NEXT: s_waitcnt lgkmcnt(0)1141 ; GCN-NEXT: v_cndmask_b32_e64 v0, v3, v2, s[12:13]1142 ; GCN-NEXT: v_fmac_f32_e32 v0, v4, v1121143 ; GCN-NEXT: ds_read_b128 v[0:3], v197 offset:17281144 ; GCN-NEXT: s_waitcnt lgkmcnt(0)1145 ; GCN-NEXT: buffer_inv sc0 sc11146 ; GCN-NEXT: ;;#ASMSTART1147 ; GCN-NEXT: s_waitcnt vmcnt(8)1148 ; GCN-NEXT: ;;#ASMEND1149 ; GCN-NEXT: v_mfma_f32_32x32x8_f16 v[32:47], v[66:67], v[50:51], v[32:47]1150 ; GCN-NEXT: s_endpgm1151 1152 attributes #0 = {"amdgpu-flat-work-group-size"="256,256"}1153 !0 = !{i64 2862105}1154 1155...1156 1157---1158name: largeInterleave1159tracksRegLiveness: true1160machineFunctionInfo:1161 stackPtrOffsetReg: '$sgpr32'1162 occupancy: 71163body: |1164 bb.0:1165 liveins: $vgpr0, $sgpr0_sgpr1, $sgpr2, $sgpr3, $sgpr41166 %11:vgpr_32 = IMPLICIT_DEF1167 %1:sgpr_512 = IMPLICIT_DEF1168 %16:vgpr_32 = IMPLICIT_DEF1169 %443:sgpr_128 = IMPLICIT_DEF1170 %18:sreg_32 = IMPLICIT_DEF1171 %25:vgpr_32 = IMPLICIT_DEF1172 %23:vgpr_32 = IMPLICIT_DEF1173 %391:vreg_128_align2 = IMPLICIT_DEF1174 %24:vgpr_32 = IMPLICIT_DEF1175 %392:vreg_128_align2 = IMPLICIT_DEF1176 %401:vreg_128_align2 = IMPLICIT_DEF1177 %406:vreg_128_align2 = IMPLICIT_DEF1178 %48:vgpr_32 = IMPLICIT_DEF1179 %473:sgpr_128 = IMPLICIT_DEF1180 %411:vreg_128_align2 = IMPLICIT_DEF1181 %416:vreg_128_align2 = IMPLICIT_DEF1182 %421:vreg_128_align2 = IMPLICIT_DEF1183 %426:vreg_128_align2 = IMPLICIT_DEF1184 %1114:sgpr_32 = IMPLICIT_DEF1185 %39:vgpr_32 = IMPLICIT_DEF1186 %484:sreg_64_xexec = IMPLICIT_DEF1187 %3346:vgpr_32 = IMPLICIT_DEF1188 %1422:sreg_32 = IMPLICIT_DEF1189 %1424:sreg_32 = IMPLICIT_DEF1190 %15:vgpr_32 = IMPLICIT_DEF1191 %494:sreg_32 = IMPLICIT_DEF1192 %47:vgpr_32 = IMPLICIT_DEF1193 %41:vgpr_32 = IMPLICIT_DEF1194 %42:vgpr_32 = IMPLICIT_DEF1195 %43:vgpr_32 = IMPLICIT_DEF1196 %44:vgpr_32 = IMPLICIT_DEF1197 %45:vgpr_32 = IMPLICIT_DEF1198 %50:sreg_32 = IMPLICIT_DEF1199 %3347:vgpr_32 = IMPLICIT_DEF1200 %3329:vgpr_32 = IMPLICIT_DEF1201 %3330:vgpr_32 = IMPLICIT_DEF1202 %3331:vgpr_32 = IMPLICIT_DEF1203 %3332:vgpr_32 = IMPLICIT_DEF1204 %3333:vgpr_32 = IMPLICIT_DEF1205 %2986:vreg_512_align2 = IMPLICIT_DEF1206 %3038:vreg_512_align2 = IMPLICIT_DEF1207 %2980:vreg_512_align2 = IMPLICIT_DEF1208 %3003:vreg_512_align2 = IMPLICIT_DEF1209 %3334:vgpr_32 = IMPLICIT_DEF1210 %3335:vgpr_32 = IMPLICIT_DEF1211 %3336:vgpr_32 = IMPLICIT_DEF1212 %3337:vgpr_32 = IMPLICIT_DEF1213 %3338:vgpr_32 = IMPLICIT_DEF1214 %3339:vgpr_32 = IMPLICIT_DEF1215 %3345:vgpr_32 = IMPLICIT_DEF1216 %3340:vgpr_32 = IMPLICIT_DEF1217 %3341:vgpr_32 = IMPLICIT_DEF1218 %3342:vgpr_32 = IMPLICIT_DEF1219 %3343:vgpr_32 = IMPLICIT_DEF1220 %3344:vgpr_32 = IMPLICIT_DEF1221 %84:vgpr_32 = COPY %33471222 %86:vgpr_32 = COPY %3347:vgpr_321223 IGLP_OPT 21224 %593:sreg_32_xm0 = V_READFIRSTLANE_B32 %11:vgpr_32, implicit $exec1225 %595:vgpr_32 = V_LSHL_ADD_U32_e64 %593:sreg_32_xm0, 4, %3329:vgpr_32, implicit $exec1226 %597:vgpr_32 = nsw V_MUL_LO_U32_e64 %595:vgpr_32, %1.sub6:sgpr_512, implicit $exec1227 %599:vgpr_32 = V_ADD_LSHL_U32_e64 %597:vgpr_32, %16:vgpr_32, 1, implicit $exec1228 %601:vreg_128_align2 = BUFFER_LOAD_DWORDX4_OFFEN %599:vgpr_32, %443:sgpr_128, 0, 0, 0, 0, implicit $exec1229 %602:vgpr_32 = V_ADD_U32_e32 %18:sreg_32, %599:vgpr_32, implicit $exec1230 %603:vreg_128_align2 = BUFFER_LOAD_DWORDX4_OFFEN %602:vgpr_32, %443:sgpr_128, 0, 0, 0, 0, implicit $exec1231 %605:sreg_32 = S_LSHL_B32 %593:sreg_32_xm0, 7, implicit-def dead $scc1232 %606:vgpr_32 = V_ADD_LSHL_U32_e64 %25:vgpr_32, %605:sreg_32, 1, implicit $exec1233 DS_WRITE_B128_gfx9 %606:vgpr_32, %601:vreg_128_align2, 0, 0, implicit $exec1234 DS_WRITE_B128_gfx9 %606:vgpr_32, %603:vreg_128_align2, 1024, 0, implicit $exec1235 %608:vreg_128_align2 = BUFFER_LOAD_DWORDX4_OFFEN %599:vgpr_32, %443:sgpr_128, 0, 64, 0, 0, implicit $exec1236 %610:vgpr_32 = V_ADD_U32_e32 64, %602:vgpr_32, implicit $exec1237 %611:vreg_128_align2 = BUFFER_LOAD_DWORDX4_OFFEN %610:vgpr_32, %443:sgpr_128, 0, 0, 0, 0, implicit $exec1238 INLINEASM &"s_waitcnt vmcnt($0)", 57 /* sideeffect mayload maystore isconvergent attdialect */, 13 /* imm */, 8, !01239 %612:vreg_128_align2 = DS_READ_B128_gfx9 %23:vgpr_32, 0, 0, implicit $exec1240 early-clobber %668:vreg_512_align2 = contract V_MFMA_F32_32X32X8F16_vgprcd_e64 %612.sub0_sub1:vreg_128_align2, %391.sub0_sub1:vreg_128_align2, 0, 0, 0, 0, implicit $mode, implicit $exec1241 %668:vreg_512_align2 = contract V_MFMA_F32_32X32X8F16_mac_vgprcd_e64 %612.sub2_sub3:vreg_128_align2, %391.sub2_sub3:vreg_128_align2, %668:vreg_512_align2, 0, 0, 0, implicit $mode, implicit $exec1242 %626:vreg_128_align2 = DS_READ_B128_gfx9 %23:vgpr_32, 512, 0, implicit $exec1243 early-clobber %679:vreg_512_align2 = contract V_MFMA_F32_32X32X8F16_vgprcd_e64 %626.sub0_sub1:vreg_128_align2, %391.sub0_sub1:vreg_128_align2, 0, 0, 0, 0, implicit $mode, implicit $exec1244 %679:vreg_512_align2 = contract V_MFMA_F32_32X32X8F16_mac_vgprcd_e64 %626.sub2_sub3:vreg_128_align2, %391.sub2_sub3:vreg_128_align2, %679:vreg_512_align2, 0, 0, 0, implicit $mode, implicit $exec1245 %638:vreg_128_align2 = DS_READ_B128_gfx9 %23:vgpr_32, 1024, 0, implicit $exec1246 early-clobber %690:vreg_512_align2 = contract V_MFMA_F32_32X32X8F16_vgprcd_e64 %638.sub0_sub1:vreg_128_align2, %391.sub0_sub1:vreg_128_align2, 0, 0, 0, 0, implicit $mode, implicit $exec1247 %690:vreg_512_align2 = contract V_MFMA_F32_32X32X8F16_mac_vgprcd_e64 %638.sub2_sub3:vreg_128_align2, %391.sub2_sub3:vreg_128_align2, %690:vreg_512_align2, 0, 0, 0, implicit $mode, implicit $exec1248 %650:vreg_128_align2 = DS_READ_B128_gfx9 %23:vgpr_32, 1536, 0, implicit $exec1249 early-clobber %701:vreg_512_align2 = contract V_MFMA_F32_32X32X8F16_vgprcd_e64 %650.sub0_sub1:vreg_128_align2, %391.sub0_sub1:vreg_128_align2, 0, 0, 0, 0, implicit $mode, implicit $exec1250 %701:vreg_512_align2 = contract V_MFMA_F32_32X32X8F16_mac_vgprcd_e64 %650.sub2_sub3:vreg_128_align2, %391.sub2_sub3:vreg_128_align2, %701:vreg_512_align2, 0, 0, 0, implicit $mode, implicit $exec1251 %662:vreg_128_align2 = DS_READ_B128_gfx9 %24:vgpr_32, 0, 0, implicit $exec1252 %668:vreg_512_align2 = contract V_MFMA_F32_32X32X8F16_mac_vgprcd_e64 %662.sub0_sub1:vreg_128_align2, %392.sub0_sub1:vreg_128_align2, %668:vreg_512_align2, 0, 0, 0, implicit $mode, implicit $exec1253 %668:vreg_512_align2 = contract V_MFMA_F32_32X32X8F16_mac_vgprcd_e64 %662.sub2_sub3:vreg_128_align2, %392.sub2_sub3:vreg_128_align2, %668:vreg_512_align2, 0, 0, 0, implicit $mode, implicit $exec1254 %673:vreg_128_align2 = DS_READ_B128_gfx9 %24:vgpr_32, 512, 0, implicit $exec1255 %679:vreg_512_align2 = contract V_MFMA_F32_32X32X8F16_mac_vgprcd_e64 %673.sub0_sub1:vreg_128_align2, %392.sub0_sub1:vreg_128_align2, %679:vreg_512_align2, 0, 0, 0, implicit $mode, implicit $exec1256 %679:vreg_512_align2 = contract V_MFMA_F32_32X32X8F16_mac_vgprcd_e64 %673.sub2_sub3:vreg_128_align2, %392.sub2_sub3:vreg_128_align2, %679:vreg_512_align2, 0, 0, 0, implicit $mode, implicit $exec1257 %684:vreg_128_align2 = DS_READ_B128_gfx9 %24:vgpr_32, 1024, 0, implicit $exec1258 %690:vreg_512_align2 = contract V_MFMA_F32_32X32X8F16_mac_vgprcd_e64 %684.sub0_sub1:vreg_128_align2, %392.sub0_sub1:vreg_128_align2, %690:vreg_512_align2, 0, 0, 0, implicit $mode, implicit $exec1259 %690:vreg_512_align2 = contract V_MFMA_F32_32X32X8F16_mac_vgprcd_e64 %684.sub2_sub3:vreg_128_align2, %392.sub2_sub3:vreg_128_align2, %690:vreg_512_align2, 0, 0, 0, implicit $mode, implicit $exec1260 %695:vreg_128_align2 = DS_READ_B128_gfx9 %24:vgpr_32, 1536, 0, implicit $exec1261 %701:vreg_512_align2 = contract V_MFMA_F32_32X32X8F16_mac_vgprcd_e64 %695.sub0_sub1:vreg_128_align2, %392.sub0_sub1:vreg_128_align2, %701:vreg_512_align2, 0, 0, 0, implicit $mode, implicit $exec1262 %701:vreg_512_align2 = contract V_MFMA_F32_32X32X8F16_mac_vgprcd_e64 %695.sub2_sub3:vreg_128_align2, %392.sub2_sub3:vreg_128_align2, %701:vreg_512_align2, 0, 0, 0, implicit $mode, implicit $exec1263 INLINEASM &"s_waitcnt vmcnt($0)", 57 /* sideeffect mayload maystore isconvergent attdialect */, 13 /* imm */, 8, !01264 DS_WRITE_B128_gfx9 %606:vgpr_32, %608:vreg_128_align2, 0, 0, implicit $exec1265 DS_WRITE_B128_gfx9 %606:vgpr_32, %611:vreg_128_align2, 1024, 0, implicit $exec1266 %706:vreg_128_align2 = BUFFER_LOAD_DWORDX4_OFFEN %599:vgpr_32, %443:sgpr_128, 0, 128, 0, 0, implicit $exec1267 %708:vgpr_32 = V_ADD_U32_e32 128, %602:vgpr_32, implicit $exec1268 %709:vreg_128_align2 = BUFFER_LOAD_DWORDX4_OFFEN %708:vgpr_32, %443:sgpr_128, 0, 0, 0, 0, implicit $exec1269 INLINEASM &"s_waitcnt vmcnt($0)", 57 /* sideeffect mayload maystore isconvergent attdialect */, 13 /* imm */, 8, !01270 %710:vreg_128_align2 = DS_READ_B128_gfx9 %23:vgpr_32, 0, 0, implicit $exec1271 %668:vreg_512_align2 = contract V_MFMA_F32_32X32X8F16_mac_vgprcd_e64 %710.sub0_sub1:vreg_128_align2, %401.sub0_sub1:vreg_128_align2, %668:vreg_512_align2, 0, 0, 0, implicit $mode, implicit $exec1272 %668:vreg_512_align2 = contract V_MFMA_F32_32X32X8F16_mac_vgprcd_e64 %710.sub2_sub3:vreg_128_align2, %401.sub2_sub3:vreg_128_align2, %668:vreg_512_align2, 0, 0, 0, implicit $mode, implicit $exec1273 %721:vreg_128_align2 = DS_READ_B128_gfx9 %23:vgpr_32, 512, 0, implicit $exec1274 %679:vreg_512_align2 = contract V_MFMA_F32_32X32X8F16_mac_vgprcd_e64 %721.sub0_sub1:vreg_128_align2, %401.sub0_sub1:vreg_128_align2, %679:vreg_512_align2, 0, 0, 0, implicit $mode, implicit $exec1275 %679:vreg_512_align2 = contract V_MFMA_F32_32X32X8F16_mac_vgprcd_e64 %721.sub2_sub3:vreg_128_align2, %401.sub2_sub3:vreg_128_align2, %679:vreg_512_align2, 0, 0, 0, implicit $mode, implicit $exec1276 %732:vreg_128_align2 = DS_READ_B128_gfx9 %23:vgpr_32, 1024, 0, implicit $exec1277 %690:vreg_512_align2 = contract V_MFMA_F32_32X32X8F16_mac_vgprcd_e64 %732.sub0_sub1:vreg_128_align2, %401.sub0_sub1:vreg_128_align2, %690:vreg_512_align2, 0, 0, 0, implicit $mode, implicit $exec1278 %690:vreg_512_align2 = contract V_MFMA_F32_32X32X8F16_mac_vgprcd_e64 %732.sub2_sub3:vreg_128_align2, %401.sub2_sub3:vreg_128_align2, %690:vreg_512_align2, 0, 0, 0, implicit $mode, implicit $exec1279 %743:vreg_128_align2 = DS_READ_B128_gfx9 %23:vgpr_32, 1536, 0, implicit $exec1280 %701:vreg_512_align2 = contract V_MFMA_F32_32X32X8F16_mac_vgprcd_e64 %743.sub0_sub1:vreg_128_align2, %401.sub0_sub1:vreg_128_align2, %701:vreg_512_align2, 0, 0, 0, implicit $mode, implicit $exec1281 %701:vreg_512_align2 = contract V_MFMA_F32_32X32X8F16_mac_vgprcd_e64 %743.sub2_sub3:vreg_128_align2, %401.sub2_sub3:vreg_128_align2, %701:vreg_512_align2, 0, 0, 0, implicit $mode, implicit $exec1282 %754:vreg_128_align2 = DS_READ_B128_gfx9 %24:vgpr_32, 0, 0, implicit $exec1283 %668:vreg_512_align2 = contract V_MFMA_F32_32X32X8F16_mac_vgprcd_e64 %754.sub0_sub1:vreg_128_align2, %406.sub0_sub1:vreg_128_align2, %668:vreg_512_align2, 0, 0, 0, implicit $mode, implicit $exec1284 %668:vreg_512_align2 = contract V_MFMA_F32_32X32X8F16_mac_vgprcd_e64 %754.sub2_sub3:vreg_128_align2, %406.sub2_sub3:vreg_128_align2, %668:vreg_512_align2, 0, 0, 0, implicit $mode, implicit $exec1285 %765:vreg_128_align2 = DS_READ_B128_gfx9 %24:vgpr_32, 512, 0, implicit $exec1286 %679:vreg_512_align2 = contract V_MFMA_F32_32X32X8F16_mac_vgprcd_e64 %765.sub0_sub1:vreg_128_align2, %406.sub0_sub1:vreg_128_align2, %679:vreg_512_align2, 0, 0, 0, implicit $mode, implicit $exec1287 %679:vreg_512_align2 = contract V_MFMA_F32_32X32X8F16_mac_vgprcd_e64 %765.sub2_sub3:vreg_128_align2, %406.sub2_sub3:vreg_128_align2, %679:vreg_512_align2, 0, 0, 0, implicit $mode, implicit $exec1288 %776:vreg_128_align2 = DS_READ_B128_gfx9 %24:vgpr_32, 1024, 0, implicit $exec1289 %690:vreg_512_align2 = contract V_MFMA_F32_32X32X8F16_mac_vgprcd_e64 %776.sub0_sub1:vreg_128_align2, %406.sub0_sub1:vreg_128_align2, %690:vreg_512_align2, 0, 0, 0, implicit $mode, implicit $exec1290 %690:vreg_512_align2 = contract V_MFMA_F32_32X32X8F16_mac_vgprcd_e64 %776.sub2_sub3:vreg_128_align2, %406.sub2_sub3:vreg_128_align2, %690:vreg_512_align2, 0, 0, 0, implicit $mode, implicit $exec1291 %787:vreg_128_align2 = DS_READ_B128_gfx9 %24:vgpr_32, 1536, 0, implicit $exec1292 %701:vreg_512_align2 = contract V_MFMA_F32_32X32X8F16_mac_vgprcd_e64 %787.sub0_sub1:vreg_128_align2, %406.sub0_sub1:vreg_128_align2, %701:vreg_512_align2, 0, 0, 0, implicit $mode, implicit $exec1293 %701:vreg_512_align2 = contract V_MFMA_F32_32X32X8F16_mac_vgprcd_e64 %787.sub2_sub3:vreg_128_align2, %406.sub2_sub3:vreg_128_align2, %701:vreg_512_align2, 0, 0, 0, implicit $mode, implicit $exec1294 INLINEASM &"s_waitcnt vmcnt($0)", 57 /* sideeffect mayload maystore isconvergent attdialect */, 13 /* imm */, 8, !01295 DS_WRITE_B128_gfx9 %606:vgpr_32, %706:vreg_128_align2, 0, 0, implicit $exec1296 DS_WRITE_B128_gfx9 %606:vgpr_32, %709:vreg_128_align2, 1024, 0, implicit $exec1297 %798:vreg_128_align2 = BUFFER_LOAD_DWORDX4_OFFEN %599:vgpr_32, %443:sgpr_128, 0, 192, 0, 0, implicit $exec1298 %800:vgpr_32 = V_ADD_U32_e32 192, %602:vgpr_32, implicit $exec1299 %801:vreg_128_align2 = BUFFER_LOAD_DWORDX4_OFFEN %800:vgpr_32, %443:sgpr_128, 0, 0, 0, 0, implicit $exec1300 %802:vgpr_32 = V_ADD_U32_e32 %48:vgpr_32, %3330:vgpr_32, implicit $exec1301 %803:vreg_64_align2 = BUFFER_LOAD_DWORDX2_OFFEN %802:vgpr_32, %473:sgpr_128, 0, 0, 0, 0, implicit $exec1302 %804:vgpr_32 = V_ADD_U32_e32 %48:vgpr_32, %3331:vgpr_32, implicit $exec1303 %805:vreg_64_align2 = BUFFER_LOAD_DWORDX2_OFFEN %804:vgpr_32, %473:sgpr_128, 0, 0, 0, 0, implicit $exec1304 %806:vgpr_32 = V_ADD_U32_e32 %48:vgpr_32, %3332:vgpr_32, implicit $exec1305 %807:vreg_64_align2 = BUFFER_LOAD_DWORDX2_OFFEN %806:vgpr_32, %473:sgpr_128, 0, 0, 0, 0, implicit $exec1306 %808:vgpr_32 = V_ADD_U32_e32 %48:vgpr_32, %3333:vgpr_32, implicit $exec1307 %809:vreg_64_align2 = BUFFER_LOAD_DWORDX2_OFFEN %808:vgpr_32, %473:sgpr_128, 0, 0, 0, 0, implicit $exec1308 INLINEASM &"s_waitcnt vmcnt($0)", 57 /* sideeffect mayload maystore isconvergent attdialect */, 13 /* imm */, 8, !01309 %810:vreg_128_align2 = DS_READ_B128_gfx9 %23:vgpr_32, 0, 0, implicit $exec1310 %668:vreg_512_align2 = contract V_MFMA_F32_32X32X8F16_mac_vgprcd_e64 %810.sub0_sub1:vreg_128_align2, %411.sub0_sub1:vreg_128_align2, %668:vreg_512_align2, 0, 0, 0, implicit $mode, implicit $exec1311 %668:vreg_512_align2 = contract V_MFMA_F32_32X32X8F16_mac_vgprcd_e64 %810.sub2_sub3:vreg_128_align2, %411.sub2_sub3:vreg_128_align2, %668:vreg_512_align2, 0, 0, 0, implicit $mode, implicit $exec1312 %821:vreg_128_align2 = DS_READ_B128_gfx9 %23:vgpr_32, 512, 0, implicit $exec1313 %679:vreg_512_align2 = contract V_MFMA_F32_32X32X8F16_mac_vgprcd_e64 %821.sub0_sub1:vreg_128_align2, %411.sub0_sub1:vreg_128_align2, %679:vreg_512_align2, 0, 0, 0, implicit $mode, implicit $exec1314 %679:vreg_512_align2 = contract V_MFMA_F32_32X32X8F16_mac_vgprcd_e64 %821.sub2_sub3:vreg_128_align2, %411.sub2_sub3:vreg_128_align2, %679:vreg_512_align2, 0, 0, 0, implicit $mode, implicit $exec1315 %832:vreg_128_align2 = DS_READ_B128_gfx9 %23:vgpr_32, 1024, 0, implicit $exec1316 %690:vreg_512_align2 = contract V_MFMA_F32_32X32X8F16_mac_vgprcd_e64 %832.sub0_sub1:vreg_128_align2, %411.sub0_sub1:vreg_128_align2, %690:vreg_512_align2, 0, 0, 0, implicit $mode, implicit $exec1317 %690:vreg_512_align2 = contract V_MFMA_F32_32X32X8F16_mac_vgprcd_e64 %832.sub2_sub3:vreg_128_align2, %411.sub2_sub3:vreg_128_align2, %690:vreg_512_align2, 0, 0, 0, implicit $mode, implicit $exec1318 %843:vreg_128_align2 = DS_READ_B128_gfx9 %23:vgpr_32, 1536, 0, implicit $exec1319 %701:vreg_512_align2 = contract V_MFMA_F32_32X32X8F16_mac_vgprcd_e64 %843.sub0_sub1:vreg_128_align2, %411.sub0_sub1:vreg_128_align2, %701:vreg_512_align2, 0, 0, 0, implicit $mode, implicit $exec1320 %701:vreg_512_align2 = contract V_MFMA_F32_32X32X8F16_mac_vgprcd_e64 %843.sub2_sub3:vreg_128_align2, %411.sub2_sub3:vreg_128_align2, %701:vreg_512_align2, 0, 0, 0, implicit $mode, implicit $exec1321 %854:vreg_128_align2 = DS_READ_B128_gfx9 %24:vgpr_32, 0, 0, implicit $exec1322 %668:vreg_512_align2 = contract V_MFMA_F32_32X32X8F16_mac_vgprcd_e64 %854.sub0_sub1:vreg_128_align2, %416.sub0_sub1:vreg_128_align2, %668:vreg_512_align2, 0, 0, 0, implicit $mode, implicit $exec1323 %668:vreg_512_align2 = contract V_MFMA_F32_32X32X8F16_mac_vgprcd_e64 %854.sub2_sub3:vreg_128_align2, %416.sub2_sub3:vreg_128_align2, %668:vreg_512_align2, 0, 0, 0, implicit $mode, implicit $exec1324 %865:vreg_128_align2 = DS_READ_B128_gfx9 %24:vgpr_32, 512, 0, implicit $exec1325 %679:vreg_512_align2 = contract V_MFMA_F32_32X32X8F16_mac_vgprcd_e64 %865.sub0_sub1:vreg_128_align2, %416.sub0_sub1:vreg_128_align2, %679:vreg_512_align2, 0, 0, 0, implicit $mode, implicit $exec1326 %679:vreg_512_align2 = contract V_MFMA_F32_32X32X8F16_mac_vgprcd_e64 %865.sub2_sub3:vreg_128_align2, %416.sub2_sub3:vreg_128_align2, %679:vreg_512_align2, 0, 0, 0, implicit $mode, implicit $exec1327 %876:vreg_128_align2 = DS_READ_B128_gfx9 %24:vgpr_32, 1024, 0, implicit $exec1328 %690:vreg_512_align2 = contract V_MFMA_F32_32X32X8F16_mac_vgprcd_e64 %876.sub0_sub1:vreg_128_align2, %416.sub0_sub1:vreg_128_align2, %690:vreg_512_align2, 0, 0, 0, implicit $mode, implicit $exec1329 %690:vreg_512_align2 = contract V_MFMA_F32_32X32X8F16_mac_vgprcd_e64 %876.sub2_sub3:vreg_128_align2, %416.sub2_sub3:vreg_128_align2, %690:vreg_512_align2, 0, 0, 0, implicit $mode, implicit $exec1330 %887:vreg_128_align2 = DS_READ_B128_gfx9 %24:vgpr_32, 1536, 0, implicit $exec1331 %701:vreg_512_align2 = contract V_MFMA_F32_32X32X8F16_mac_vgprcd_e64 %887.sub0_sub1:vreg_128_align2, %416.sub0_sub1:vreg_128_align2, %701:vreg_512_align2, 0, 0, 0, implicit $mode, implicit $exec1332 %701:vreg_512_align2 = contract V_MFMA_F32_32X32X8F16_mac_vgprcd_e64 %887.sub2_sub3:vreg_128_align2, %416.sub2_sub3:vreg_128_align2, %701:vreg_512_align2, 0, 0, 0, implicit $mode, implicit $exec1333 INLINEASM &"s_waitcnt vmcnt($0)", 57 /* sideeffect mayload maystore isconvergent attdialect */, 13 /* imm */, 8, !01334 DS_WRITE_B128_gfx9 %606:vgpr_32, %798:vreg_128_align2, 0, 0, implicit $exec1335 DS_WRITE_B128_gfx9 %606:vgpr_32, %801:vreg_128_align2, 1024, 0, implicit $exec1336 INLINEASM &"s_waitcnt vmcnt($0)", 57 /* sideeffect mayload maystore isconvergent attdialect */, 13 /* imm */, 8, !01337 %898:vreg_128_align2 = DS_READ_B128_gfx9 %23:vgpr_32, 0, 0, implicit $exec1338 %668:vreg_512_align2 = contract V_MFMA_F32_32X32X8F16_mac_vgprcd_e64 %898.sub0_sub1:vreg_128_align2, %421.sub0_sub1:vreg_128_align2, %668:vreg_512_align2, 0, 0, 0, implicit $mode, implicit $exec1339 %668:vreg_512_align2 = contract V_MFMA_F32_32X32X8F16_mac_vgprcd_e64 %898.sub2_sub3:vreg_128_align2, %421.sub2_sub3:vreg_128_align2, %668:vreg_512_align2, 0, 0, 0, implicit $mode, implicit $exec1340 %909:vreg_128_align2 = DS_READ_B128_gfx9 %23:vgpr_32, 512, 0, implicit $exec1341 %679:vreg_512_align2 = contract V_MFMA_F32_32X32X8F16_mac_vgprcd_e64 %909.sub0_sub1:vreg_128_align2, %421.sub0_sub1:vreg_128_align2, %679:vreg_512_align2, 0, 0, 0, implicit $mode, implicit $exec1342 %679:vreg_512_align2 = contract V_MFMA_F32_32X32X8F16_mac_vgprcd_e64 %909.sub2_sub3:vreg_128_align2, %421.sub2_sub3:vreg_128_align2, %679:vreg_512_align2, 0, 0, 0, implicit $mode, implicit $exec1343 %920:vreg_128_align2 = DS_READ_B128_gfx9 %23:vgpr_32, 1024, 0, implicit $exec1344 %690:vreg_512_align2 = contract V_MFMA_F32_32X32X8F16_mac_vgprcd_e64 %920.sub0_sub1:vreg_128_align2, %421.sub0_sub1:vreg_128_align2, %690:vreg_512_align2, 0, 0, 0, implicit $mode, implicit $exec1345 %690:vreg_512_align2 = contract V_MFMA_F32_32X32X8F16_mac_vgprcd_e64 %920.sub2_sub3:vreg_128_align2, %421.sub2_sub3:vreg_128_align2, %690:vreg_512_align2, 0, 0, 0, implicit $mode, implicit $exec1346 %931:vreg_128_align2 = DS_READ_B128_gfx9 %23:vgpr_32, 1536, 0, implicit $exec1347 %701:vreg_512_align2 = contract V_MFMA_F32_32X32X8F16_mac_vgprcd_e64 %931.sub0_sub1:vreg_128_align2, %421.sub0_sub1:vreg_128_align2, %701:vreg_512_align2, 0, 0, 0, implicit $mode, implicit $exec1348 %701:vreg_512_align2 = contract V_MFMA_F32_32X32X8F16_mac_vgprcd_e64 %931.sub2_sub3:vreg_128_align2, %421.sub2_sub3:vreg_128_align2, %701:vreg_512_align2, 0, 0, 0, implicit $mode, implicit $exec1349 %942:vreg_128_align2 = DS_READ_B128_gfx9 %24:vgpr_32, 0, 0, implicit $exec1350 %668:vreg_512_align2 = contract V_MFMA_F32_32X32X8F16_mac_vgprcd_e64 %942.sub0_sub1:vreg_128_align2, %426.sub0_sub1:vreg_128_align2, %668:vreg_512_align2, 0, 0, 0, implicit $mode, implicit $exec1351 %668:vreg_512_align2 = contract V_MFMA_F32_32X32X8F16_mac_vgprcd_e64 %942.sub2_sub3:vreg_128_align2, %426.sub2_sub3:vreg_128_align2, %668:vreg_512_align2, 0, 0, 0, implicit $mode, implicit $exec1352 %969:vreg_128_align2 = DS_READ_B128_gfx9 %24:vgpr_32, 512, 0, implicit $exec1353 %679:vreg_512_align2 = contract V_MFMA_F32_32X32X8F16_mac_vgprcd_e64 %969.sub0_sub1:vreg_128_align2, %426.sub0_sub1:vreg_128_align2, %679:vreg_512_align2, 0, 0, 0, implicit $mode, implicit $exec1354 %679:vreg_512_align2 = contract V_MFMA_F32_32X32X8F16_mac_vgprcd_e64 %969.sub2_sub3:vreg_128_align2, %426.sub2_sub3:vreg_128_align2, %679:vreg_512_align2, 0, 0, 0, implicit $mode, implicit $exec1355 %996:vreg_128_align2 = DS_READ_B128_gfx9 %24:vgpr_32, 1024, 0, implicit $exec1356 %690:vreg_512_align2 = contract V_MFMA_F32_32X32X8F16_mac_vgprcd_e64 %996.sub0_sub1:vreg_128_align2, %426.sub0_sub1:vreg_128_align2, %690:vreg_512_align2, 0, 0, 0, implicit $mode, implicit $exec1357 %690:vreg_512_align2 = contract V_MFMA_F32_32X32X8F16_mac_vgprcd_e64 %996.sub2_sub3:vreg_128_align2, %426.sub2_sub3:vreg_128_align2, %690:vreg_512_align2, 0, 0, 0, implicit $mode, implicit $exec1358 %1023:vreg_128_align2 = DS_READ_B128_gfx9 %24:vgpr_32, 1536, 0, implicit $exec1359 %701:vreg_512_align2 = contract V_MFMA_F32_32X32X8F16_mac_vgprcd_e64 %1023.sub0_sub1:vreg_128_align2, %426.sub0_sub1:vreg_128_align2, %701:vreg_512_align2, 0, 0, 0, implicit $mode, implicit $exec1360 %701:vreg_512_align2 = contract V_MFMA_F32_32X32X8F16_mac_vgprcd_e64 %1023.sub2_sub3:vreg_128_align2, %426.sub2_sub3:vreg_128_align2, %701:vreg_512_align2, 0, 0, 0, implicit $mode, implicit $exec1361 %1050:vgpr_32 = contract nofpexcept V_MUL_F32_e32 %1.sub4:sgpr_512, %668.sub0:vreg_512_align2, implicit $mode, implicit $exec1362 %1051:vgpr_32 = contract nofpexcept V_MUL_F32_e32 %1.sub4:sgpr_512, %668.sub1:vreg_512_align2, implicit $mode, implicit $exec1363 %1052:vgpr_32 = contract nofpexcept V_MUL_F32_e32 %1.sub4:sgpr_512, %668.sub2:vreg_512_align2, implicit $mode, implicit $exec1364 %1053:vgpr_32 = contract nofpexcept V_MUL_F32_e32 %1.sub4:sgpr_512, %668.sub3:vreg_512_align2, implicit $mode, implicit $exec1365 %1054:vgpr_32 = contract nofpexcept V_MUL_F32_e32 %1.sub4:sgpr_512, %668.sub4:vreg_512_align2, implicit $mode, implicit $exec1366 %1055:vgpr_32 = contract nofpexcept V_MUL_F32_e32 %1.sub4:sgpr_512, %668.sub5:vreg_512_align2, implicit $mode, implicit $exec1367 %1056:vgpr_32 = contract nofpexcept V_MUL_F32_e32 %1.sub4:sgpr_512, %668.sub6:vreg_512_align2, implicit $mode, implicit $exec1368 %1057:vgpr_32 = contract nofpexcept V_MUL_F32_e32 %1.sub4:sgpr_512, %668.sub7:vreg_512_align2, implicit $mode, implicit $exec1369 %1058:vgpr_32 = contract nofpexcept V_MUL_F32_e32 %1.sub4:sgpr_512, %668.sub8:vreg_512_align2, implicit $mode, implicit $exec1370 %1059:vgpr_32 = contract nofpexcept V_MUL_F32_e32 %1.sub4:sgpr_512, %668.sub9:vreg_512_align2, implicit $mode, implicit $exec1371 %1060:vgpr_32 = contract nofpexcept V_MUL_F32_e32 %1.sub4:sgpr_512, %668.sub10:vreg_512_align2, implicit $mode, implicit $exec1372 %1061:vgpr_32 = contract nofpexcept V_MUL_F32_e32 %1.sub4:sgpr_512, %668.sub11:vreg_512_align2, implicit $mode, implicit $exec1373 %1062:vgpr_32 = contract nofpexcept V_MUL_F32_e32 %1.sub4:sgpr_512, %668.sub12:vreg_512_align2, implicit $mode, implicit $exec1374 %1063:vgpr_32 = contract nofpexcept V_MUL_F32_e32 %1.sub4:sgpr_512, %668.sub13:vreg_512_align2, implicit $mode, implicit $exec1375 %1064:vgpr_32 = contract nofpexcept V_MUL_F32_e32 %1.sub4:sgpr_512, %668.sub14:vreg_512_align2, implicit $mode, implicit $exec1376 %1065:vgpr_32 = contract nofpexcept V_MUL_F32_e32 %1.sub4:sgpr_512, %668.sub15:vreg_512_align2, implicit $mode, implicit $exec1377 %1066:vgpr_32 = contract nofpexcept V_MUL_F32_e32 %1.sub4:sgpr_512, %679.sub0:vreg_512_align2, implicit $mode, implicit $exec1378 %1067:vgpr_32 = contract nofpexcept V_MUL_F32_e32 %1.sub4:sgpr_512, %679.sub1:vreg_512_align2, implicit $mode, implicit $exec1379 %1068:vgpr_32 = contract nofpexcept V_MUL_F32_e32 %1.sub4:sgpr_512, %679.sub2:vreg_512_align2, implicit $mode, implicit $exec1380 %1069:vgpr_32 = contract nofpexcept V_MUL_F32_e32 %1.sub4:sgpr_512, %679.sub3:vreg_512_align2, implicit $mode, implicit $exec1381 %1070:vgpr_32 = contract nofpexcept V_MUL_F32_e32 %1.sub4:sgpr_512, %679.sub4:vreg_512_align2, implicit $mode, implicit $exec1382 %1071:vgpr_32 = contract nofpexcept V_MUL_F32_e32 %1.sub4:sgpr_512, %679.sub5:vreg_512_align2, implicit $mode, implicit $exec1383 %1072:vgpr_32 = contract nofpexcept V_MUL_F32_e32 %1.sub4:sgpr_512, %679.sub6:vreg_512_align2, implicit $mode, implicit $exec1384 %1073:vgpr_32 = contract nofpexcept V_MUL_F32_e32 %1.sub4:sgpr_512, %679.sub7:vreg_512_align2, implicit $mode, implicit $exec1385 %1074:vgpr_32 = contract nofpexcept V_MUL_F32_e32 %1.sub4:sgpr_512, %679.sub8:vreg_512_align2, implicit $mode, implicit $exec1386 %1075:vgpr_32 = contract nofpexcept V_MUL_F32_e32 %1.sub4:sgpr_512, %679.sub9:vreg_512_align2, implicit $mode, implicit $exec1387 %1076:vgpr_32 = contract nofpexcept V_MUL_F32_e32 %1.sub4:sgpr_512, %679.sub10:vreg_512_align2, implicit $mode, implicit $exec1388 %1077:vgpr_32 = contract nofpexcept V_MUL_F32_e32 %1.sub4:sgpr_512, %679.sub11:vreg_512_align2, implicit $mode, implicit $exec1389 %1078:vgpr_32 = contract nofpexcept V_MUL_F32_e32 %1.sub4:sgpr_512, %679.sub12:vreg_512_align2, implicit $mode, implicit $exec1390 %1079:vgpr_32 = contract nofpexcept V_MUL_F32_e32 %1.sub4:sgpr_512, %679.sub13:vreg_512_align2, implicit $mode, implicit $exec1391 %1080:vgpr_32 = contract nofpexcept V_MUL_F32_e32 %1.sub4:sgpr_512, %679.sub14:vreg_512_align2, implicit $mode, implicit $exec1392 %1081:vgpr_32 = contract nofpexcept V_MUL_F32_e32 %1.sub4:sgpr_512, %679.sub15:vreg_512_align2, implicit $mode, implicit $exec1393 %1082:vgpr_32 = contract nofpexcept V_MUL_F32_e32 %1.sub4:sgpr_512, %690.sub0:vreg_512_align2, implicit $mode, implicit $exec1394 %1083:vgpr_32 = contract nofpexcept V_MUL_F32_e32 %1.sub4:sgpr_512, %690.sub1:vreg_512_align2, implicit $mode, implicit $exec1395 %1084:vgpr_32 = contract nofpexcept V_MUL_F32_e32 %1.sub4:sgpr_512, %690.sub2:vreg_512_align2, implicit $mode, implicit $exec1396 %1085:vgpr_32 = contract nofpexcept V_MUL_F32_e32 %1.sub4:sgpr_512, %690.sub3:vreg_512_align2, implicit $mode, implicit $exec1397 %1086:vgpr_32 = contract nofpexcept V_MUL_F32_e32 %1.sub4:sgpr_512, %690.sub4:vreg_512_align2, implicit $mode, implicit $exec1398 %1087:vgpr_32 = contract nofpexcept V_MUL_F32_e32 %1.sub4:sgpr_512, %690.sub5:vreg_512_align2, implicit $mode, implicit $exec1399 %1088:vgpr_32 = contract nofpexcept V_MUL_F32_e32 %1.sub4:sgpr_512, %690.sub6:vreg_512_align2, implicit $mode, implicit $exec1400 %1089:vgpr_32 = contract nofpexcept V_MUL_F32_e32 %1.sub4:sgpr_512, %690.sub7:vreg_512_align2, implicit $mode, implicit $exec1401 %1090:vgpr_32 = contract nofpexcept V_MUL_F32_e32 %1.sub4:sgpr_512, %690.sub8:vreg_512_align2, implicit $mode, implicit $exec1402 %1091:vgpr_32 = contract nofpexcept V_MUL_F32_e32 %1.sub4:sgpr_512, %690.sub9:vreg_512_align2, implicit $mode, implicit $exec1403 %1092:vgpr_32 = contract nofpexcept V_MUL_F32_e32 %1.sub4:sgpr_512, %690.sub10:vreg_512_align2, implicit $mode, implicit $exec1404 %1093:vgpr_32 = contract nofpexcept V_MUL_F32_e32 %1.sub4:sgpr_512, %690.sub11:vreg_512_align2, implicit $mode, implicit $exec1405 %1094:vgpr_32 = contract nofpexcept V_MUL_F32_e32 %1.sub4:sgpr_512, %690.sub12:vreg_512_align2, implicit $mode, implicit $exec1406 %1095:vgpr_32 = contract nofpexcept V_MUL_F32_e32 %1.sub4:sgpr_512, %690.sub13:vreg_512_align2, implicit $mode, implicit $exec1407 %1096:vgpr_32 = contract nofpexcept V_MUL_F32_e32 %1.sub4:sgpr_512, %690.sub14:vreg_512_align2, implicit $mode, implicit $exec1408 %1097:vgpr_32 = contract nofpexcept V_MUL_F32_e32 %1.sub4:sgpr_512, %690.sub15:vreg_512_align2, implicit $mode, implicit $exec1409 %1098:vgpr_32 = contract nofpexcept V_MUL_F32_e32 %1.sub4:sgpr_512, %701.sub0:vreg_512_align2, implicit $mode, implicit $exec1410 %1099:vgpr_32 = contract nofpexcept V_MUL_F32_e32 %1.sub4:sgpr_512, %701.sub1:vreg_512_align2, implicit $mode, implicit $exec1411 %1100:vgpr_32 = contract nofpexcept V_MUL_F32_e32 %1.sub4:sgpr_512, %701.sub2:vreg_512_align2, implicit $mode, implicit $exec1412 %1101:vgpr_32 = contract nofpexcept V_MUL_F32_e32 %1.sub4:sgpr_512, %701.sub3:vreg_512_align2, implicit $mode, implicit $exec1413 %1102:vgpr_32 = contract nofpexcept V_MUL_F32_e32 %1.sub4:sgpr_512, %701.sub4:vreg_512_align2, implicit $mode, implicit $exec1414 %1103:vgpr_32 = contract nofpexcept V_MUL_F32_e32 %1.sub4:sgpr_512, %701.sub5:vreg_512_align2, implicit $mode, implicit $exec1415 %1104:vgpr_32 = contract nofpexcept V_MUL_F32_e32 %1.sub4:sgpr_512, %701.sub6:vreg_512_align2, implicit $mode, implicit $exec1416 %1105:vgpr_32 = contract nofpexcept V_MUL_F32_e32 %1.sub4:sgpr_512, %701.sub7:vreg_512_align2, implicit $mode, implicit $exec1417 %1106:vgpr_32 = contract nofpexcept V_MUL_F32_e32 %1.sub4:sgpr_512, %701.sub8:vreg_512_align2, implicit $mode, implicit $exec1418 %1107:vgpr_32 = contract nofpexcept V_MUL_F32_e32 %1.sub4:sgpr_512, %701.sub9:vreg_512_align2, implicit $mode, implicit $exec1419 %1108:vgpr_32 = contract nofpexcept V_MUL_F32_e32 %1.sub4:sgpr_512, %701.sub10:vreg_512_align2, implicit $mode, implicit $exec1420 %1109:vgpr_32 = contract nofpexcept V_MUL_F32_e32 %1.sub4:sgpr_512, %701.sub11:vreg_512_align2, implicit $mode, implicit $exec1421 %1110:vgpr_32 = contract nofpexcept V_MUL_F32_e32 %1.sub4:sgpr_512, %701.sub12:vreg_512_align2, implicit $mode, implicit $exec1422 %1111:vgpr_32 = contract nofpexcept V_MUL_F32_e32 %1.sub4:sgpr_512, %701.sub13:vreg_512_align2, implicit $mode, implicit $exec1423 %1112:vgpr_32 = contract nofpexcept V_MUL_F32_e32 %1.sub4:sgpr_512, %701.sub14:vreg_512_align2, implicit $mode, implicit $exec1424 %1113:vgpr_32 = contract nofpexcept V_MUL_F32_e32 %1.sub4:sgpr_512, %701.sub15:vreg_512_align2, implicit $mode, implicit $exec1425 %1115:vgpr_32 = V_MAX3_F32_e64 0, %1050:vgpr_32, 0, %1114:sgpr_32, 0, %1051:vgpr_32, 0, 0, implicit $mode, implicit $exec1426 %1116:vgpr_32 = V_MAX3_F32_e64 0, %1115:vgpr_32, 0, %1052:vgpr_32, 0, %1053:vgpr_32, 0, 0, implicit $mode, implicit $exec1427 %1117:vgpr_32 = V_MAX3_F32_e64 0, %1116:vgpr_32, 0, %1054:vgpr_32, 0, %1055:vgpr_32, 0, 0, implicit $mode, implicit $exec1428 %1118:vgpr_32 = V_MAX3_F32_e64 0, %1117:vgpr_32, 0, %1056:vgpr_32, 0, %1057:vgpr_32, 0, 0, implicit $mode, implicit $exec1429 %1119:vgpr_32 = V_MAX3_F32_e64 0, %1118:vgpr_32, 0, %1058:vgpr_32, 0, %1059:vgpr_32, 0, 0, implicit $mode, implicit $exec1430 %1120:vgpr_32 = V_MAX3_F32_e64 0, %1119:vgpr_32, 0, %1060:vgpr_32, 0, %1061:vgpr_32, 0, 0, implicit $mode, implicit $exec1431 %1121:vgpr_32 = V_MAX3_F32_e64 0, %1120:vgpr_32, 0, %1062:vgpr_32, 0, %1063:vgpr_32, 0, 0, implicit $mode, implicit $exec1432 %1122:vgpr_32 = V_MAX3_F32_e64 0, %1121:vgpr_32, 0, %1064:vgpr_32, 0, %1065:vgpr_32, 0, 0, implicit $mode, implicit $exec1433 %1123:vgpr_32 = V_MAX3_F32_e64 0, %1122:vgpr_32, 0, %1066:vgpr_32, 0, %1067:vgpr_32, 0, 0, implicit $mode, implicit $exec1434 %1124:vgpr_32 = V_MAX3_F32_e64 0, %1123:vgpr_32, 0, %1068:vgpr_32, 0, %1069:vgpr_32, 0, 0, implicit $mode, implicit $exec1435 %1125:vgpr_32 = V_MAX3_F32_e64 0, %1124:vgpr_32, 0, %1070:vgpr_32, 0, %1071:vgpr_32, 0, 0, implicit $mode, implicit $exec1436 %1126:vgpr_32 = V_MAX3_F32_e64 0, %1125:vgpr_32, 0, %1072:vgpr_32, 0, %1073:vgpr_32, 0, 0, implicit $mode, implicit $exec1437 %1127:vgpr_32 = V_MAX3_F32_e64 0, %1126:vgpr_32, 0, %1074:vgpr_32, 0, %1075:vgpr_32, 0, 0, implicit $mode, implicit $exec1438 %1128:vgpr_32 = V_MAX3_F32_e64 0, %1127:vgpr_32, 0, %1076:vgpr_32, 0, %1077:vgpr_32, 0, 0, implicit $mode, implicit $exec1439 %1129:vgpr_32 = V_MAX3_F32_e64 0, %1128:vgpr_32, 0, %1078:vgpr_32, 0, %1079:vgpr_32, 0, 0, implicit $mode, implicit $exec1440 %1130:vgpr_32 = V_MAX3_F32_e64 0, %1129:vgpr_32, 0, %1080:vgpr_32, 0, %1081:vgpr_32, 0, 0, implicit $mode, implicit $exec1441 %1131:vgpr_32 = V_MAX3_F32_e64 0, %1130:vgpr_32, 0, %1082:vgpr_32, 0, %1083:vgpr_32, 0, 0, implicit $mode, implicit $exec1442 %1132:vgpr_32 = V_MAX3_F32_e64 0, %1131:vgpr_32, 0, %1084:vgpr_32, 0, %1085:vgpr_32, 0, 0, implicit $mode, implicit $exec1443 %1133:vgpr_32 = V_MAX3_F32_e64 0, %1132:vgpr_32, 0, %1086:vgpr_32, 0, %1087:vgpr_32, 0, 0, implicit $mode, implicit $exec1444 %1134:vgpr_32 = V_MAX3_F32_e64 0, %1133:vgpr_32, 0, %1088:vgpr_32, 0, %1089:vgpr_32, 0, 0, implicit $mode, implicit $exec1445 %1135:vgpr_32 = V_MAX3_F32_e64 0, %1134:vgpr_32, 0, %1090:vgpr_32, 0, %1091:vgpr_32, 0, 0, implicit $mode, implicit $exec1446 %1136:vgpr_32 = V_MAX3_F32_e64 0, %1135:vgpr_32, 0, %1092:vgpr_32, 0, %1093:vgpr_32, 0, 0, implicit $mode, implicit $exec1447 %1137:vgpr_32 = V_MAX3_F32_e64 0, %1136:vgpr_32, 0, %1094:vgpr_32, 0, %1095:vgpr_32, 0, 0, implicit $mode, implicit $exec1448 %1138:vgpr_32 = V_MAX3_F32_e64 0, %1137:vgpr_32, 0, %1096:vgpr_32, 0, %1097:vgpr_32, 0, 0, implicit $mode, implicit $exec1449 %1139:vgpr_32 = V_MAX3_F32_e64 0, %1138:vgpr_32, 0, %1098:vgpr_32, 0, %1099:vgpr_32, 0, 0, implicit $mode, implicit $exec1450 %1140:vgpr_32 = V_MAX3_F32_e64 0, %1139:vgpr_32, 0, %1100:vgpr_32, 0, %1101:vgpr_32, 0, 0, implicit $mode, implicit $exec1451 %1141:vgpr_32 = V_MAX3_F32_e64 0, %1140:vgpr_32, 0, %1102:vgpr_32, 0, %1103:vgpr_32, 0, 0, implicit $mode, implicit $exec1452 %1142:vgpr_32 = V_MAX3_F32_e64 0, %1141:vgpr_32, 0, %1104:vgpr_32, 0, %1105:vgpr_32, 0, 0, implicit $mode, implicit $exec1453 %1143:vgpr_32 = V_MAX3_F32_e64 0, %1142:vgpr_32, 0, %1106:vgpr_32, 0, %1107:vgpr_32, 0, 0, implicit $mode, implicit $exec1454 %1144:vgpr_32 = V_MAX3_F32_e64 0, %1143:vgpr_32, 0, %1108:vgpr_32, 0, %1109:vgpr_32, 0, 0, implicit $mode, implicit $exec1455 %1145:vgpr_32 = V_MAX3_F32_e64 0, %1144:vgpr_32, 0, %1110:vgpr_32, 0, %1111:vgpr_32, 0, 0, implicit $mode, implicit $exec1456 %1146:vgpr_32 = V_MAX3_F32_e64 0, %1145:vgpr_32, 0, %1112:vgpr_32, 0, %1113:vgpr_32, 0, 0, implicit $mode, implicit $exec1457 %1147:vgpr_32 = DS_BPERMUTE_B32 %39:vgpr_32, %1146:vgpr_32, 0, implicit $exec1458 %1148:vgpr_32 = contract nofpexcept V_MAX_F32_e32 %1147:vgpr_32, %1147:vgpr_32, implicit $mode, implicit $exec1459 %1149:vgpr_32 = contract nofpexcept V_MAX_F32_e32 %1146:vgpr_32, %1148:vgpr_32, implicit $mode, implicit $exec1460 %1150:vgpr_32 = DS_BPERMUTE_B32 %39:vgpr_32, %1149:vgpr_32, 0, implicit $exec1461 %1151:vgpr_32 = V_CNDMASK_B32_e64 0, %1150:vgpr_32, 0, %1149:vgpr_32, %484:sreg_64_xexec, implicit $exec1462 %1153:vgpr_32 = contract nofpexcept V_MAX_F32_e32 %1151:vgpr_32, %1151:vgpr_32, implicit $mode, implicit $exec1463 %1154:vgpr_32 = contract nofpexcept V_MAX_F32_e32 %3346:vgpr_32, %3346:vgpr_32, implicit $mode, implicit $exec1464 %151:vgpr_32 = contract nofpexcept V_MAX_F32_e32 %1154:vgpr_32, %1153:vgpr_32, implicit $mode, implicit $exec1465 %1155:vgpr_32 = contract nofpexcept V_FMA_F32_e64 0, %1.sub4:sgpr_512, 0, %668.sub0:vreg_512_align2, 1, %151:vgpr_32, 0, 0, implicit $mode, implicit $exec1466 %1157:vgpr_32 = afn nofpexcept V_MUL_F32_e32 1069066811, %1155:vgpr_32, implicit $mode, implicit $exec1467 %1158:vgpr_32 = afn nofpexcept V_EXP_F32_e32 %1157:vgpr_32, implicit $mode, implicit $exec1468 %1159:vgpr_32 = contract nofpexcept V_FMA_F32_e64 0, %1.sub4:sgpr_512, 0, %668.sub1:vreg_512_align2, 1, %151:vgpr_32, 0, 0, implicit $mode, implicit $exec1469 %1160:vgpr_32 = afn nofpexcept V_MUL_F32_e32 1069066811, %1159:vgpr_32, implicit $mode, implicit $exec1470 %1161:vgpr_32 = afn nofpexcept V_EXP_F32_e32 %1160:vgpr_32, implicit $mode, implicit $exec1471 %1162:vgpr_32 = contract nofpexcept V_FMA_F32_e64 0, %1.sub4:sgpr_512, 0, %668.sub2:vreg_512_align2, 1, %151:vgpr_32, 0, 0, implicit $mode, implicit $exec1472 %1163:vgpr_32 = afn nofpexcept V_MUL_F32_e32 1069066811, %1162:vgpr_32, implicit $mode, implicit $exec1473 %1164:vgpr_32 = afn nofpexcept V_EXP_F32_e32 %1163:vgpr_32, implicit $mode, implicit $exec1474 %1165:vgpr_32 = contract nofpexcept V_FMA_F32_e64 0, %1.sub4:sgpr_512, 0, %668.sub3:vreg_512_align2, 1, %151:vgpr_32, 0, 0, implicit $mode, implicit $exec1475 %1166:vgpr_32 = afn nofpexcept V_MUL_F32_e32 1069066811, %1165:vgpr_32, implicit $mode, implicit $exec1476 %1167:vgpr_32 = afn nofpexcept V_EXP_F32_e32 %1166:vgpr_32, implicit $mode, implicit $exec1477 %1168:vgpr_32 = contract nofpexcept V_FMA_F32_e64 0, %1.sub4:sgpr_512, 0, %668.sub4:vreg_512_align2, 1, %151:vgpr_32, 0, 0, implicit $mode, implicit $exec1478 %1169:vgpr_32 = afn nofpexcept V_MUL_F32_e32 1069066811, %1168:vgpr_32, implicit $mode, implicit $exec1479 %1170:vgpr_32 = afn nofpexcept V_EXP_F32_e32 %1169:vgpr_32, implicit $mode, implicit $exec1480 %1171:vgpr_32 = contract nofpexcept V_FMA_F32_e64 0, %1.sub4:sgpr_512, 0, %668.sub5:vreg_512_align2, 1, %151:vgpr_32, 0, 0, implicit $mode, implicit $exec1481 %1172:vgpr_32 = afn nofpexcept V_MUL_F32_e32 1069066811, %1171:vgpr_32, implicit $mode, implicit $exec1482 %1173:vgpr_32 = afn nofpexcept V_EXP_F32_e32 %1172:vgpr_32, implicit $mode, implicit $exec1483 %1174:vgpr_32 = contract nofpexcept V_FMA_F32_e64 0, %1.sub4:sgpr_512, 0, %668.sub6:vreg_512_align2, 1, %151:vgpr_32, 0, 0, implicit $mode, implicit $exec1484 %1175:vgpr_32 = afn nofpexcept V_MUL_F32_e32 1069066811, %1174:vgpr_32, implicit $mode, implicit $exec1485 %1176:vgpr_32 = afn nofpexcept V_EXP_F32_e32 %1175:vgpr_32, implicit $mode, implicit $exec1486 %1177:vgpr_32 = contract nofpexcept V_FMA_F32_e64 0, %1.sub4:sgpr_512, 0, %668.sub7:vreg_512_align2, 1, %151:vgpr_32, 0, 0, implicit $mode, implicit $exec1487 %1178:vgpr_32 = afn nofpexcept V_MUL_F32_e32 1069066811, %1177:vgpr_32, implicit $mode, implicit $exec1488 %1179:vgpr_32 = afn nofpexcept V_EXP_F32_e32 %1178:vgpr_32, implicit $mode, implicit $exec1489 %1180:vgpr_32 = contract nofpexcept V_FMA_F32_e64 0, %1.sub4:sgpr_512, 0, %668.sub8:vreg_512_align2, 1, %151:vgpr_32, 0, 0, implicit $mode, implicit $exec1490 %1181:vgpr_32 = afn nofpexcept V_MUL_F32_e32 1069066811, %1180:vgpr_32, implicit $mode, implicit $exec1491 %1182:vgpr_32 = afn nofpexcept V_EXP_F32_e32 %1181:vgpr_32, implicit $mode, implicit $exec1492 %1183:vgpr_32 = contract nofpexcept V_FMA_F32_e64 0, %1.sub4:sgpr_512, 0, %668.sub9:vreg_512_align2, 1, %151:vgpr_32, 0, 0, implicit $mode, implicit $exec1493 %1184:vgpr_32 = afn nofpexcept V_MUL_F32_e32 1069066811, %1183:vgpr_32, implicit $mode, implicit $exec1494 %1185:vgpr_32 = afn nofpexcept V_EXP_F32_e32 %1184:vgpr_32, implicit $mode, implicit $exec1495 %1186:vgpr_32 = contract nofpexcept V_FMA_F32_e64 0, %1.sub4:sgpr_512, 0, %668.sub10:vreg_512_align2, 1, %151:vgpr_32, 0, 0, implicit $mode, implicit $exec1496 %1187:vgpr_32 = afn nofpexcept V_MUL_F32_e32 1069066811, %1186:vgpr_32, implicit $mode, implicit $exec1497 %1188:vgpr_32 = afn nofpexcept V_EXP_F32_e32 %1187:vgpr_32, implicit $mode, implicit $exec1498 %1189:vgpr_32 = contract nofpexcept V_FMA_F32_e64 0, %1.sub4:sgpr_512, 0, %668.sub11:vreg_512_align2, 1, %151:vgpr_32, 0, 0, implicit $mode, implicit $exec1499 %1190:vgpr_32 = afn nofpexcept V_MUL_F32_e32 1069066811, %1189:vgpr_32, implicit $mode, implicit $exec1500 %1191:vgpr_32 = afn nofpexcept V_EXP_F32_e32 %1190:vgpr_32, implicit $mode, implicit $exec1501 %1192:vgpr_32 = contract nofpexcept V_FMA_F32_e64 0, %1.sub4:sgpr_512, 0, %668.sub12:vreg_512_align2, 1, %151:vgpr_32, 0, 0, implicit $mode, implicit $exec1502 %1193:vgpr_32 = afn nofpexcept V_MUL_F32_e32 1069066811, %1192:vgpr_32, implicit $mode, implicit $exec1503 %1194:vgpr_32 = afn nofpexcept V_EXP_F32_e32 %1193:vgpr_32, implicit $mode, implicit $exec1504 %1195:vgpr_32 = contract nofpexcept V_FMA_F32_e64 0, %1.sub4:sgpr_512, 0, %668.sub13:vreg_512_align2, 1, %151:vgpr_32, 0, 0, implicit $mode, implicit $exec1505 %1196:vgpr_32 = afn nofpexcept V_MUL_F32_e32 1069066811, %1195:vgpr_32, implicit $mode, implicit $exec1506 %1197:vgpr_32 = afn nofpexcept V_EXP_F32_e32 %1196:vgpr_32, implicit $mode, implicit $exec1507 %1198:vgpr_32 = contract nofpexcept V_FMA_F32_e64 0, %1.sub4:sgpr_512, 0, %668.sub14:vreg_512_align2, 1, %151:vgpr_32, 0, 0, implicit $mode, implicit $exec1508 %1199:vgpr_32 = afn nofpexcept V_MUL_F32_e32 1069066811, %1198:vgpr_32, implicit $mode, implicit $exec1509 %1200:vgpr_32 = afn nofpexcept V_EXP_F32_e32 %1199:vgpr_32, implicit $mode, implicit $exec1510 %1201:vgpr_32 = contract nofpexcept V_FMA_F32_e64 0, %1.sub4:sgpr_512, 0, %668.sub15:vreg_512_align2, 1, %151:vgpr_32, 0, 0, implicit $mode, implicit $exec1511 %1202:vgpr_32 = afn nofpexcept V_MUL_F32_e32 1069066811, %1201:vgpr_32, implicit $mode, implicit $exec1512 %1203:vgpr_32 = afn nofpexcept V_EXP_F32_e32 %1202:vgpr_32, implicit $mode, implicit $exec1513 %1204:vgpr_32 = contract nofpexcept V_FMA_F32_e64 0, %1.sub4:sgpr_512, 0, %679.sub0:vreg_512_align2, 1, %151:vgpr_32, 0, 0, implicit $mode, implicit $exec1514 %1205:vgpr_32 = afn nofpexcept V_MUL_F32_e32 1069066811, %1204:vgpr_32, implicit $mode, implicit $exec1515 %1206:vgpr_32 = afn nofpexcept V_EXP_F32_e32 %1205:vgpr_32, implicit $mode, implicit $exec1516 %1207:vgpr_32 = contract nofpexcept V_FMA_F32_e64 0, %1.sub4:sgpr_512, 0, %679.sub1:vreg_512_align2, 1, %151:vgpr_32, 0, 0, implicit $mode, implicit $exec1517 %1208:vgpr_32 = afn nofpexcept V_MUL_F32_e32 1069066811, %1207:vgpr_32, implicit $mode, implicit $exec1518 %1209:vgpr_32 = afn nofpexcept V_EXP_F32_e32 %1208:vgpr_32, implicit $mode, implicit $exec1519 %1210:vgpr_32 = contract nofpexcept V_FMA_F32_e64 0, %1.sub4:sgpr_512, 0, %679.sub2:vreg_512_align2, 1, %151:vgpr_32, 0, 0, implicit $mode, implicit $exec1520 %1211:vgpr_32 = afn nofpexcept V_MUL_F32_e32 1069066811, %1210:vgpr_32, implicit $mode, implicit $exec1521 %1212:vgpr_32 = afn nofpexcept V_EXP_F32_e32 %1211:vgpr_32, implicit $mode, implicit $exec1522 %1213:vgpr_32 = contract nofpexcept V_FMA_F32_e64 0, %1.sub4:sgpr_512, 0, %679.sub3:vreg_512_align2, 1, %151:vgpr_32, 0, 0, implicit $mode, implicit $exec1523 %1214:vgpr_32 = afn nofpexcept V_MUL_F32_e32 1069066811, %1213:vgpr_32, implicit $mode, implicit $exec1524 %1215:vgpr_32 = afn nofpexcept V_EXP_F32_e32 %1214:vgpr_32, implicit $mode, implicit $exec1525 %1216:vgpr_32 = contract nofpexcept V_FMA_F32_e64 0, %1.sub4:sgpr_512, 0, %679.sub4:vreg_512_align2, 1, %151:vgpr_32, 0, 0, implicit $mode, implicit $exec1526 %1217:vgpr_32 = afn nofpexcept V_MUL_F32_e32 1069066811, %1216:vgpr_32, implicit $mode, implicit $exec1527 %1218:vgpr_32 = afn nofpexcept V_EXP_F32_e32 %1217:vgpr_32, implicit $mode, implicit $exec1528 %1219:vgpr_32 = contract nofpexcept V_FMA_F32_e64 0, %1.sub4:sgpr_512, 0, %679.sub5:vreg_512_align2, 1, %151:vgpr_32, 0, 0, implicit $mode, implicit $exec1529 %1220:vgpr_32 = afn nofpexcept V_MUL_F32_e32 1069066811, %1219:vgpr_32, implicit $mode, implicit $exec1530 %1221:vgpr_32 = afn nofpexcept V_EXP_F32_e32 %1220:vgpr_32, implicit $mode, implicit $exec1531 %1222:vgpr_32 = contract nofpexcept V_FMA_F32_e64 0, %1.sub4:sgpr_512, 0, %679.sub6:vreg_512_align2, 1, %151:vgpr_32, 0, 0, implicit $mode, implicit $exec1532 %1223:vgpr_32 = afn nofpexcept V_MUL_F32_e32 1069066811, %1222:vgpr_32, implicit $mode, implicit $exec1533 %1224:vgpr_32 = afn nofpexcept V_EXP_F32_e32 %1223:vgpr_32, implicit $mode, implicit $exec1534 %1225:vgpr_32 = contract nofpexcept V_FMA_F32_e64 0, %1.sub4:sgpr_512, 0, %679.sub7:vreg_512_align2, 1, %151:vgpr_32, 0, 0, implicit $mode, implicit $exec1535 %1226:vgpr_32 = afn nofpexcept V_MUL_F32_e32 1069066811, %1225:vgpr_32, implicit $mode, implicit $exec1536 %1227:vgpr_32 = afn nofpexcept V_EXP_F32_e32 %1226:vgpr_32, implicit $mode, implicit $exec1537 %1228:vgpr_32 = contract nofpexcept V_FMA_F32_e64 0, %1.sub4:sgpr_512, 0, %679.sub8:vreg_512_align2, 1, %151:vgpr_32, 0, 0, implicit $mode, implicit $exec1538 %1229:vgpr_32 = afn nofpexcept V_MUL_F32_e32 1069066811, %1228:vgpr_32, implicit $mode, implicit $exec1539 %1230:vgpr_32 = afn nofpexcept V_EXP_F32_e32 %1229:vgpr_32, implicit $mode, implicit $exec1540 %1231:vgpr_32 = contract nofpexcept V_FMA_F32_e64 0, %1.sub4:sgpr_512, 0, %679.sub9:vreg_512_align2, 1, %151:vgpr_32, 0, 0, implicit $mode, implicit $exec1541 %1232:vgpr_32 = afn nofpexcept V_MUL_F32_e32 1069066811, %1231:vgpr_32, implicit $mode, implicit $exec1542 %1233:vgpr_32 = afn nofpexcept V_EXP_F32_e32 %1232:vgpr_32, implicit $mode, implicit $exec1543 %1234:vgpr_32 = contract nofpexcept V_FMA_F32_e64 0, %1.sub4:sgpr_512, 0, %679.sub10:vreg_512_align2, 1, %151:vgpr_32, 0, 0, implicit $mode, implicit $exec1544 %1235:vgpr_32 = afn nofpexcept V_MUL_F32_e32 1069066811, %1234:vgpr_32, implicit $mode, implicit $exec1545 %1236:vgpr_32 = afn nofpexcept V_EXP_F32_e32 %1235:vgpr_32, implicit $mode, implicit $exec1546 %1237:vgpr_32 = contract nofpexcept V_FMA_F32_e64 0, %1.sub4:sgpr_512, 0, %679.sub11:vreg_512_align2, 1, %151:vgpr_32, 0, 0, implicit $mode, implicit $exec1547 %1238:vgpr_32 = afn nofpexcept V_MUL_F32_e32 1069066811, %1237:vgpr_32, implicit $mode, implicit $exec1548 %1239:vgpr_32 = afn nofpexcept V_EXP_F32_e32 %1238:vgpr_32, implicit $mode, implicit $exec1549 %1240:vgpr_32 = contract nofpexcept V_FMA_F32_e64 0, %1.sub4:sgpr_512, 0, %679.sub12:vreg_512_align2, 1, %151:vgpr_32, 0, 0, implicit $mode, implicit $exec1550 %1241:vgpr_32 = afn nofpexcept V_MUL_F32_e32 1069066811, %1240:vgpr_32, implicit $mode, implicit $exec1551 %1242:vgpr_32 = afn nofpexcept V_EXP_F32_e32 %1241:vgpr_32, implicit $mode, implicit $exec1552 %1243:vgpr_32 = contract nofpexcept V_FMA_F32_e64 0, %1.sub4:sgpr_512, 0, %679.sub13:vreg_512_align2, 1, %151:vgpr_32, 0, 0, implicit $mode, implicit $exec1553 %1244:vgpr_32 = afn nofpexcept V_MUL_F32_e32 1069066811, %1243:vgpr_32, implicit $mode, implicit $exec1554 %1245:vgpr_32 = afn nofpexcept V_EXP_F32_e32 %1244:vgpr_32, implicit $mode, implicit $exec1555 %1246:vgpr_32 = contract nofpexcept V_FMA_F32_e64 0, %1.sub4:sgpr_512, 0, %679.sub14:vreg_512_align2, 1, %151:vgpr_32, 0, 0, implicit $mode, implicit $exec1556 %1247:vgpr_32 = afn nofpexcept V_MUL_F32_e32 1069066811, %1246:vgpr_32, implicit $mode, implicit $exec1557 %1248:vgpr_32 = afn nofpexcept V_EXP_F32_e32 %1247:vgpr_32, implicit $mode, implicit $exec1558 %1249:vgpr_32 = contract nofpexcept V_FMA_F32_e64 0, %1.sub4:sgpr_512, 0, %679.sub15:vreg_512_align2, 1, %151:vgpr_32, 0, 0, implicit $mode, implicit $exec1559 %1250:vgpr_32 = afn nofpexcept V_MUL_F32_e32 1069066811, %1249:vgpr_32, implicit $mode, implicit $exec1560 %1251:vgpr_32 = afn nofpexcept V_EXP_F32_e32 %1250:vgpr_32, implicit $mode, implicit $exec1561 %1252:vgpr_32 = contract nofpexcept V_FMA_F32_e64 0, %1.sub4:sgpr_512, 0, %690.sub0:vreg_512_align2, 1, %151:vgpr_32, 0, 0, implicit $mode, implicit $exec1562 %1253:vgpr_32 = afn nofpexcept V_MUL_F32_e32 1069066811, %1252:vgpr_32, implicit $mode, implicit $exec1563 %1254:vgpr_32 = afn nofpexcept V_EXP_F32_e32 %1253:vgpr_32, implicit $mode, implicit $exec1564 %1255:vgpr_32 = contract nofpexcept V_FMA_F32_e64 0, %1.sub4:sgpr_512, 0, %690.sub1:vreg_512_align2, 1, %151:vgpr_32, 0, 0, implicit $mode, implicit $exec1565 %1256:vgpr_32 = afn nofpexcept V_MUL_F32_e32 1069066811, %1255:vgpr_32, implicit $mode, implicit $exec1566 %1257:vgpr_32 = afn nofpexcept V_EXP_F32_e32 %1256:vgpr_32, implicit $mode, implicit $exec1567 %1258:vgpr_32 = contract nofpexcept V_FMA_F32_e64 0, %1.sub4:sgpr_512, 0, %690.sub2:vreg_512_align2, 1, %151:vgpr_32, 0, 0, implicit $mode, implicit $exec1568 %1259:vgpr_32 = afn nofpexcept V_MUL_F32_e32 1069066811, %1258:vgpr_32, implicit $mode, implicit $exec1569 %1260:vgpr_32 = afn nofpexcept V_EXP_F32_e32 %1259:vgpr_32, implicit $mode, implicit $exec1570 %1261:vgpr_32 = contract nofpexcept V_FMA_F32_e64 0, %1.sub4:sgpr_512, 0, %690.sub3:vreg_512_align2, 1, %151:vgpr_32, 0, 0, implicit $mode, implicit $exec1571 %1262:vgpr_32 = afn nofpexcept V_MUL_F32_e32 1069066811, %1261:vgpr_32, implicit $mode, implicit $exec1572 %1263:vgpr_32 = afn nofpexcept V_EXP_F32_e32 %1262:vgpr_32, implicit $mode, implicit $exec1573 %1264:vgpr_32 = contract nofpexcept V_FMA_F32_e64 0, %1.sub4:sgpr_512, 0, %690.sub4:vreg_512_align2, 1, %151:vgpr_32, 0, 0, implicit $mode, implicit $exec1574 %1265:vgpr_32 = afn nofpexcept V_MUL_F32_e32 1069066811, %1264:vgpr_32, implicit $mode, implicit $exec1575 %1266:vgpr_32 = afn nofpexcept V_EXP_F32_e32 %1265:vgpr_32, implicit $mode, implicit $exec1576 %1267:vgpr_32 = contract nofpexcept V_FMA_F32_e64 0, %1.sub4:sgpr_512, 0, %690.sub5:vreg_512_align2, 1, %151:vgpr_32, 0, 0, implicit $mode, implicit $exec1577 %1268:vgpr_32 = afn nofpexcept V_MUL_F32_e32 1069066811, %1267:vgpr_32, implicit $mode, implicit $exec1578 %1269:vgpr_32 = afn nofpexcept V_EXP_F32_e32 %1268:vgpr_32, implicit $mode, implicit $exec1579 %1270:vgpr_32 = contract nofpexcept V_FMA_F32_e64 0, %1.sub4:sgpr_512, 0, %690.sub6:vreg_512_align2, 1, %151:vgpr_32, 0, 0, implicit $mode, implicit $exec1580 %1271:vgpr_32 = afn nofpexcept V_MUL_F32_e32 1069066811, %1270:vgpr_32, implicit $mode, implicit $exec1581 %1272:vgpr_32 = afn nofpexcept V_EXP_F32_e32 %1271:vgpr_32, implicit $mode, implicit $exec1582 %1273:vgpr_32 = contract nofpexcept V_FMA_F32_e64 0, %1.sub4:sgpr_512, 0, %690.sub7:vreg_512_align2, 1, %151:vgpr_32, 0, 0, implicit $mode, implicit $exec1583 %1274:vgpr_32 = afn nofpexcept V_MUL_F32_e32 1069066811, %1273:vgpr_32, implicit $mode, implicit $exec1584 %1275:vgpr_32 = afn nofpexcept V_EXP_F32_e32 %1274:vgpr_32, implicit $mode, implicit $exec1585 %1276:vgpr_32 = contract nofpexcept V_FMA_F32_e64 0, %1.sub4:sgpr_512, 0, %690.sub8:vreg_512_align2, 1, %151:vgpr_32, 0, 0, implicit $mode, implicit $exec1586 %1277:vgpr_32 = afn nofpexcept V_MUL_F32_e32 1069066811, %1276:vgpr_32, implicit $mode, implicit $exec1587 %1278:vgpr_32 = afn nofpexcept V_EXP_F32_e32 %1277:vgpr_32, implicit $mode, implicit $exec1588 %1279:vgpr_32 = contract nofpexcept V_FMA_F32_e64 0, %1.sub4:sgpr_512, 0, %690.sub9:vreg_512_align2, 1, %151:vgpr_32, 0, 0, implicit $mode, implicit $exec1589 %1280:vgpr_32 = afn nofpexcept V_MUL_F32_e32 1069066811, %1279:vgpr_32, implicit $mode, implicit $exec1590 %1281:vgpr_32 = afn nofpexcept V_EXP_F32_e32 %1280:vgpr_32, implicit $mode, implicit $exec1591 %1282:vgpr_32 = contract nofpexcept V_FMA_F32_e64 0, %1.sub4:sgpr_512, 0, %690.sub10:vreg_512_align2, 1, %151:vgpr_32, 0, 0, implicit $mode, implicit $exec1592 %1283:vgpr_32 = afn nofpexcept V_MUL_F32_e32 1069066811, %1282:vgpr_32, implicit $mode, implicit $exec1593 %1284:vgpr_32 = afn nofpexcept V_EXP_F32_e32 %1283:vgpr_32, implicit $mode, implicit $exec1594 %1285:vgpr_32 = contract nofpexcept V_FMA_F32_e64 0, %1.sub4:sgpr_512, 0, %690.sub11:vreg_512_align2, 1, %151:vgpr_32, 0, 0, implicit $mode, implicit $exec1595 %1286:vgpr_32 = afn nofpexcept V_MUL_F32_e32 1069066811, %1285:vgpr_32, implicit $mode, implicit $exec1596 %1287:vgpr_32 = afn nofpexcept V_EXP_F32_e32 %1286:vgpr_32, implicit $mode, implicit $exec1597 %1288:vgpr_32 = contract nofpexcept V_FMA_F32_e64 0, %1.sub4:sgpr_512, 0, %690.sub12:vreg_512_align2, 1, %151:vgpr_32, 0, 0, implicit $mode, implicit $exec1598 %1289:vgpr_32 = afn nofpexcept V_MUL_F32_e32 1069066811, %1288:vgpr_32, implicit $mode, implicit $exec1599 %1290:vgpr_32 = afn nofpexcept V_EXP_F32_e32 %1289:vgpr_32, implicit $mode, implicit $exec1600 %1291:vgpr_32 = contract nofpexcept V_FMA_F32_e64 0, %1.sub4:sgpr_512, 0, %690.sub13:vreg_512_align2, 1, %151:vgpr_32, 0, 0, implicit $mode, implicit $exec1601 %1292:vgpr_32 = afn nofpexcept V_MUL_F32_e32 1069066811, %1291:vgpr_32, implicit $mode, implicit $exec1602 %1293:vgpr_32 = afn nofpexcept V_EXP_F32_e32 %1292:vgpr_32, implicit $mode, implicit $exec1603 %1294:vgpr_32 = contract nofpexcept V_FMA_F32_e64 0, %1.sub4:sgpr_512, 0, %690.sub14:vreg_512_align2, 1, %151:vgpr_32, 0, 0, implicit $mode, implicit $exec1604 %1295:vgpr_32 = afn nofpexcept V_MUL_F32_e32 1069066811, %1294:vgpr_32, implicit $mode, implicit $exec1605 %1296:vgpr_32 = afn nofpexcept V_EXP_F32_e32 %1295:vgpr_32, implicit $mode, implicit $exec1606 %1297:vgpr_32 = contract nofpexcept V_FMA_F32_e64 0, %1.sub4:sgpr_512, 0, %690.sub15:vreg_512_align2, 1, %151:vgpr_32, 0, 0, implicit $mode, implicit $exec1607 %1298:vgpr_32 = afn nofpexcept V_MUL_F32_e32 1069066811, %1297:vgpr_32, implicit $mode, implicit $exec1608 %1299:vgpr_32 = afn nofpexcept V_EXP_F32_e32 %1298:vgpr_32, implicit $mode, implicit $exec1609 %1300:vgpr_32 = contract nofpexcept V_FMA_F32_e64 0, %1.sub4:sgpr_512, 0, %701.sub0:vreg_512_align2, 1, %151:vgpr_32, 0, 0, implicit $mode, implicit $exec1610 %1301:vgpr_32 = afn nofpexcept V_MUL_F32_e32 1069066811, %1300:vgpr_32, implicit $mode, implicit $exec1611 %1302:vgpr_32 = afn nofpexcept V_EXP_F32_e32 %1301:vgpr_32, implicit $mode, implicit $exec1612 %1303:vgpr_32 = contract nofpexcept V_FMA_F32_e64 0, %1.sub4:sgpr_512, 0, %701.sub1:vreg_512_align2, 1, %151:vgpr_32, 0, 0, implicit $mode, implicit $exec1613 %1304:vgpr_32 = afn nofpexcept V_MUL_F32_e32 1069066811, %1303:vgpr_32, implicit $mode, implicit $exec1614 %1305:vgpr_32 = afn nofpexcept V_EXP_F32_e32 %1304:vgpr_32, implicit $mode, implicit $exec1615 %1306:vgpr_32 = contract nofpexcept V_FMA_F32_e64 0, %1.sub4:sgpr_512, 0, %701.sub2:vreg_512_align2, 1, %151:vgpr_32, 0, 0, implicit $mode, implicit $exec1616 %1307:vgpr_32 = afn nofpexcept V_MUL_F32_e32 1069066811, %1306:vgpr_32, implicit $mode, implicit $exec1617 %1308:vgpr_32 = afn nofpexcept V_EXP_F32_e32 %1307:vgpr_32, implicit $mode, implicit $exec1618 %1309:vgpr_32 = contract nofpexcept V_FMA_F32_e64 0, %1.sub4:sgpr_512, 0, %701.sub3:vreg_512_align2, 1, %151:vgpr_32, 0, 0, implicit $mode, implicit $exec1619 %1310:vgpr_32 = afn nofpexcept V_MUL_F32_e32 1069066811, %1309:vgpr_32, implicit $mode, implicit $exec1620 %1311:vgpr_32 = afn nofpexcept V_EXP_F32_e32 %1310:vgpr_32, implicit $mode, implicit $exec1621 %1312:vgpr_32 = contract nofpexcept V_FMA_F32_e64 0, %1.sub4:sgpr_512, 0, %701.sub4:vreg_512_align2, 1, %151:vgpr_32, 0, 0, implicit $mode, implicit $exec1622 %1313:vgpr_32 = afn nofpexcept V_MUL_F32_e32 1069066811, %1312:vgpr_32, implicit $mode, implicit $exec1623 %1314:vgpr_32 = afn nofpexcept V_EXP_F32_e32 %1313:vgpr_32, implicit $mode, implicit $exec1624 %1315:vgpr_32 = contract nofpexcept V_FMA_F32_e64 0, %1.sub4:sgpr_512, 0, %701.sub5:vreg_512_align2, 1, %151:vgpr_32, 0, 0, implicit $mode, implicit $exec1625 %1316:vgpr_32 = afn nofpexcept V_MUL_F32_e32 1069066811, %1315:vgpr_32, implicit $mode, implicit $exec1626 %1317:vgpr_32 = afn nofpexcept V_EXP_F32_e32 %1316:vgpr_32, implicit $mode, implicit $exec1627 %1318:vgpr_32 = contract nofpexcept V_FMA_F32_e64 0, %1.sub4:sgpr_512, 0, %701.sub6:vreg_512_align2, 1, %151:vgpr_32, 0, 0, implicit $mode, implicit $exec1628 %1319:vgpr_32 = afn nofpexcept V_MUL_F32_e32 1069066811, %1318:vgpr_32, implicit $mode, implicit $exec1629 %1320:vgpr_32 = afn nofpexcept V_EXP_F32_e32 %1319:vgpr_32, implicit $mode, implicit $exec1630 %1321:vgpr_32 = contract nofpexcept V_FMA_F32_e64 0, %1.sub4:sgpr_512, 0, %701.sub7:vreg_512_align2, 1, %151:vgpr_32, 0, 0, implicit $mode, implicit $exec1631 %1322:vgpr_32 = afn nofpexcept V_MUL_F32_e32 1069066811, %1321:vgpr_32, implicit $mode, implicit $exec1632 %1323:vgpr_32 = afn nofpexcept V_EXP_F32_e32 %1322:vgpr_32, implicit $mode, implicit $exec1633 %1324:vgpr_32 = contract nofpexcept V_FMA_F32_e64 0, %1.sub4:sgpr_512, 0, %701.sub8:vreg_512_align2, 1, %151:vgpr_32, 0, 0, implicit $mode, implicit $exec1634 %1325:vgpr_32 = afn nofpexcept V_MUL_F32_e32 1069066811, %1324:vgpr_32, implicit $mode, implicit $exec1635 %1326:vgpr_32 = afn nofpexcept V_EXP_F32_e32 %1325:vgpr_32, implicit $mode, implicit $exec1636 %1327:vgpr_32 = contract nofpexcept V_FMA_F32_e64 0, %1.sub4:sgpr_512, 0, %701.sub9:vreg_512_align2, 1, %151:vgpr_32, 0, 0, implicit $mode, implicit $exec1637 %1328:vgpr_32 = afn nofpexcept V_MUL_F32_e32 1069066811, %1327:vgpr_32, implicit $mode, implicit $exec1638 %1329:vgpr_32 = afn nofpexcept V_EXP_F32_e32 %1328:vgpr_32, implicit $mode, implicit $exec1639 %1330:vgpr_32 = contract nofpexcept V_FMA_F32_e64 0, %1.sub4:sgpr_512, 0, %701.sub10:vreg_512_align2, 1, %151:vgpr_32, 0, 0, implicit $mode, implicit $exec1640 %1331:vgpr_32 = afn nofpexcept V_MUL_F32_e32 1069066811, %1330:vgpr_32, implicit $mode, implicit $exec1641 %1332:vgpr_32 = afn nofpexcept V_EXP_F32_e32 %1331:vgpr_32, implicit $mode, implicit $exec1642 %1333:vgpr_32 = contract nofpexcept V_FMA_F32_e64 0, %1.sub4:sgpr_512, 0, %701.sub11:vreg_512_align2, 1, %151:vgpr_32, 0, 0, implicit $mode, implicit $exec1643 %1334:vgpr_32 = afn nofpexcept V_MUL_F32_e32 1069066811, %1333:vgpr_32, implicit $mode, implicit $exec1644 %1335:vgpr_32 = afn nofpexcept V_EXP_F32_e32 %1334:vgpr_32, implicit $mode, implicit $exec1645 %1336:vgpr_32 = contract nofpexcept V_FMA_F32_e64 0, %1.sub4:sgpr_512, 0, %701.sub12:vreg_512_align2, 1, %151:vgpr_32, 0, 0, implicit $mode, implicit $exec1646 %1337:vgpr_32 = afn nofpexcept V_MUL_F32_e32 1069066811, %1336:vgpr_32, implicit $mode, implicit $exec1647 %1338:vgpr_32 = afn nofpexcept V_EXP_F32_e32 %1337:vgpr_32, implicit $mode, implicit $exec1648 %1339:vgpr_32 = contract nofpexcept V_FMA_F32_e64 0, %1.sub4:sgpr_512, 0, %701.sub13:vreg_512_align2, 1, %151:vgpr_32, 0, 0, implicit $mode, implicit $exec1649 %1340:vgpr_32 = afn nofpexcept V_MUL_F32_e32 1069066811, %1339:vgpr_32, implicit $mode, implicit $exec1650 %1341:vgpr_32 = afn nofpexcept V_EXP_F32_e32 %1340:vgpr_32, implicit $mode, implicit $exec1651 %1342:vgpr_32 = contract nofpexcept V_FMA_F32_e64 0, %1.sub4:sgpr_512, 0, %701.sub14:vreg_512_align2, 1, %151:vgpr_32, 0, 0, implicit $mode, implicit $exec1652 %1343:vgpr_32 = afn nofpexcept V_MUL_F32_e32 1069066811, %1342:vgpr_32, implicit $mode, implicit $exec1653 %1344:vgpr_32 = afn nofpexcept V_EXP_F32_e32 %1343:vgpr_32, implicit $mode, implicit $exec1654 %1345:vgpr_32 = contract nofpexcept V_FMA_F32_e64 0, %1.sub4:sgpr_512, 0, %701.sub15:vreg_512_align2, 1, %151:vgpr_32, 0, 0, implicit $mode, implicit $exec1655 %1346:vgpr_32 = afn nofpexcept V_MUL_F32_e32 1069066811, %1345:vgpr_32, implicit $mode, implicit $exec1656 %1347:vgpr_32 = afn nofpexcept V_EXP_F32_e32 %1346:vgpr_32, implicit $mode, implicit $exec1657 %1348:vgpr_32 = contract nofpexcept V_ADD_F32_e32 0, %1158:vgpr_32, implicit $mode, implicit $exec1658 %1349:vgpr_32 = contract nofpexcept V_ADD_F32_e32 %1161:vgpr_32, %1348:vgpr_32, implicit $mode, implicit $exec1659 %1350:vgpr_32 = contract nofpexcept V_ADD_F32_e32 %1164:vgpr_32, %1349:vgpr_32, implicit $mode, implicit $exec1660 %1351:vgpr_32 = contract nofpexcept V_ADD_F32_e32 %1167:vgpr_32, %1350:vgpr_32, implicit $mode, implicit $exec1661 %1352:vgpr_32 = contract nofpexcept V_ADD_F32_e32 %1170:vgpr_32, %1351:vgpr_32, implicit $mode, implicit $exec1662 %1353:vgpr_32 = contract nofpexcept V_ADD_F32_e32 %1173:vgpr_32, %1352:vgpr_32, implicit $mode, implicit $exec1663 %1354:vgpr_32 = contract nofpexcept V_ADD_F32_e32 %1176:vgpr_32, %1353:vgpr_32, implicit $mode, implicit $exec1664 %1355:vgpr_32 = contract nofpexcept V_ADD_F32_e32 %1179:vgpr_32, %1354:vgpr_32, implicit $mode, implicit $exec1665 %1356:vgpr_32 = contract nofpexcept V_ADD_F32_e32 %1182:vgpr_32, %1355:vgpr_32, implicit $mode, implicit $exec1666 %1357:vgpr_32 = contract nofpexcept V_ADD_F32_e32 %1185:vgpr_32, %1356:vgpr_32, implicit $mode, implicit $exec1667 %1358:vgpr_32 = contract nofpexcept V_ADD_F32_e32 %1188:vgpr_32, %1357:vgpr_32, implicit $mode, implicit $exec1668 %1359:vgpr_32 = contract nofpexcept V_ADD_F32_e32 %1191:vgpr_32, %1358:vgpr_32, implicit $mode, implicit $exec1669 %1360:vgpr_32 = contract nofpexcept V_ADD_F32_e32 %1194:vgpr_32, %1359:vgpr_32, implicit $mode, implicit $exec1670 %1361:vgpr_32 = contract nofpexcept V_ADD_F32_e32 %1197:vgpr_32, %1360:vgpr_32, implicit $mode, implicit $exec1671 %1362:vgpr_32 = contract nofpexcept V_ADD_F32_e32 %1200:vgpr_32, %1361:vgpr_32, implicit $mode, implicit $exec1672 %1363:vgpr_32 = contract nofpexcept V_ADD_F32_e32 %1203:vgpr_32, %1362:vgpr_32, implicit $mode, implicit $exec1673 %1364:vgpr_32 = contract nofpexcept V_ADD_F32_e32 %1206:vgpr_32, %1363:vgpr_32, implicit $mode, implicit $exec1674 %1365:vgpr_32 = contract nofpexcept V_ADD_F32_e32 %1209:vgpr_32, %1364:vgpr_32, implicit $mode, implicit $exec1675 %1366:vgpr_32 = contract nofpexcept V_ADD_F32_e32 %1212:vgpr_32, %1365:vgpr_32, implicit $mode, implicit $exec1676 %1367:vgpr_32 = contract nofpexcept V_ADD_F32_e32 %1215:vgpr_32, %1366:vgpr_32, implicit $mode, implicit $exec1677 %1368:vgpr_32 = contract nofpexcept V_ADD_F32_e32 %1218:vgpr_32, %1367:vgpr_32, implicit $mode, implicit $exec1678 %1369:vgpr_32 = contract nofpexcept V_ADD_F32_e32 %1221:vgpr_32, %1368:vgpr_32, implicit $mode, implicit $exec1679 %1370:vgpr_32 = contract nofpexcept V_ADD_F32_e32 %1224:vgpr_32, %1369:vgpr_32, implicit $mode, implicit $exec1680 %1371:vgpr_32 = contract nofpexcept V_ADD_F32_e32 %1227:vgpr_32, %1370:vgpr_32, implicit $mode, implicit $exec1681 %1372:vgpr_32 = contract nofpexcept V_ADD_F32_e32 %1230:vgpr_32, %1371:vgpr_32, implicit $mode, implicit $exec1682 %1373:vgpr_32 = contract nofpexcept V_ADD_F32_e32 %1233:vgpr_32, %1372:vgpr_32, implicit $mode, implicit $exec1683 %1374:vgpr_32 = contract nofpexcept V_ADD_F32_e32 %1236:vgpr_32, %1373:vgpr_32, implicit $mode, implicit $exec1684 %1375:vgpr_32 = contract nofpexcept V_ADD_F32_e32 %1239:vgpr_32, %1374:vgpr_32, implicit $mode, implicit $exec1685 %1376:vgpr_32 = contract nofpexcept V_ADD_F32_e32 %1242:vgpr_32, %1375:vgpr_32, implicit $mode, implicit $exec1686 %1377:vgpr_32 = contract nofpexcept V_ADD_F32_e32 %1245:vgpr_32, %1376:vgpr_32, implicit $mode, implicit $exec1687 %1378:vgpr_32 = contract nofpexcept V_ADD_F32_e32 %1248:vgpr_32, %1377:vgpr_32, implicit $mode, implicit $exec1688 %1379:vgpr_32 = contract nofpexcept V_ADD_F32_e32 %1251:vgpr_32, %1378:vgpr_32, implicit $mode, implicit $exec1689 %1380:vgpr_32 = contract nofpexcept V_ADD_F32_e32 %1254:vgpr_32, %1379:vgpr_32, implicit $mode, implicit $exec1690 %1381:vgpr_32 = contract nofpexcept V_ADD_F32_e32 %1257:vgpr_32, %1380:vgpr_32, implicit $mode, implicit $exec1691 %1382:vgpr_32 = contract nofpexcept V_ADD_F32_e32 %1260:vgpr_32, %1381:vgpr_32, implicit $mode, implicit $exec1692 %1383:vgpr_32 = contract nofpexcept V_ADD_F32_e32 %1263:vgpr_32, %1382:vgpr_32, implicit $mode, implicit $exec1693 %1384:vgpr_32 = contract nofpexcept V_ADD_F32_e32 %1266:vgpr_32, %1383:vgpr_32, implicit $mode, implicit $exec1694 %1385:vgpr_32 = contract nofpexcept V_ADD_F32_e32 %1269:vgpr_32, %1384:vgpr_32, implicit $mode, implicit $exec1695 %1386:vgpr_32 = contract nofpexcept V_ADD_F32_e32 %1272:vgpr_32, %1385:vgpr_32, implicit $mode, implicit $exec1696 %1387:vgpr_32 = contract nofpexcept V_ADD_F32_e32 %1275:vgpr_32, %1386:vgpr_32, implicit $mode, implicit $exec1697 %1388:vgpr_32 = contract nofpexcept V_ADD_F32_e32 %1278:vgpr_32, %1387:vgpr_32, implicit $mode, implicit $exec1698 %1389:vgpr_32 = contract nofpexcept V_ADD_F32_e32 %1281:vgpr_32, %1388:vgpr_32, implicit $mode, implicit $exec1699 %1390:vgpr_32 = contract nofpexcept V_ADD_F32_e32 %1284:vgpr_32, %1389:vgpr_32, implicit $mode, implicit $exec1700 %1391:vgpr_32 = contract nofpexcept V_ADD_F32_e32 %1287:vgpr_32, %1390:vgpr_32, implicit $mode, implicit $exec1701 %1392:vgpr_32 = contract nofpexcept V_ADD_F32_e32 %1290:vgpr_32, %1391:vgpr_32, implicit $mode, implicit $exec1702 %1393:vgpr_32 = contract nofpexcept V_ADD_F32_e32 %1293:vgpr_32, %1392:vgpr_32, implicit $mode, implicit $exec1703 %1394:vgpr_32 = contract nofpexcept V_ADD_F32_e32 %1296:vgpr_32, %1393:vgpr_32, implicit $mode, implicit $exec1704 %1395:vgpr_32 = contract nofpexcept V_ADD_F32_e32 %1299:vgpr_32, %1394:vgpr_32, implicit $mode, implicit $exec1705 %1396:vgpr_32 = contract nofpexcept V_ADD_F32_e32 %1302:vgpr_32, %1395:vgpr_32, implicit $mode, implicit $exec1706 %1397:vgpr_32 = contract nofpexcept V_ADD_F32_e32 %1305:vgpr_32, %1396:vgpr_32, implicit $mode, implicit $exec1707 %1398:vgpr_32 = contract nofpexcept V_ADD_F32_e32 %1308:vgpr_32, %1397:vgpr_32, implicit $mode, implicit $exec1708 %1399:vgpr_32 = contract nofpexcept V_ADD_F32_e32 %1311:vgpr_32, %1398:vgpr_32, implicit $mode, implicit $exec1709 %1400:vgpr_32 = contract nofpexcept V_ADD_F32_e32 %1314:vgpr_32, %1399:vgpr_32, implicit $mode, implicit $exec1710 %1401:vgpr_32 = contract nofpexcept V_ADD_F32_e32 %1317:vgpr_32, %1400:vgpr_32, implicit $mode, implicit $exec1711 %1402:vgpr_32 = contract nofpexcept V_ADD_F32_e32 %1320:vgpr_32, %1401:vgpr_32, implicit $mode, implicit $exec1712 %1403:vgpr_32 = contract nofpexcept V_ADD_F32_e32 %1323:vgpr_32, %1402:vgpr_32, implicit $mode, implicit $exec1713 %1404:vgpr_32 = contract nofpexcept V_ADD_F32_e32 %1326:vgpr_32, %1403:vgpr_32, implicit $mode, implicit $exec1714 %1405:vgpr_32 = contract nofpexcept V_ADD_F32_e32 %1329:vgpr_32, %1404:vgpr_32, implicit $mode, implicit $exec1715 %1406:vgpr_32 = contract nofpexcept V_ADD_F32_e32 %1332:vgpr_32, %1405:vgpr_32, implicit $mode, implicit $exec1716 %1407:vgpr_32 = contract nofpexcept V_ADD_F32_e32 %1335:vgpr_32, %1406:vgpr_32, implicit $mode, implicit $exec1717 %1408:vgpr_32 = contract nofpexcept V_ADD_F32_e32 %1338:vgpr_32, %1407:vgpr_32, implicit $mode, implicit $exec1718 %1409:vgpr_32 = contract nofpexcept V_ADD_F32_e32 %1341:vgpr_32, %1408:vgpr_32, implicit $mode, implicit $exec1719 %1410:vgpr_32 = contract nofpexcept V_ADD_F32_e32 %1344:vgpr_32, %1409:vgpr_32, implicit $mode, implicit $exec1720 %1411:vgpr_32 = contract nofpexcept V_ADD_F32_e32 %1347:vgpr_32, %1410:vgpr_32, implicit $mode, implicit $exec1721 %1412:vgpr_32 = DS_BPERMUTE_B32 %39:vgpr_32, %1411:vgpr_32, 0, implicit $exec1722 %1413:vgpr_32 = contract nofpexcept V_ADD_F32_e32 %1411:vgpr_32, %1412:vgpr_32, implicit $mode, implicit $exec1723 %1414:vgpr_32 = DS_BPERMUTE_B32 %39:vgpr_32, %1413:vgpr_32, 0, implicit $exec1724 %3347:vgpr_32 = V_CNDMASK_B32_e64 0, %1414:vgpr_32, 0, %1413:vgpr_32, %484:sreg_64_xexec, implicit $exec1725 %1417:vgpr_32 = contract nofpexcept V_SUB_F32_e32 %3346:vgpr_32, %151:vgpr_32, implicit $mode, implicit $exec1726 %1418:vgpr_32 = afn nofpexcept V_MUL_F32_e32 1069066811, %1417:vgpr_32, implicit $mode, implicit $exec1727 undef %1455.sub0:vreg_64_align2 = afn nofpexcept V_EXP_F32_e32 %1418:vgpr_32, implicit $mode, implicit $exec1728 INLINEASM &"s_waitcnt vmcnt($0)", 57 /* sideeffect mayload maystore isconvergent attdialect */, 13 /* imm */, 8, !01729 undef %3037.sub0:vreg_64_align2 = V_PERM_B32_e64 %805.sub0:vreg_64_align2, %803.sub0:vreg_64_align2, %1422:sreg_32, implicit $exec1730 undef %3021.sub0:vreg_64_align2 = V_PERM_B32_e64 %805.sub0:vreg_64_align2, %803.sub0:vreg_64_align2, %1424:sreg_32, implicit $exec1731 %3037.sub1:vreg_64_align2 = V_PERM_B32_e64 %809.sub0:vreg_64_align2, %807.sub0:vreg_64_align2, %1422:sreg_32, implicit $exec1732 %3021.sub1:vreg_64_align2 = V_PERM_B32_e64 %809.sub0:vreg_64_align2, %807.sub0:vreg_64_align2, %1424:sreg_32, implicit $exec1733 undef %3005.sub0:vreg_64_align2 = V_PERM_B32_e64 %805.sub1:vreg_64_align2, %803.sub1:vreg_64_align2, %1422:sreg_32, implicit $exec1734 undef %2978.sub0:vreg_64_align2 = V_PERM_B32_e64 %805.sub1:vreg_64_align2, %803.sub1:vreg_64_align2, %1424:sreg_32, implicit $exec1735 %3005.sub1:vreg_64_align2 = V_PERM_B32_e64 %809.sub1:vreg_64_align2, %807.sub1:vreg_64_align2, %1422:sreg_32, implicit $exec1736 %2978.sub1:vreg_64_align2 = V_PERM_B32_e64 %809.sub1:vreg_64_align2, %807.sub1:vreg_64_align2, %1424:sreg_32, implicit $exec1737 %1442:vgpr_32 = V_ADD_U32_e32 %593:sreg_32_xm0, %15:vgpr_32, implicit $exec1738 %1444:vgpr_32 = V_AND_B32_e32 536870911, %1442:vgpr_32, implicit $exec1739 %1446:vgpr_32 = nsw V_MUL_LO_U32_e64 %1444:vgpr_32, %494:sreg_32, implicit $exec1740 %1447:vgpr_32 = V_ADD_LSHL_U32_e64 %47:vgpr_32, %1446:vgpr_32, 1, implicit $exec1741 DS_WRITE_B64_gfx9 %1447:vgpr_32, %3037:vreg_64_align2, 0, 0, implicit $exec1742 %1449:vgpr_32 = V_LSHL_ADD_U32_e64 %41:vgpr_32, 1, %1447:vgpr_32, implicit $exec1743 DS_WRITE_B64_gfx9 %1449:vgpr_32, %3021:vreg_64_align2, 0, 0, implicit $exec1744 %1451:vgpr_32 = V_LSHL_ADD_U32_e64 %42:vgpr_32, 1, %1449:vgpr_32, implicit $exec1745 DS_WRITE_B64_gfx9 %1451:vgpr_32, %3005:vreg_64_align2, 0, 0, implicit $exec1746 %1453:vgpr_32 = V_LSHL_ADD_U32_e64 %43:vgpr_32, 1, %1451:vgpr_32, implicit $exec1747 DS_WRITE_B64_gfx9 %1453:vgpr_32, %2978:vreg_64_align2, 0, 0, implicit $exec1748 %3347:vgpr_32 = contract nofpexcept V_FMAC_F32_e32 %86:vgpr_32, %1455.sub0:vreg_64_align2, %3347:vgpr_32, implicit $mode, implicit $exec1749 %2986.sub0_sub1:vreg_512_align2 = contract nofpexcept V_PK_MUL_F32 8, %2986.sub0_sub1:vreg_512_align2, 0, %1455:vreg_64_align2, 0, 0, 0, 0, 0, implicit $mode, implicit $exec1750 %2986.sub2_sub3:vreg_512_align2 = contract nofpexcept V_PK_MUL_F32 8, %2986.sub2_sub3:vreg_512_align2, 0, %1455:vreg_64_align2, 0, 0, 0, 0, 0, implicit $mode, implicit $exec1751 %2986.sub4_sub5:vreg_512_align2 = contract nofpexcept V_PK_MUL_F32 8, %2986.sub4_sub5:vreg_512_align2, 0, %1455:vreg_64_align2, 0, 0, 0, 0, 0, implicit $mode, implicit $exec1752 %2986.sub6_sub7:vreg_512_align2 = contract nofpexcept V_PK_MUL_F32 8, %2986.sub6_sub7:vreg_512_align2, 0, %1455:vreg_64_align2, 0, 0, 0, 0, 0, implicit $mode, implicit $exec1753 %2986.sub8_sub9:vreg_512_align2 = contract nofpexcept V_PK_MUL_F32 8, %2986.sub8_sub9:vreg_512_align2, 0, %1455:vreg_64_align2, 0, 0, 0, 0, 0, implicit $mode, implicit $exec1754 %2986.sub10_sub11:vreg_512_align2 = contract nofpexcept V_PK_MUL_F32 8, %2986.sub10_sub11:vreg_512_align2, 0, %1455:vreg_64_align2, 0, 0, 0, 0, 0, implicit $mode, implicit $exec1755 %2986.sub12_sub13:vreg_512_align2 = contract nofpexcept V_PK_MUL_F32 8, %2986.sub12_sub13:vreg_512_align2, 0, %1455:vreg_64_align2, 0, 0, 0, 0, 0, implicit $mode, implicit $exec1756 %2986.sub14_sub15:vreg_512_align2 = contract nofpexcept V_PK_MUL_F32 8, %2986.sub14_sub15:vreg_512_align2, 0, %1455:vreg_64_align2, 0, 0, 0, 0, 0, implicit $mode, implicit $exec1757 %3038.sub0_sub1:vreg_512_align2 = contract nofpexcept V_PK_MUL_F32 8, %3038.sub0_sub1:vreg_512_align2, 0, %1455:vreg_64_align2, 0, 0, 0, 0, 0, implicit $mode, implicit $exec1758 %3038.sub2_sub3:vreg_512_align2 = contract nofpexcept V_PK_MUL_F32 8, %3038.sub2_sub3:vreg_512_align2, 0, %1455:vreg_64_align2, 0, 0, 0, 0, 0, implicit $mode, implicit $exec1759 %3038.sub4_sub5:vreg_512_align2 = contract nofpexcept V_PK_MUL_F32 8, %3038.sub4_sub5:vreg_512_align2, 0, %1455:vreg_64_align2, 0, 0, 0, 0, 0, implicit $mode, implicit $exec1760 %3038.sub6_sub7:vreg_512_align2 = contract nofpexcept V_PK_MUL_F32 8, %3038.sub6_sub7:vreg_512_align2, 0, %1455:vreg_64_align2, 0, 0, 0, 0, 0, implicit $mode, implicit $exec1761 %3038.sub8_sub9:vreg_512_align2 = contract nofpexcept V_PK_MUL_F32 8, %3038.sub8_sub9:vreg_512_align2, 0, %1455:vreg_64_align2, 0, 0, 0, 0, 0, implicit $mode, implicit $exec1762 %3038.sub10_sub11:vreg_512_align2 = contract nofpexcept V_PK_MUL_F32 8, %3038.sub10_sub11:vreg_512_align2, 0, %1455:vreg_64_align2, 0, 0, 0, 0, 0, implicit $mode, implicit $exec1763 %3038.sub12_sub13:vreg_512_align2 = contract nofpexcept V_PK_MUL_F32 8, %3038.sub12_sub13:vreg_512_align2, 0, %1455:vreg_64_align2, 0, 0, 0, 0, 0, implicit $mode, implicit $exec1764 %3038.sub14_sub15:vreg_512_align2 = contract nofpexcept V_PK_MUL_F32 8, %3038.sub14_sub15:vreg_512_align2, 0, %1455:vreg_64_align2, 0, 0, 0, 0, 0, implicit $mode, implicit $exec1765 %2980.sub0_sub1:vreg_512_align2 = contract nofpexcept V_PK_MUL_F32 8, %2980.sub0_sub1:vreg_512_align2, 0, %1455:vreg_64_align2, 0, 0, 0, 0, 0, implicit $mode, implicit $exec1766 %2980.sub2_sub3:vreg_512_align2 = contract nofpexcept V_PK_MUL_F32 8, %2980.sub2_sub3:vreg_512_align2, 0, %1455:vreg_64_align2, 0, 0, 0, 0, 0, implicit $mode, implicit $exec1767 %2980.sub4_sub5:vreg_512_align2 = contract nofpexcept V_PK_MUL_F32 8, %2980.sub4_sub5:vreg_512_align2, 0, %1455:vreg_64_align2, 0, 0, 0, 0, 0, implicit $mode, implicit $exec1768 %2980.sub6_sub7:vreg_512_align2 = contract nofpexcept V_PK_MUL_F32 8, %2980.sub6_sub7:vreg_512_align2, 0, %1455:vreg_64_align2, 0, 0, 0, 0, 0, implicit $mode, implicit $exec1769 %2980.sub8_sub9:vreg_512_align2 = contract nofpexcept V_PK_MUL_F32 8, %2980.sub8_sub9:vreg_512_align2, 0, %1455:vreg_64_align2, 0, 0, 0, 0, 0, implicit $mode, implicit $exec1770 %2980.sub10_sub11:vreg_512_align2 = contract nofpexcept V_PK_MUL_F32 8, %2980.sub10_sub11:vreg_512_align2, 0, %1455:vreg_64_align2, 0, 0, 0, 0, 0, implicit $mode, implicit $exec1771 %2980.sub12_sub13:vreg_512_align2 = contract nofpexcept V_PK_MUL_F32 8, %2980.sub12_sub13:vreg_512_align2, 0, %1455:vreg_64_align2, 0, 0, 0, 0, 0, implicit $mode, implicit $exec1772 %2980.sub14_sub15:vreg_512_align2 = contract nofpexcept V_PK_MUL_F32 8, %2980.sub14_sub15:vreg_512_align2, 0, %1455:vreg_64_align2, 0, 0, 0, 0, 0, implicit $mode, implicit $exec1773 %3003.sub0_sub1:vreg_512_align2 = contract nofpexcept V_PK_MUL_F32 8, %3003.sub0_sub1:vreg_512_align2, 0, %1455:vreg_64_align2, 0, 0, 0, 0, 0, implicit $mode, implicit $exec1774 %3003.sub2_sub3:vreg_512_align2 = contract nofpexcept V_PK_MUL_F32 8, %3003.sub2_sub3:vreg_512_align2, 0, %1455:vreg_64_align2, 0, 0, 0, 0, 0, implicit $mode, implicit $exec1775 %3003.sub4_sub5:vreg_512_align2 = contract nofpexcept V_PK_MUL_F32 8, %3003.sub4_sub5:vreg_512_align2, 0, %1455:vreg_64_align2, 0, 0, 0, 0, 0, implicit $mode, implicit $exec1776 %3003.sub6_sub7:vreg_512_align2 = contract nofpexcept V_PK_MUL_F32 8, %3003.sub6_sub7:vreg_512_align2, 0, %1455:vreg_64_align2, 0, 0, 0, 0, 0, implicit $mode, implicit $exec1777 %3003.sub8_sub9:vreg_512_align2 = contract nofpexcept V_PK_MUL_F32 8, %3003.sub8_sub9:vreg_512_align2, 0, %1455:vreg_64_align2, 0, 0, 0, 0, 0, implicit $mode, implicit $exec1778 %3003.sub10_sub11:vreg_512_align2 = contract nofpexcept V_PK_MUL_F32 8, %3003.sub10_sub11:vreg_512_align2, 0, %1455:vreg_64_align2, 0, 0, 0, 0, 0, implicit $mode, implicit $exec1779 %3003.sub12_sub13:vreg_512_align2 = contract nofpexcept V_PK_MUL_F32 8, %3003.sub12_sub13:vreg_512_align2, 0, %1455:vreg_64_align2, 0, 0, 0, 0, 0, implicit $mode, implicit $exec1780 %3003.sub14_sub15:vreg_512_align2 = contract nofpexcept V_PK_MUL_F32 8, %3003.sub14_sub15:vreg_512_align2, 0, %1455:vreg_64_align2, 0, 0, 0, 0, 0, implicit $mode, implicit $exec1781 %1554:vgpr_32 = nofpexcept V_CVT_F16_F32_e32 %1158:vgpr_32, implicit $mode, implicit $exec1782 %1555:vgpr_32 = nofpexcept V_CVT_F16_F32_e32 %1161:vgpr_32, implicit $mode, implicit $exec1783 %1556:vgpr_32 = nofpexcept V_CVT_F16_F32_e32 %1164:vgpr_32, implicit $mode, implicit $exec1784 %1557:vgpr_32 = nofpexcept V_CVT_F16_F32_e32 %1170:vgpr_32, implicit $mode, implicit $exec1785 %1558:vgpr_32 = nofpexcept V_CVT_F16_F32_e32 %1173:vgpr_32, implicit $mode, implicit $exec1786 %1559:vgpr_32 = nofpexcept V_CVT_F16_F32_e32 %1176:vgpr_32, implicit $mode, implicit $exec1787 %1560:vgpr_32 = nofpexcept V_CVT_F16_F32_e32 %1182:vgpr_32, implicit $mode, implicit $exec1788 %1561:vgpr_32 = nofpexcept V_CVT_F16_F32_e32 %1185:vgpr_32, implicit $mode, implicit $exec1789 %1562:vgpr_32 = nofpexcept V_CVT_F16_F32_e32 %1188:vgpr_32, implicit $mode, implicit $exec1790 %1563:vgpr_32 = nofpexcept V_CVT_F16_F32_e32 %1194:vgpr_32, implicit $mode, implicit $exec1791 %1564:vgpr_32 = nofpexcept V_CVT_F16_F32_e32 %1197:vgpr_32, implicit $mode, implicit $exec1792 %1565:vgpr_32 = nofpexcept V_CVT_F16_F32_e32 %1200:vgpr_32, implicit $mode, implicit $exec1793 %1566:vgpr_32 = nofpexcept V_CVT_F16_F32_e32 %1206:vgpr_32, implicit $mode, implicit $exec1794 %1567:vgpr_32 = nofpexcept V_CVT_F16_F32_e32 %1209:vgpr_32, implicit $mode, implicit $exec1795 %1568:vgpr_32 = nofpexcept V_CVT_F16_F32_e32 %1212:vgpr_32, implicit $mode, implicit $exec1796 %1569:vgpr_32 = nofpexcept V_CVT_F16_F32_e32 %1218:vgpr_32, implicit $mode, implicit $exec1797 %1570:vgpr_32 = nofpexcept V_CVT_F16_F32_e32 %1221:vgpr_32, implicit $mode, implicit $exec1798 %1571:vgpr_32 = nofpexcept V_CVT_F16_F32_e32 %1224:vgpr_32, implicit $mode, implicit $exec1799 %1572:vgpr_32 = nofpexcept V_CVT_F16_F32_e32 %1230:vgpr_32, implicit $mode, implicit $exec1800 %1573:vgpr_32 = nofpexcept V_CVT_F16_F32_e32 %1233:vgpr_32, implicit $mode, implicit $exec1801 %1574:vgpr_32 = nofpexcept V_CVT_F16_F32_e32 %1236:vgpr_32, implicit $mode, implicit $exec1802 %1575:vgpr_32 = nofpexcept V_CVT_F16_F32_e32 %1242:vgpr_32, implicit $mode, implicit $exec1803 %1576:vgpr_32 = nofpexcept V_CVT_F16_F32_e32 %1245:vgpr_32, implicit $mode, implicit $exec1804 %1577:vgpr_32 = nofpexcept V_CVT_F16_F32_e32 %1248:vgpr_32, implicit $mode, implicit $exec1805 %1578:vgpr_32 = nofpexcept V_CVT_F16_F32_e32 %1254:vgpr_32, implicit $mode, implicit $exec1806 %1579:vgpr_32 = nofpexcept V_CVT_F16_F32_e32 %1257:vgpr_32, implicit $mode, implicit $exec1807 %1580:vgpr_32 = nofpexcept V_CVT_F16_F32_e32 %1260:vgpr_32, implicit $mode, implicit $exec1808 %1581:vgpr_32 = nofpexcept V_CVT_F16_F32_e32 %1266:vgpr_32, implicit $mode, implicit $exec1809 %1582:vgpr_32 = nofpexcept V_CVT_F16_F32_e32 %1269:vgpr_32, implicit $mode, implicit $exec1810 %1583:vgpr_32 = nofpexcept V_CVT_F16_F32_e32 %1272:vgpr_32, implicit $mode, implicit $exec1811 %1584:vgpr_32 = nofpexcept V_CVT_F16_F32_e32 %1278:vgpr_32, implicit $mode, implicit $exec1812 %1585:vgpr_32 = nofpexcept V_CVT_F16_F32_e32 %1281:vgpr_32, implicit $mode, implicit $exec1813 %1586:vgpr_32 = nofpexcept V_CVT_F16_F32_e32 %1284:vgpr_32, implicit $mode, implicit $exec1814 %1587:vgpr_32 = nofpexcept V_CVT_F16_F32_e32 %1290:vgpr_32, implicit $mode, implicit $exec1815 %1588:vgpr_32 = nofpexcept V_CVT_F16_F32_e32 %1293:vgpr_32, implicit $mode, implicit $exec1816 %1589:vgpr_32 = nofpexcept V_CVT_F16_F32_e32 %1296:vgpr_32, implicit $mode, implicit $exec1817 %1590:vgpr_32 = V_ADD_U32_e32 %48:vgpr_32, %3345:vgpr_32, implicit $exec1818 %1591:vreg_64_align2 = BUFFER_LOAD_DWORDX2_OFFEN %1590:vgpr_32, %473:sgpr_128, 0, 0, 0, 0, implicit $exec1819 %1592:vgpr_32 = V_ADD_U32_e32 %48:vgpr_32, %3334:vgpr_32, implicit $exec1820 %1593:vreg_64_align2 = BUFFER_LOAD_DWORDX2_OFFEN %1592:vgpr_32, %473:sgpr_128, 0, 0, 0, 0, implicit $exec1821 %1594:vgpr_32 = V_ADD_U32_e32 %48:vgpr_32, %3335:vgpr_32, implicit $exec1822 %1595:vreg_64_align2 = BUFFER_LOAD_DWORDX2_OFFEN %1594:vgpr_32, %473:sgpr_128, 0, 0, 0, 0, implicit $exec1823 %1596:vgpr_32 = V_ADD_U32_e32 %48:vgpr_32, %3336:vgpr_32, implicit $exec1824 %1597:vreg_64_align2 = BUFFER_LOAD_DWORDX2_OFFEN %1596:vgpr_32, %473:sgpr_128, 0, 0, 0, 0, implicit $exec1825 INLINEASM &"s_waitcnt vmcnt($0)", 57 /* sideeffect mayload maystore isconvergent attdialect */, 13 /* imm */, 8, !01826 %1598:vreg_128_align2 = DS_READ_B128_gfx9 %44:vgpr_32, 0, 0, implicit $exec1827 %1605:vreg_128_align2 = DS_READ_B128_gfx9 %44:vgpr_32, 576, 0, implicit $exec1828 %1612:vreg_128_align2 = DS_READ_B128_gfx9 %44:vgpr_32, 1152, 0, implicit $exec1829 %1619:vreg_128_align2 = DS_READ_B128_gfx9 %44:vgpr_32, 1728, 0, implicit $exec1830 %1626:vreg_128_align2 = DS_READ_B128_gfx9 %45:vgpr_32, 0, 0, implicit $exec1831 %1633:vreg_128_align2 = DS_READ_B128_gfx9 %45:vgpr_32, 576, 0, implicit $exec1832 %1640:vreg_128_align2 = DS_READ_B128_gfx9 %45:vgpr_32, 1152, 0, implicit $exec1833 %1647:vreg_128_align2 = DS_READ_B128_gfx9 %45:vgpr_32, 1728, 0, implicit $exec1834 INLINEASM &"s_waitcnt vmcnt($0)", 57 /* sideeffect mayload maystore isconvergent attdialect */, 13 /* imm */, 8, !01835 undef %3161.sub0:vreg_64_align2 = V_PERM_B32_e64 %1593.sub0:vreg_64_align2, %1591.sub0:vreg_64_align2, %1422:sreg_32, implicit $exec1836 undef %3145.sub0:vreg_64_align2 = V_PERM_B32_e64 %1593.sub0:vreg_64_align2, %1591.sub0:vreg_64_align2, %1424:sreg_32, implicit $exec1837 %3161.sub1:vreg_64_align2 = V_PERM_B32_e64 %1597.sub0:vreg_64_align2, %1595.sub0:vreg_64_align2, %1422:sreg_32, implicit $exec1838 %3145.sub1:vreg_64_align2 = V_PERM_B32_e64 %1597.sub0:vreg_64_align2, %1595.sub0:vreg_64_align2, %1424:sreg_32, implicit $exec1839 undef %3129.sub0:vreg_64_align2 = V_PERM_B32_e64 %1593.sub1:vreg_64_align2, %1591.sub1:vreg_64_align2, %1422:sreg_32, implicit $exec1840 undef %3113.sub0:vreg_64_align2 = V_PERM_B32_e64 %1593.sub1:vreg_64_align2, %1591.sub1:vreg_64_align2, %1424:sreg_32, implicit $exec1841 %3129.sub1:vreg_64_align2 = V_PERM_B32_e64 %1597.sub1:vreg_64_align2, %1595.sub1:vreg_64_align2, %1422:sreg_32, implicit $exec1842 %3113.sub1:vreg_64_align2 = V_PERM_B32_e64 %1597.sub1:vreg_64_align2, %1595.sub1:vreg_64_align2, %1424:sreg_32, implicit $exec1843 DS_WRITE_B64_gfx9 %1447:vgpr_32, %3161:vreg_64_align2, 0, 0, implicit $exec1844 DS_WRITE_B64_gfx9 %1449:vgpr_32, %3145:vreg_64_align2, 0, 0, implicit $exec1845 DS_WRITE_B64_gfx9 %1451:vgpr_32, %3129:vreg_64_align2, 0, 0, implicit $exec1846 DS_WRITE_B64_gfx9 %1453:vgpr_32, %3113:vreg_64_align2, 0, 0, implicit $exec1847 %1678:vgpr_32 = V_ADD_U32_e32 %48:vgpr_32, %3344:vgpr_32, implicit $exec1848 %1679:vreg_64_align2 = BUFFER_LOAD_DWORDX2_OFFEN %1678:vgpr_32, %473:sgpr_128, 0, 0, 0, 0, implicit $exec1849 %1680:vgpr_32 = V_ADD_U32_e32 %48:vgpr_32, %3337:vgpr_32, implicit $exec1850 %1681:vreg_64_align2 = BUFFER_LOAD_DWORDX2_OFFEN %1680:vgpr_32, %473:sgpr_128, 0, 0, 0, 0, implicit $exec1851 %1682:vgpr_32 = V_ADD_U32_e32 %48:vgpr_32, %3338:vgpr_32, implicit $exec1852 %1683:vreg_64_align2 = BUFFER_LOAD_DWORDX2_OFFEN %1682:vgpr_32, %473:sgpr_128, 0, 0, 0, 0, implicit $exec1853 %1684:vgpr_32 = V_ADD_U32_e32 %48:vgpr_32, %3339:vgpr_32, implicit $exec1854 %1685:vreg_64_align2 = BUFFER_LOAD_DWORDX2_OFFEN %1684:vgpr_32, %473:sgpr_128, 0, 0, 0, 0, implicit $exec1855 INLINEASM &"s_waitcnt vmcnt($0)", 57 /* sideeffect mayload maystore isconvergent attdialect */, 13 /* imm */, 8, !01856 %1686:vreg_128_align2 = DS_READ_B128_gfx9 %44:vgpr_32, 0, 0, implicit $exec1857 %1693:vreg_128_align2 = DS_READ_B128_gfx9 %44:vgpr_32, 576, 0, implicit $exec1858 %1700:vreg_128_align2 = DS_READ_B128_gfx9 %44:vgpr_32, 1152, 0, implicit $exec1859 %1707:vreg_128_align2 = DS_READ_B128_gfx9 %44:vgpr_32, 1728, 0, implicit $exec1860 %1714:vreg_128_align2 = DS_READ_B128_gfx9 %45:vgpr_32, 0, 0, implicit $exec1861 %1721:vreg_128_align2 = DS_READ_B128_gfx9 %45:vgpr_32, 576, 0, implicit $exec1862 %1728:vreg_128_align2 = DS_READ_B128_gfx9 %45:vgpr_32, 1152, 0, implicit $exec1863 %1735:vreg_128_align2 = DS_READ_B128_gfx9 %45:vgpr_32, 1728, 0, implicit $exec1864 INLINEASM &"s_waitcnt vmcnt($0)", 57 /* sideeffect mayload maystore isconvergent attdialect */, 13 /* imm */, 8, !01865 undef %3062.sub0:vreg_64_align2 = V_PERM_B32_e64 %1681.sub0:vreg_64_align2, %1679.sub0:vreg_64_align2, %1422:sreg_32, implicit $exec1866 undef %3046.sub0:vreg_64_align2 = V_PERM_B32_e64 %1681.sub0:vreg_64_align2, %1679.sub0:vreg_64_align2, %1424:sreg_32, implicit $exec1867 %3062.sub1:vreg_64_align2 = V_PERM_B32_e64 %1685.sub0:vreg_64_align2, %1683.sub0:vreg_64_align2, %1422:sreg_32, implicit $exec1868 %3046.sub1:vreg_64_align2 = V_PERM_B32_e64 %1685.sub0:vreg_64_align2, %1683.sub0:vreg_64_align2, %1424:sreg_32, implicit $exec1869 undef %3029.sub0:vreg_64_align2 = V_PERM_B32_e64 %1681.sub1:vreg_64_align2, %1679.sub1:vreg_64_align2, %1422:sreg_32, implicit $exec1870 undef %3013.sub0:vreg_64_align2 = V_PERM_B32_e64 %1681.sub1:vreg_64_align2, %1679.sub1:vreg_64_align2, %1424:sreg_32, implicit $exec1871 %3029.sub1:vreg_64_align2 = V_PERM_B32_e64 %1685.sub1:vreg_64_align2, %1683.sub1:vreg_64_align2, %1422:sreg_32, implicit $exec1872 %3013.sub1:vreg_64_align2 = V_PERM_B32_e64 %1685.sub1:vreg_64_align2, %1683.sub1:vreg_64_align2, %1424:sreg_32, implicit $exec1873 DS_WRITE_B64_gfx9 %1447:vgpr_32, %3062:vreg_64_align2, 0, 0, implicit $exec1874 DS_WRITE_B64_gfx9 %1449:vgpr_32, %3046:vreg_64_align2, 0, 0, implicit $exec1875 DS_WRITE_B64_gfx9 %1451:vgpr_32, %3029:vreg_64_align2, 0, 0, implicit $exec1876 DS_WRITE_B64_gfx9 %1453:vgpr_32, %3013:vreg_64_align2, 0, 0, implicit $exec1877 %1766:vgpr_32 = V_ADD_U32_e32 %48:vgpr_32, %3343:vgpr_32, implicit $exec1878 %1767:vreg_64_align2 = BUFFER_LOAD_DWORDX2_OFFEN %1766:vgpr_32, %473:sgpr_128, 0, 0, 0, 0, implicit $exec1879 %1768:vgpr_32 = V_ADD_U32_e32 %48:vgpr_32, %3340:vgpr_32, implicit $exec1880 %1769:vreg_64_align2 = BUFFER_LOAD_DWORDX2_OFFEN %1768:vgpr_32, %473:sgpr_128, 0, 0, 0, 0, implicit $exec1881 %1770:vgpr_32 = V_ADD_U32_e32 %48:vgpr_32, %3341:vgpr_32, implicit $exec1882 %1771:vreg_64_align2 = BUFFER_LOAD_DWORDX2_OFFEN %1770:vgpr_32, %473:sgpr_128, 0, 0, 0, 0, implicit $exec1883 %1772:vgpr_32 = V_ADD_U32_e32 %48:vgpr_32, %3342:vgpr_32, implicit $exec1884 %1773:vreg_64_align2 = BUFFER_LOAD_DWORDX2_OFFEN %1772:vgpr_32, %473:sgpr_128, 0, 0, 0, 0, implicit $exec1885 INLINEASM &"s_waitcnt vmcnt($0)", 57 /* sideeffect mayload maystore isconvergent attdialect */, 13 /* imm */, 8, !01886 %1774:vreg_128_align2 = DS_READ_B128_gfx9 %44:vgpr_32, 0, 0, implicit $exec1887 %1781:vreg_128_align2 = DS_READ_B128_gfx9 %44:vgpr_32, 576, 0, implicit $exec1888 %1788:vreg_128_align2 = DS_READ_B128_gfx9 %44:vgpr_32, 1152, 0, implicit $exec1889 %1795:vreg_128_align2 = DS_READ_B128_gfx9 %44:vgpr_32, 1728, 0, implicit $exec1890 %1802:vreg_128_align2 = DS_READ_B128_gfx9 %45:vgpr_32, 0, 0, implicit $exec1891 %1809:vreg_128_align2 = DS_READ_B128_gfx9 %45:vgpr_32, 576, 0, implicit $exec1892 %1816:vreg_128_align2 = DS_READ_B128_gfx9 %45:vgpr_32, 1152, 0, implicit $exec1893 %1823:vreg_128_align2 = DS_READ_B128_gfx9 %45:vgpr_32, 1728, 0, implicit $exec1894 INLINEASM &"s_waitcnt vmcnt($0)", 57 /* sideeffect mayload maystore isconvergent attdialect */, 13 /* imm */, 8, !01895 undef %3185.sub0:vreg_64_align2 = V_PERM_B32_e64 %1769.sub0:vreg_64_align2, %1767.sub0:vreg_64_align2, %1422:sreg_32, implicit $exec1896 undef %3169.sub0:vreg_64_align2 = V_PERM_B32_e64 %1769.sub0:vreg_64_align2, %1767.sub0:vreg_64_align2, %1424:sreg_32, implicit $exec1897 %3185.sub1:vreg_64_align2 = V_PERM_B32_e64 %1773.sub0:vreg_64_align2, %1771.sub0:vreg_64_align2, %1422:sreg_32, implicit $exec1898 %3169.sub1:vreg_64_align2 = V_PERM_B32_e64 %1773.sub0:vreg_64_align2, %1771.sub0:vreg_64_align2, %1424:sreg_32, implicit $exec1899 undef %3153.sub0:vreg_64_align2 = V_PERM_B32_e64 %1769.sub1:vreg_64_align2, %1767.sub1:vreg_64_align2, %1422:sreg_32, implicit $exec1900 undef %3137.sub0:vreg_64_align2 = V_PERM_B32_e64 %1769.sub1:vreg_64_align2, %1767.sub1:vreg_64_align2, %1424:sreg_32, implicit $exec1901 %3153.sub1:vreg_64_align2 = V_PERM_B32_e64 %1773.sub1:vreg_64_align2, %1771.sub1:vreg_64_align2, %1422:sreg_32, implicit $exec1902 %3137.sub1:vreg_64_align2 = V_PERM_B32_e64 %1773.sub1:vreg_64_align2, %1771.sub1:vreg_64_align2, %1424:sreg_32, implicit $exec1903 DS_WRITE_B64_gfx9 %1447:vgpr_32, %3185:vreg_64_align2, 0, 0, implicit $exec1904 DS_WRITE_B64_gfx9 %1449:vgpr_32, %3169:vreg_64_align2, 0, 0, implicit $exec1905 DS_WRITE_B64_gfx9 %1451:vgpr_32, %3153:vreg_64_align2, 0, 0, implicit $exec1906 DS_WRITE_B64_gfx9 %1453:vgpr_32, %3137:vreg_64_align2, 0, 0, implicit $exec1907 %1854:vgpr_32 = nofpexcept V_CVT_F16_F32_e32 %1167:vgpr_32, implicit $mode, implicit $exec1908 %1855:vgpr_32 = nofpexcept V_CVT_F16_F32_e32 %1179:vgpr_32, implicit $mode, implicit $exec1909 %1856:vgpr_32 = nofpexcept V_CVT_F16_F32_e32 %1191:vgpr_32, implicit $mode, implicit $exec1910 %1857:vgpr_32 = nofpexcept V_CVT_F16_F32_e32 %1203:vgpr_32, implicit $mode, implicit $exec1911 %1858:vgpr_32 = nofpexcept V_CVT_F16_F32_e32 %1215:vgpr_32, implicit $mode, implicit $exec1912 %1859:vgpr_32 = nofpexcept V_CVT_F16_F32_e32 %1227:vgpr_32, implicit $mode, implicit $exec1913 %1860:vgpr_32 = nofpexcept V_CVT_F16_F32_e32 %1239:vgpr_32, implicit $mode, implicit $exec1914 %1861:vgpr_32 = nofpexcept V_CVT_F16_F32_e32 %1251:vgpr_32, implicit $mode, implicit $exec1915 %1862:vgpr_32 = nofpexcept V_CVT_F16_F32_e32 %1263:vgpr_32, implicit $mode, implicit $exec1916 %1863:vgpr_32 = nofpexcept V_CVT_F16_F32_e32 %1275:vgpr_32, implicit $mode, implicit $exec1917 %1864:vgpr_32 = nofpexcept V_CVT_F16_F32_e32 %1287:vgpr_32, implicit $mode, implicit $exec1918 %1865:vgpr_32 = nofpexcept V_CVT_F16_F32_e32 %1299:vgpr_32, implicit $mode, implicit $exec1919 undef %3121.sub1:vreg_64_align2 = nofpexcept V_PACK_B32_F16_e64 0, %1556:vgpr_32, 0, %1854:vgpr_32, 0, 0, implicit $mode, implicit $exec1920 %3121.sub0:vreg_64_align2 = nofpexcept V_PACK_B32_F16_e64 0, %1554:vgpr_32, 0, %1555:vgpr_32, 0, 0, implicit $mode, implicit $exec1921 undef %3105.sub1:vreg_64_align2 = nofpexcept V_PACK_B32_F16_e64 0, %1559:vgpr_32, 0, %1855:vgpr_32, 0, 0, implicit $mode, implicit $exec1922 %3105.sub0:vreg_64_align2 = nofpexcept V_PACK_B32_F16_e64 0, %1557:vgpr_32, 0, %1558:vgpr_32, 0, 0, implicit $mode, implicit $exec1923 undef %3089.sub1:vreg_64_align2 = nofpexcept V_PACK_B32_F16_e64 0, %1562:vgpr_32, 0, %1856:vgpr_32, 0, 0, implicit $mode, implicit $exec1924 %3089.sub0:vreg_64_align2 = nofpexcept V_PACK_B32_F16_e64 0, %1560:vgpr_32, 0, %1561:vgpr_32, 0, 0, implicit $mode, implicit $exec1925 undef %3073.sub1:vreg_64_align2 = nofpexcept V_PACK_B32_F16_e64 0, %1565:vgpr_32, 0, %1857:vgpr_32, 0, 0, implicit $mode, implicit $exec1926 %3073.sub0:vreg_64_align2 = nofpexcept V_PACK_B32_F16_e64 0, %1563:vgpr_32, 0, %1564:vgpr_32, 0, 0, implicit $mode, implicit $exec1927 %2986:vreg_512_align2 = contract V_MFMA_F32_32X32X8F16_mac_vgprcd_e64 %1598.sub0_sub1:vreg_128_align2, %3121:vreg_64_align2, %2986:vreg_512_align2, 0, 0, 0, implicit $mode, implicit $exec1928 %2986:vreg_512_align2 = contract V_MFMA_F32_32X32X8F16_mac_vgprcd_e64 %1598.sub2_sub3:vreg_128_align2, %3105:vreg_64_align2, %2986:vreg_512_align2, 0, 0, 0, implicit $mode, implicit $exec1929 %3038:vreg_512_align2 = contract V_MFMA_F32_32X32X8F16_mac_vgprcd_e64 %1605.sub0_sub1:vreg_128_align2, %3121:vreg_64_align2, %3038:vreg_512_align2, 0, 0, 0, implicit $mode, implicit $exec1930 %3038:vreg_512_align2 = contract V_MFMA_F32_32X32X8F16_mac_vgprcd_e64 %1605.sub2_sub3:vreg_128_align2, %3105:vreg_64_align2, %3038:vreg_512_align2, 0, 0, 0, implicit $mode, implicit $exec1931 %2980:vreg_512_align2 = contract V_MFMA_F32_32X32X8F16_mac_vgprcd_e64 %1612.sub0_sub1:vreg_128_align2, %3121:vreg_64_align2, %2980:vreg_512_align2, 0, 0, 0, implicit $mode, implicit $exec1932 %2980:vreg_512_align2 = contract V_MFMA_F32_32X32X8F16_mac_vgprcd_e64 %1612.sub2_sub3:vreg_128_align2, %3105:vreg_64_align2, %2980:vreg_512_align2, 0, 0, 0, implicit $mode, implicit $exec1933 %3003:vreg_512_align2 = contract V_MFMA_F32_32X32X8F16_mac_vgprcd_e64 %1619.sub0_sub1:vreg_128_align2, %3121:vreg_64_align2, %3003:vreg_512_align2, 0, 0, 0, implicit $mode, implicit $exec1934 %3003:vreg_512_align2 = contract V_MFMA_F32_32X32X8F16_mac_vgprcd_e64 %1619.sub2_sub3:vreg_128_align2, %3105:vreg_64_align2, %3003:vreg_512_align2, 0, 0, 0, implicit $mode, implicit $exec1935 %2986:vreg_512_align2 = contract V_MFMA_F32_32X32X8F16_mac_vgprcd_e64 %1626.sub0_sub1:vreg_128_align2, %3089:vreg_64_align2, %2986:vreg_512_align2, 0, 0, 0, implicit $mode, implicit $exec1936 %2986:vreg_512_align2 = contract V_MFMA_F32_32X32X8F16_mac_vgprcd_e64 %1626.sub2_sub3:vreg_128_align2, %3073:vreg_64_align2, %2986:vreg_512_align2, 0, 0, 0, implicit $mode, implicit $exec1937 %3038:vreg_512_align2 = contract V_MFMA_F32_32X32X8F16_mac_vgprcd_e64 %1633.sub0_sub1:vreg_128_align2, %3089:vreg_64_align2, %3038:vreg_512_align2, 0, 0, 0, implicit $mode, implicit $exec1938 %3038:vreg_512_align2 = contract V_MFMA_F32_32X32X8F16_mac_vgprcd_e64 %1633.sub2_sub3:vreg_128_align2, %3073:vreg_64_align2, %3038:vreg_512_align2, 0, 0, 0, implicit $mode, implicit $exec1939 %2980:vreg_512_align2 = contract V_MFMA_F32_32X32X8F16_mac_vgprcd_e64 %1640.sub0_sub1:vreg_128_align2, %3089:vreg_64_align2, %2980:vreg_512_align2, 0, 0, 0, implicit $mode, implicit $exec1940 %2980:vreg_512_align2 = contract V_MFMA_F32_32X32X8F16_mac_vgprcd_e64 %1640.sub2_sub3:vreg_128_align2, %3073:vreg_64_align2, %2980:vreg_512_align2, 0, 0, 0, implicit $mode, implicit $exec1941 %3003:vreg_512_align2 = contract V_MFMA_F32_32X32X8F16_mac_vgprcd_e64 %1647.sub0_sub1:vreg_128_align2, %3089:vreg_64_align2, %3003:vreg_512_align2, 0, 0, 0, implicit $mode, implicit $exec1942 %3003:vreg_512_align2 = contract V_MFMA_F32_32X32X8F16_mac_vgprcd_e64 %1647.sub2_sub3:vreg_128_align2, %3073:vreg_64_align2, %3003:vreg_512_align2, 0, 0, 0, implicit $mode, implicit $exec1943 undef %2993.sub1:vreg_64_align2 = nofpexcept V_PACK_B32_F16_e64 0, %1568:vgpr_32, 0, %1858:vgpr_32, 0, 0, implicit $mode, implicit $exec1944 %2993.sub0:vreg_64_align2 = nofpexcept V_PACK_B32_F16_e64 0, %1566:vgpr_32, 0, %1567:vgpr_32, 0, 0, implicit $mode, implicit $exec1945 undef %3195.sub1:vreg_64_align2 = nofpexcept V_PACK_B32_F16_e64 0, %1571:vgpr_32, 0, %1859:vgpr_32, 0, 0, implicit $mode, implicit $exec1946 %3195.sub0:vreg_64_align2 = nofpexcept V_PACK_B32_F16_e64 0, %1569:vgpr_32, 0, %1570:vgpr_32, 0, 0, implicit $mode, implicit $exec1947 undef %3178.sub1:vreg_64_align2 = nofpexcept V_PACK_B32_F16_e64 0, %1574:vgpr_32, 0, %1860:vgpr_32, 0, 0, implicit $mode, implicit $exec1948 %3178.sub0:vreg_64_align2 = nofpexcept V_PACK_B32_F16_e64 0, %1572:vgpr_32, 0, %1573:vgpr_32, 0, 0, implicit $mode, implicit $exec1949 undef %3162.sub1:vreg_64_align2 = nofpexcept V_PACK_B32_F16_e64 0, %1577:vgpr_32, 0, %1861:vgpr_32, 0, 0, implicit $mode, implicit $exec1950 %3162.sub0:vreg_64_align2 = nofpexcept V_PACK_B32_F16_e64 0, %1575:vgpr_32, 0, %1576:vgpr_32, 0, 0, implicit $mode, implicit $exec1951 %2986:vreg_512_align2 = contract V_MFMA_F32_32X32X8F16_mac_vgprcd_e64 %1686.sub0_sub1:vreg_128_align2, %2993:vreg_64_align2, %2986:vreg_512_align2, 0, 0, 0, implicit $mode, implicit $exec1952 %2986:vreg_512_align2 = contract V_MFMA_F32_32X32X8F16_mac_vgprcd_e64 %1686.sub2_sub3:vreg_128_align2, %3195:vreg_64_align2, %2986:vreg_512_align2, 0, 0, 0, implicit $mode, implicit $exec1953 %3038:vreg_512_align2 = contract V_MFMA_F32_32X32X8F16_mac_vgprcd_e64 %1693.sub0_sub1:vreg_128_align2, %2993:vreg_64_align2, %3038:vreg_512_align2, 0, 0, 0, implicit $mode, implicit $exec1954 %3038:vreg_512_align2 = contract V_MFMA_F32_32X32X8F16_mac_vgprcd_e64 %1693.sub2_sub3:vreg_128_align2, %3195:vreg_64_align2, %3038:vreg_512_align2, 0, 0, 0, implicit $mode, implicit $exec1955 %2980:vreg_512_align2 = contract V_MFMA_F32_32X32X8F16_mac_vgprcd_e64 %1700.sub0_sub1:vreg_128_align2, %2993:vreg_64_align2, %2980:vreg_512_align2, 0, 0, 0, implicit $mode, implicit $exec1956 %2980:vreg_512_align2 = contract V_MFMA_F32_32X32X8F16_mac_vgprcd_e64 %1700.sub2_sub3:vreg_128_align2, %3195:vreg_64_align2, %2980:vreg_512_align2, 0, 0, 0, implicit $mode, implicit $exec1957 %3003:vreg_512_align2 = contract V_MFMA_F32_32X32X8F16_mac_vgprcd_e64 %1707.sub0_sub1:vreg_128_align2, %2993:vreg_64_align2, %3003:vreg_512_align2, 0, 0, 0, implicit $mode, implicit $exec1958 %3003:vreg_512_align2 = contract V_MFMA_F32_32X32X8F16_mac_vgprcd_e64 %1707.sub2_sub3:vreg_128_align2, %3195:vreg_64_align2, %3003:vreg_512_align2, 0, 0, 0, implicit $mode, implicit $exec1959 %2986:vreg_512_align2 = contract V_MFMA_F32_32X32X8F16_mac_vgprcd_e64 %1714.sub0_sub1:vreg_128_align2, %3178:vreg_64_align2, %2986:vreg_512_align2, 0, 0, 0, implicit $mode, implicit $exec1960 %2986:vreg_512_align2 = contract V_MFMA_F32_32X32X8F16_mac_vgprcd_e64 %1714.sub2_sub3:vreg_128_align2, %3162:vreg_64_align2, %2986:vreg_512_align2, 0, 0, 0, implicit $mode, implicit $exec1961 %3038:vreg_512_align2 = contract V_MFMA_F32_32X32X8F16_mac_vgprcd_e64 %1721.sub0_sub1:vreg_128_align2, %3178:vreg_64_align2, %3038:vreg_512_align2, 0, 0, 0, implicit $mode, implicit $exec1962 %3038:vreg_512_align2 = contract V_MFMA_F32_32X32X8F16_mac_vgprcd_e64 %1721.sub2_sub3:vreg_128_align2, %3162:vreg_64_align2, %3038:vreg_512_align2, 0, 0, 0, implicit $mode, implicit $exec1963 %2980:vreg_512_align2 = contract V_MFMA_F32_32X32X8F16_mac_vgprcd_e64 %1728.sub0_sub1:vreg_128_align2, %3178:vreg_64_align2, %2980:vreg_512_align2, 0, 0, 0, implicit $mode, implicit $exec1964 %2980:vreg_512_align2 = contract V_MFMA_F32_32X32X8F16_mac_vgprcd_e64 %1728.sub2_sub3:vreg_128_align2, %3162:vreg_64_align2, %2980:vreg_512_align2, 0, 0, 0, implicit $mode, implicit $exec1965 %3003:vreg_512_align2 = contract V_MFMA_F32_32X32X8F16_mac_vgprcd_e64 %1735.sub0_sub1:vreg_128_align2, %3178:vreg_64_align2, %3003:vreg_512_align2, 0, 0, 0, implicit $mode, implicit $exec1966 %3003:vreg_512_align2 = contract V_MFMA_F32_32X32X8F16_mac_vgprcd_e64 %1735.sub2_sub3:vreg_128_align2, %3162:vreg_64_align2, %3003:vreg_512_align2, 0, 0, 0, implicit $mode, implicit $exec1967 undef %3146.sub1:vreg_64_align2 = nofpexcept V_PACK_B32_F16_e64 0, %1580:vgpr_32, 0, %1862:vgpr_32, 0, 0, implicit $mode, implicit $exec1968 %3146.sub0:vreg_64_align2 = nofpexcept V_PACK_B32_F16_e64 0, %1578:vgpr_32, 0, %1579:vgpr_32, 0, 0, implicit $mode, implicit $exec1969 undef %3130.sub1:vreg_64_align2 = nofpexcept V_PACK_B32_F16_e64 0, %1583:vgpr_32, 0, %1863:vgpr_32, 0, 0, implicit $mode, implicit $exec1970 %3130.sub0:vreg_64_align2 = nofpexcept V_PACK_B32_F16_e64 0, %1581:vgpr_32, 0, %1582:vgpr_32, 0, 0, implicit $mode, implicit $exec1971 undef %3114.sub1:vreg_64_align2 = nofpexcept V_PACK_B32_F16_e64 0, %1586:vgpr_32, 0, %1864:vgpr_32, 0, 0, implicit $mode, implicit $exec1972 %3114.sub0:vreg_64_align2 = nofpexcept V_PACK_B32_F16_e64 0, %1584:vgpr_32, 0, %1585:vgpr_32, 0, 0, implicit $mode, implicit $exec1973 undef %3098.sub1:vreg_64_align2 = nofpexcept V_PACK_B32_F16_e64 0, %1589:vgpr_32, 0, %1865:vgpr_32, 0, 0, implicit $mode, implicit $exec1974 %3098.sub0:vreg_64_align2 = nofpexcept V_PACK_B32_F16_e64 0, %1587:vgpr_32, 0, %1588:vgpr_32, 0, 0, implicit $mode, implicit $exec1975 %2986:vreg_512_align2 = contract V_MFMA_F32_32X32X8F16_mac_vgprcd_e64 %1774.sub0_sub1:vreg_128_align2, %3146:vreg_64_align2, %2986:vreg_512_align2, 0, 0, 0, implicit $mode, implicit $exec1976 %2986:vreg_512_align2 = contract V_MFMA_F32_32X32X8F16_mac_vgprcd_e64 %1774.sub2_sub3:vreg_128_align2, %3130:vreg_64_align2, %2986:vreg_512_align2, 0, 0, 0, implicit $mode, implicit $exec1977 %3038:vreg_512_align2 = contract V_MFMA_F32_32X32X8F16_mac_vgprcd_e64 %1781.sub0_sub1:vreg_128_align2, %3146:vreg_64_align2, %3038:vreg_512_align2, 0, 0, 0, implicit $mode, implicit $exec1978 %3038:vreg_512_align2 = contract V_MFMA_F32_32X32X8F16_mac_vgprcd_e64 %1781.sub2_sub3:vreg_128_align2, %3130:vreg_64_align2, %3038:vreg_512_align2, 0, 0, 0, implicit $mode, implicit $exec1979 %2980:vreg_512_align2 = contract V_MFMA_F32_32X32X8F16_mac_vgprcd_e64 %1788.sub0_sub1:vreg_128_align2, %3146:vreg_64_align2, %2980:vreg_512_align2, 0, 0, 0, implicit $mode, implicit $exec1980 %2980:vreg_512_align2 = contract V_MFMA_F32_32X32X8F16_mac_vgprcd_e64 %1788.sub2_sub3:vreg_128_align2, %3130:vreg_64_align2, %2980:vreg_512_align2, 0, 0, 0, implicit $mode, implicit $exec1981 %3003:vreg_512_align2 = contract V_MFMA_F32_32X32X8F16_mac_vgprcd_e64 %1795.sub0_sub1:vreg_128_align2, %3146:vreg_64_align2, %3003:vreg_512_align2, 0, 0, 0, implicit $mode, implicit $exec1982 %3003:vreg_512_align2 = contract V_MFMA_F32_32X32X8F16_mac_vgprcd_e64 %1795.sub2_sub3:vreg_128_align2, %3130:vreg_64_align2, %3003:vreg_512_align2, 0, 0, 0, implicit $mode, implicit $exec1983 %2986:vreg_512_align2 = contract V_MFMA_F32_32X32X8F16_mac_vgprcd_e64 %1802.sub0_sub1:vreg_128_align2, %3114:vreg_64_align2, %2986:vreg_512_align2, 0, 0, 0, implicit $mode, implicit $exec1984 %2986:vreg_512_align2 = contract V_MFMA_F32_32X32X8F16_mac_vgprcd_e64 %1802.sub2_sub3:vreg_128_align2, %3098:vreg_64_align2, %2986:vreg_512_align2, 0, 0, 0, implicit $mode, implicit $exec1985 %3038:vreg_512_align2 = contract V_MFMA_F32_32X32X8F16_mac_vgprcd_e64 %1809.sub0_sub1:vreg_128_align2, %3114:vreg_64_align2, %3038:vreg_512_align2, 0, 0, 0, implicit $mode, implicit $exec1986 %3038:vreg_512_align2 = contract V_MFMA_F32_32X32X8F16_mac_vgprcd_e64 %1809.sub2_sub3:vreg_128_align2, %3098:vreg_64_align2, %3038:vreg_512_align2, 0, 0, 0, implicit $mode, implicit $exec1987 %2980:vreg_512_align2 = contract V_MFMA_F32_32X32X8F16_mac_vgprcd_e64 %1816.sub0_sub1:vreg_128_align2, %3114:vreg_64_align2, %2980:vreg_512_align2, 0, 0, 0, implicit $mode, implicit $exec1988 %2980:vreg_512_align2 = contract V_MFMA_F32_32X32X8F16_mac_vgprcd_e64 %1816.sub2_sub3:vreg_128_align2, %3098:vreg_64_align2, %2980:vreg_512_align2, 0, 0, 0, implicit $mode, implicit $exec1989 %3003:vreg_512_align2 = contract V_MFMA_F32_32X32X8F16_mac_vgprcd_e64 %1823.sub0_sub1:vreg_128_align2, %3114:vreg_64_align2, %3003:vreg_512_align2, 0, 0, 0, implicit $mode, implicit $exec1990 %3003:vreg_512_align2 = contract V_MFMA_F32_32X32X8F16_mac_vgprcd_e64 %1823.sub2_sub3:vreg_128_align2, %3098:vreg_64_align2, %3003:vreg_512_align2, 0, 0, 0, implicit $mode, implicit $exec1991 %2054:vgpr_32 = nofpexcept V_CVT_F16_F32_e32 %1347:vgpr_32, implicit $mode, implicit $exec1992 %2055:vgpr_32 = nofpexcept V_CVT_F16_F32_e32 %1341:vgpr_32, implicit $mode, implicit $exec1993 %2056:vgpr_32 = nofpexcept V_CVT_F16_F32_e32 %1335:vgpr_32, implicit $mode, implicit $exec1994 %2057:vgpr_32 = nofpexcept V_CVT_F16_F32_e32 %1329:vgpr_32, implicit $mode, implicit $exec1995 %2058:vgpr_32 = nofpexcept V_CVT_F16_F32_e32 %1323:vgpr_32, implicit $mode, implicit $exec1996 %2059:vgpr_32 = nofpexcept V_CVT_F16_F32_e32 %1317:vgpr_32, implicit $mode, implicit $exec1997 %2060:vgpr_32 = nofpexcept V_CVT_F16_F32_e32 %1311:vgpr_32, implicit $mode, implicit $exec1998 %2061:vgpr_32 = nofpexcept V_CVT_F16_F32_e32 %1305:vgpr_32, implicit $mode, implicit $exec1999 %2062:vgpr_32 = nofpexcept V_CVT_F16_F32_e32 %1344:vgpr_32, implicit $mode, implicit $exec2000 %2063:vgpr_32 = nofpexcept V_CVT_F16_F32_e32 %1338:vgpr_32, implicit $mode, implicit $exec2001 %2064:vgpr_32 = nofpexcept V_CVT_F16_F32_e32 %1332:vgpr_32, implicit $mode, implicit $exec2002 %2065:vgpr_32 = nofpexcept V_CVT_F16_F32_e32 %1326:vgpr_32, implicit $mode, implicit $exec2003 %2066:vgpr_32 = nofpexcept V_CVT_F16_F32_e32 %1320:vgpr_32, implicit $mode, implicit $exec2004 %2067:vgpr_32 = nofpexcept V_CVT_F16_F32_e32 %1314:vgpr_32, implicit $mode, implicit $exec2005 %2068:vgpr_32 = nofpexcept V_CVT_F16_F32_e32 %1308:vgpr_32, implicit $mode, implicit $exec2006 %2069:vgpr_32 = nofpexcept V_CVT_F16_F32_e32 %1302:vgpr_32, implicit $mode, implicit $exec2007 INLINEASM &"s_waitcnt vmcnt($0)", 57 /* sideeffect mayload maystore isconvergent attdialect */, 13 /* imm */, 8, !02008 undef %3082.sub1:vreg_64_align2 = nofpexcept V_PACK_B32_F16_e64 0, %2068:vgpr_32, 0, %2060:vgpr_32, 0, 0, implicit $mode, implicit $exec2009 %3082.sub0:vreg_64_align2 = nofpexcept V_PACK_B32_F16_e64 0, %2069:vgpr_32, 0, %2061:vgpr_32, 0, 0, implicit $mode, implicit $exec2010 undef %3066.sub1:vreg_64_align2 = nofpexcept V_PACK_B32_F16_e64 0, %2066:vgpr_32, 0, %2058:vgpr_32, 0, 0, implicit $mode, implicit $exec2011 %3066.sub0:vreg_64_align2 = nofpexcept V_PACK_B32_F16_e64 0, %2067:vgpr_32, 0, %2059:vgpr_32, 0, 0, implicit $mode, implicit $exec2012 undef %3050.sub1:vreg_64_align2 = nofpexcept V_PACK_B32_F16_e64 0, %2064:vgpr_32, 0, %2056:vgpr_32, 0, 0, implicit $mode, implicit $exec2013 %3050.sub0:vreg_64_align2 = nofpexcept V_PACK_B32_F16_e64 0, %2065:vgpr_32, 0, %2057:vgpr_32, 0, 0, implicit $mode, implicit $exec2014 undef %3033.sub1:vreg_64_align2 = nofpexcept V_PACK_B32_F16_e64 0, %2062:vgpr_32, 0, %2054:vgpr_32, 0, 0, implicit $mode, implicit $exec2015 %3033.sub0:vreg_64_align2 = nofpexcept V_PACK_B32_F16_e64 0, %2063:vgpr_32, 0, %2055:vgpr_32, 0, 0, implicit $mode, implicit $exec2016 %2082:vreg_128_align2 = DS_READ_B128_gfx9 %44:vgpr_32, 0, 0, implicit $exec2017 %2986:vreg_512_align2 = contract V_MFMA_F32_32X32X8F16_mac_vgprcd_e64 %2082.sub0_sub1:vreg_128_align2, %3082:vreg_64_align2, %2986:vreg_512_align2, 0, 0, 0, implicit $mode, implicit $exec2018 %2986:vreg_512_align2 = contract V_MFMA_F32_32X32X8F16_mac_vgprcd_e64 %2082.sub2_sub3:vreg_128_align2, %3066:vreg_64_align2, %2986:vreg_512_align2, 0, 0, 0, implicit $mode, implicit $exec2019 %2095:vreg_128_align2 = DS_READ_B128_gfx9 %44:vgpr_32, 576, 0, implicit $exec2020 %3038:vreg_512_align2 = contract V_MFMA_F32_32X32X8F16_mac_vgprcd_e64 %2095.sub0_sub1:vreg_128_align2, %3082:vreg_64_align2, %3038:vreg_512_align2, 0, 0, 0, implicit $mode, implicit $exec2021 %3038:vreg_512_align2 = contract V_MFMA_F32_32X32X8F16_mac_vgprcd_e64 %2095.sub2_sub3:vreg_128_align2, %3066:vreg_64_align2, %3038:vreg_512_align2, 0, 0, 0, implicit $mode, implicit $exec2022 %2108:vreg_128_align2 = DS_READ_B128_gfx9 %44:vgpr_32, 1152, 0, implicit $exec2023 %2980:vreg_512_align2 = contract V_MFMA_F32_32X32X8F16_mac_vgprcd_e64 %2108.sub0_sub1:vreg_128_align2, %3082:vreg_64_align2, %2980:vreg_512_align2, 0, 0, 0, implicit $mode, implicit $exec2024 %2980:vreg_512_align2 = contract V_MFMA_F32_32X32X8F16_mac_vgprcd_e64 %2108.sub2_sub3:vreg_128_align2, %3066:vreg_64_align2, %2980:vreg_512_align2, 0, 0, 0, implicit $mode, implicit $exec2025 %2121:vreg_128_align2 = DS_READ_B128_gfx9 %44:vgpr_32, 1728, 0, implicit $exec2026 %3003:vreg_512_align2 = contract V_MFMA_F32_32X32X8F16_mac_vgprcd_e64 %2121.sub0_sub1:vreg_128_align2, %3082:vreg_64_align2, %3003:vreg_512_align2, 0, 0, 0, implicit $mode, implicit $exec2027 %3003:vreg_512_align2 = contract V_MFMA_F32_32X32X8F16_mac_vgprcd_e64 %2121.sub2_sub3:vreg_128_align2, %3066:vreg_64_align2, %3003:vreg_512_align2, 0, 0, 0, implicit $mode, implicit $exec2028 %2134:vreg_128_align2 = DS_READ_B128_gfx9 %45:vgpr_32, 0, 0, implicit $exec2029 %2986:vreg_512_align2 = contract V_MFMA_F32_32X32X8F16_mac_vgprcd_e64 %2134.sub0_sub1:vreg_128_align2, %3050:vreg_64_align2, %2986:vreg_512_align2, 0, 0, 0, implicit $mode, implicit $exec2030 %2986:vreg_512_align2 = contract V_MFMA_F32_32X32X8F16_mac_vgprcd_e64 %2134.sub2_sub3:vreg_128_align2, %3033:vreg_64_align2, %2986:vreg_512_align2, 0, 0, 0, implicit $mode, implicit $exec2031 %2146:vreg_128_align2 = DS_READ_B128_gfx9 %45:vgpr_32, 576, 0, implicit $exec2032 %3038:vreg_512_align2 = contract V_MFMA_F32_32X32X8F16_mac_vgprcd_e64 %2146.sub0_sub1:vreg_128_align2, %3050:vreg_64_align2, %3038:vreg_512_align2, 0, 0, 0, implicit $mode, implicit $exec2033 %3038:vreg_512_align2 = contract V_MFMA_F32_32X32X8F16_mac_vgprcd_e64 %2146.sub2_sub3:vreg_128_align2, %3033:vreg_64_align2, %3038:vreg_512_align2, 0, 0, 0, implicit $mode, implicit $exec2034 %2158:vreg_128_align2 = DS_READ_B128_gfx9 %45:vgpr_32, 1152, 0, implicit $exec2035 %2980:vreg_512_align2 = contract V_MFMA_F32_32X32X8F16_mac_vgprcd_e64 %2158.sub0_sub1:vreg_128_align2, %3050:vreg_64_align2, %2980:vreg_512_align2, 0, 0, 0, implicit $mode, implicit $exec2036 %2980:vreg_512_align2 = contract V_MFMA_F32_32X32X8F16_mac_vgprcd_e64 %2158.sub2_sub3:vreg_128_align2, %3033:vreg_64_align2, %2980:vreg_512_align2, 0, 0, 0, implicit $mode, implicit $exec2037 %2170:vreg_128_align2 = DS_READ_B128_gfx9 %45:vgpr_32, 1728, 0, implicit $exec2038 %3003:vreg_512_align2 = contract V_MFMA_F32_32X32X8F16_mac_vgprcd_e64 %2170.sub0_sub1:vreg_128_align2, %3050:vreg_64_align2, %3003:vreg_512_align2, 0, 0, 0, implicit $mode, implicit $exec2039 %3003:vreg_512_align2 = contract V_MFMA_F32_32X32X8F16_mac_vgprcd_e64 %2170.sub2_sub3:vreg_128_align2, %3033:vreg_64_align2, %3003:vreg_512_align2, 0, 0, 0, implicit $mode, implicit $exec2040 INLINEASM &"s_waitcnt vmcnt($0)", 57 /* sideeffect mayload maystore isconvergent attdialect */, 13 /* imm */, 8, !02041 %3345:vgpr_32 = V_ADD_U32_e32 %50:sreg_32, %3345:vgpr_32, implicit $exec2042 %3344:vgpr_32 = V_ADD_U32_e32 %50:sreg_32, %3344:vgpr_32, implicit $exec2043 %3343:vgpr_32 = V_ADD_U32_e32 %50:sreg_32, %3343:vgpr_32, implicit $exec2044 %3342:vgpr_32 = V_ADD_U32_e32 %50:sreg_32, %3342:vgpr_32, implicit $exec2045 %3341:vgpr_32 = V_ADD_U32_e32 %50:sreg_32, %3341:vgpr_32, implicit $exec2046 %3340:vgpr_32 = V_ADD_U32_e32 %50:sreg_32, %3340:vgpr_32, implicit $exec2047 %3339:vgpr_32 = V_ADD_U32_e32 %50:sreg_32, %3339:vgpr_32, implicit $exec2048 %3338:vgpr_32 = V_ADD_U32_e32 %50:sreg_32, %3338:vgpr_32, implicit $exec2049 %3337:vgpr_32 = V_ADD_U32_e32 %50:sreg_32, %3337:vgpr_32, implicit $exec2050 %3336:vgpr_32 = V_ADD_U32_e32 %50:sreg_32, %3336:vgpr_32, implicit $exec2051 %3335:vgpr_32 = V_ADD_U32_e32 %50:sreg_32, %3335:vgpr_32, implicit $exec2052 %3334:vgpr_32 = V_ADD_U32_e32 %50:sreg_32, %3334:vgpr_32, implicit $exec2053 %3333:vgpr_32 = V_ADD_U32_e32 %50:sreg_32, %3333:vgpr_32, implicit $exec2054 %3332:vgpr_32 = V_ADD_U32_e32 %50:sreg_32, %3332:vgpr_32, implicit $exec2055 %3331:vgpr_32 = V_ADD_U32_e32 %50:sreg_32, %3331:vgpr_32, implicit $exec2056 %3330:vgpr_32 = V_ADD_U32_e32 %50:sreg_32, %3330:vgpr_32, implicit $exec2057 %3329:vgpr_32 = nuw V_ADD_U32_e32 128, %3329:vgpr_32, implicit $exec2058 S_ENDPGM 02059...2060