llvm/test/CodeGen/AMDGPU/llvm.amdgcn.iglp.opt.exp.small.mir

   1 # NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 2
   2 # RUN: llc -mtriple=amdgcn-amd-amdhsa -mcpu=gfx942 -start-before=machine-scheduler -verify-misched -o - %s | FileCheck -check-prefix=GCN %s
   3 # 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 %s
   4
   5 --- |
   6   define amdgpu_kernel void @smallInterleave() #0 { ret void }
   7   ; GCN-LABEL: smallInterleave:
   8   ; GCN:       ; %bb.0:
   9   ; GCN-NEXT:    ; implicit-def: $vgpr2
  10   ; GCN-NEXT:    s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)
  11   ; GCN-NEXT:    v_readfirstlane_b32 s20, v2
  12   ; GCN-NEXT:    ; implicit-def: $sgpr4
  13   ; GCN-NEXT:    ; implicit-def: $vgpr3
  14   ; GCN-NEXT:    ; implicit-def: $vgpr0_vgpr1
  15   ; GCN-NEXT:    ; implicit-def: $sgpr0_sgpr1_sgpr2_sgpr3
  16   ; GCN-NEXT:    ; implicit-def: $vgpr50
  17   ; GCN-NEXT:    ; implicit-def: $sgpr16_sgpr17_sgpr18_sgpr19
  18   ; GCN-NEXT:    ; implicit-def: $vgpr49
  19   ; GCN-NEXT:    ; implicit-def: $vgpr40_vgpr41_vgpr42_vgpr43
  20   ; GCN-NEXT:    ; implicit-def: $vgpr51
  21   ; GCN-NEXT:    ; implicit-def: $vgpr62_vgpr63_vgpr64_vgpr65
  22   ; GCN-NEXT:    ; implicit-def: $vgpr76
  23   ; GCN-NEXT:    ; implicit-def: $vgpr77
  24   ; GCN-NEXT:    ; implicit-def: $vgpr78
  25   ; GCN-NEXT:    ; implicit-def: $vgpr79
  26   ; GCN-NEXT:    ; implicit-def: $vgpr80
  27   ; GCN-NEXT:    ; implicit-def: $vgpr91
  28   ; GCN-NEXT:    ; kill: killed $sgpr16_sgpr17_sgpr18_sgpr19
  29   ; GCN-NEXT:    ; iglp_opt mask(0x00000002)
  30   ; GCN-NEXT:    s_nop 1
  31   ; GCN-NEXT:    v_lshl_add_u32 v2, s20, 4, v3
  32   ; GCN-NEXT:    v_mad_u64_u32 v[4:5], s[4:5], s4, v2, v[0:1]
  33   ; GCN-NEXT:    buffer_load_dwordx4 v[0:3], v4, s[0:3], 0 offen sc0 sc1
  34   ; GCN-NEXT:    s_waitcnt vmcnt(0)
  35   ; GCN-NEXT:    buffer_inv sc0 sc1
  36   ; GCN-NEXT:    s_lshl_b32 s4, s20, 7
  37   ; GCN-NEXT:    ; implicit-def: $vgpr5
  38   ; GCN-NEXT:    v_add_lshl_u32 v48, v5, s4, 1
  39   ; GCN-NEXT:    v_add_u32_e32 v76, s20, v76
  40   ; GCN-NEXT:    v_and_b32_e32 v76, 0x1fffffff, v76
  41   ; GCN-NEXT:    buffer_wbl2 sc0 sc1
  42   ; GCN-NEXT:    ds_write_b128 v48, v[0:3]
  43   ; GCN-NEXT:    s_waitcnt lgkmcnt(0)
  44   ; GCN-NEXT:    buffer_load_dwordx4 v[32:35], v4, s[0:3], 0 offen offset:64 sc0 sc1
  45   ; GCN-NEXT:    s_waitcnt vmcnt(0)
  46   ; GCN-NEXT:    buffer_inv sc0 sc1
  47   ; GCN-NEXT:    ; implicit-def: $vgpr0
  48   ; GCN-NEXT:    ; implicit-def: $vgpr1
  49   ; GCN-NEXT:    ; implicit-def: $sgpr0_sgpr1_sgpr2_sgpr3_sgpr4_sgpr5_sgpr6_sgpr7_sgpr8_sgpr9_sgpr10_sgpr11_sgpr12_sgpr13_sgpr14_sgpr15
  50   ; GCN-NEXT:    ; implicit-def: $sgpr6
  51   ; GCN-NEXT:    v_add_u32_e32 v0, v0, v50
  52   ; GCN-NEXT:    v_add_u32_e32 v1, v1, v50
  53   ; GCN-NEXT:    buffer_load_dwordx2 v[72:73], v0, s[16:19], 0 offen sc0 sc1
  54   ; GCN-NEXT:    s_waitcnt vmcnt(0)
  55   ; GCN-NEXT:    buffer_inv sc0 sc1
  56   ; GCN-NEXT:    buffer_load_dwordx2 v[74:75], v1, s[16:19], 0 offen sc0 sc1
  57   ; GCN-NEXT:    s_waitcnt vmcnt(0)
  58   ; GCN-NEXT:    buffer_inv sc0 sc1
  59   ; GCN-NEXT:    ;;#ASMSTART
  60   ; GCN-NEXT:    s_waitcnt vmcnt(8)
  61   ; GCN-NEXT:    ;;#ASMEND
  62   ; GCN-NEXT:    ds_read_b128 v[36:39], v49
  63   ; GCN-NEXT:    s_waitcnt lgkmcnt(0)
  64   ; GCN-NEXT:    buffer_inv sc0 sc1
  65   ; GCN-NEXT:    ds_read_b128 v[44:47], v49 offset:512
  66   ; GCN-NEXT:    s_waitcnt lgkmcnt(0)
  67   ; GCN-NEXT:    buffer_inv sc0 sc1
  68   ; GCN-NEXT:    v_mfma_f32_32x32x8_f16 v[16:31], v[36:37], v[40:41], 0
  69   ; GCN-NEXT:    ; kill: killed $vgpr1
  70   ; GCN-NEXT:    ; kill: killed $vgpr0
  71   ; GCN-NEXT:    v_mul_lo_u32 v76, v76, s6
  72   ; GCN-NEXT:    v_add_lshl_u32 v76, v77, v76, 1
  73   ; GCN-NEXT:    v_lshl_add_u32 v77, v78, 1, v76
  74   ; GCN-NEXT:    ; implicit-def: $sgpr5
  75   ; GCN-NEXT:    v_lshl_add_u32 v78, v79, 1, v77
  76   ; GCN-NEXT:    ; implicit-def: $sgpr2
  77   ; GCN-NEXT:    ; implicit-def: $sgpr3
  78   ; GCN-NEXT:    v_lshl_add_u32 v79, v80, 1, v78
  79   ; GCN-NEXT:    ; implicit-def: $sgpr0_sgpr1
  80   ; GCN-NEXT:    v_mfma_f32_32x32x8_f16 v[0:15], v[44:45], v[40:41], 0
  81   ; GCN-NEXT:    v_mfma_f32_32x32x8_f16 v[16:31], v[38:39], v[42:43], v[16:31]
  82   ; GCN-NEXT:    ds_read_b128 v[36:39], v51
  83   ; GCN-NEXT:    s_waitcnt lgkmcnt(0)
  84   ; GCN-NEXT:    buffer_inv sc0 sc1
  85   ; GCN-NEXT:    v_mfma_f32_32x32x8_f16 v[0:15], v[46:47], v[42:43], v[0:15]
  86   ; GCN-NEXT:    ds_read_b128 v[44:47], v51 offset:512
  87   ; GCN-NEXT:    s_waitcnt lgkmcnt(0)
  88   ; GCN-NEXT:    buffer_inv sc0 sc1
  89   ; GCN-NEXT:    ; implicit-def: $vgpr40_vgpr41_vgpr42_vgpr43
  90   ; GCN-NEXT:    ;;#ASMSTART
  91   ; GCN-NEXT:    s_waitcnt vmcnt(8)
  92   ; GCN-NEXT:    ;;#ASMEND
  93   ; GCN-NEXT:    buffer_wbl2 sc0 sc1
  94   ; GCN-NEXT:    ds_write_b128 v48, v[32:35]
  95   ; GCN-NEXT:    v_mfma_f32_32x32x8_f16 v[16:31], v[36:37], v[40:41], v[16:31]
  96   ; GCN-NEXT:    ;;#ASMSTART
  97   ; GCN-NEXT:    s_waitcnt vmcnt(8)
  98   ; GCN-NEXT:    ;;#ASMEND
  99   ; GCN-NEXT:    s_waitcnt lgkmcnt(0)
 100   ; GCN-NEXT:    ds_read_b128 v[32:35], v49
 101   ; GCN-NEXT:    s_waitcnt lgkmcnt(0)
 102   ; GCN-NEXT:    buffer_inv sc0 sc1
 103   ; GCN-NEXT:    v_mfma_f32_32x32x8_f16 v[0:15], v[44:45], v[40:41], v[0:15]
 104   ; GCN-NEXT:    v_mfma_f32_32x32x8_f16 v[16:31], v[38:39], v[42:43], v[16:31]
 105   ; GCN-NEXT:    ; implicit-def: $vgpr36_vgpr37_vgpr38_vgpr39
 106   ; GCN-NEXT:    v_mfma_f32_32x32x8_f16 v[0:15], v[46:47], v[42:43], v[0:15]
 107   ; GCN-NEXT:    ds_read_b128 v[40:43], v49 offset:512
 108   ; GCN-NEXT:    s_waitcnt lgkmcnt(0)
 109   ; GCN-NEXT:    buffer_inv sc0 sc1
 110   ; GCN-NEXT:    ds_read_b128 v[68:71], v51
 111   ; GCN-NEXT:    s_waitcnt lgkmcnt(0)
 112   ; GCN-NEXT:    buffer_inv sc0 sc1
 113   ; GCN-NEXT:    v_mfma_f32_32x32x8_f16 v[16:31], v[32:33], v[36:37], v[16:31]
 114   ; GCN-NEXT:    ; implicit-def: $vgpr32
 115   ; GCN-NEXT:    ; implicit-def: $vgpr33
 116   ; GCN-NEXT:    v_add_u32_e32 v82, v32, v50
 117   ; GCN-NEXT:    v_add_u32_e32 v83, v33, v50
 118   ; GCN-NEXT:    ; kill: killed $vgpr82
 119   ; GCN-NEXT:    ; kill: killed $vgpr83
 120   ; GCN-NEXT:    v_mfma_f32_32x32x8_f16 v[16:31], v[34:35], v[38:39], v[16:31]
 121   ; GCN-NEXT:    v_mfma_f32_32x32x8_f16 v[0:15], v[40:41], v[36:37], v[0:15]
 122   ; GCN-NEXT:    v_mfma_f32_32x32x8_f16 v[16:31], v[68:69], v[62:63], v[16:31]
 123   ; GCN-NEXT:    ds_read_b128 v[66:69], v51 offset:512
 124   ; GCN-NEXT:    s_waitcnt lgkmcnt(0)
 125   ; GCN-NEXT:    buffer_inv sc0 sc1
 126   ; GCN-NEXT:    ;;#ASMSTART
 127   ; GCN-NEXT:    s_waitcnt vmcnt(8)
 128   ; GCN-NEXT:    ;;#ASMEND
 129   ; GCN-NEXT:    v_mfma_f32_32x32x8_f16 v[0:15], v[42:43], v[38:39], v[0:15]
 130   ; GCN-NEXT:    ; implicit-def: $vgpr32_vgpr33_vgpr34_vgpr35_vgpr36_vgpr37_vgpr38_vgpr39_vgpr40_vgpr41_vgpr42_vgpr43_vgpr44_vgpr45_vgpr46_vgpr47
 131   ; GCN-NEXT:    v_mfma_f32_32x32x8_f16 v[0:15], v[66:67], v[62:63], v[0:15]
 132   ; GCN-NEXT:    ; implicit-def: $vgpr66
 133   ; GCN-NEXT:    ; implicit-def: $vgpr67
 134   ; GCN-NEXT:    v_max_f32_e32 v81, v67, v67
 135   ; GCN-NEXT:    ; implicit-def: $vgpr48_vgpr49_vgpr50_vgpr51_vgpr52_vgpr53_vgpr54_vgpr55_vgpr56_vgpr57_vgpr58_vgpr59_vgpr60_vgpr61_vgpr62_vgpr63
 136   ; GCN-NEXT:    v_mfma_f32_32x32x8_f16 v[16:31], v[70:71], v[64:65], v[16:31]
 137   ; GCN-NEXT:    v_perm_b32 v70, v74, v72, s2
 138   ; GCN-NEXT:    v_perm_b32 v71, v74, v72, s3
 139   ; GCN-NEXT:    v_perm_b32 v72, v75, v73, s2
 140   ; GCN-NEXT:    buffer_wbl2 sc0 sc1
 141   ; GCN-NEXT:    ds_write_b32 v76, v70
 142   ; GCN-NEXT:    buffer_wbl2 sc0 sc1
 143   ; GCN-NEXT:    s_waitcnt lgkmcnt(0)
 144   ; GCN-NEXT:    ds_write_b32 v77, v71
 145   ; GCN-NEXT:    buffer_wbl2 sc0 sc1
 146   ; GCN-NEXT:    s_waitcnt lgkmcnt(0)
 147   ; GCN-NEXT:    ds_write_b32 v78, v72
 148   ; GCN-NEXT:    v_mul_f32_e32 v74, s4, v20
 149   ; GCN-NEXT:    v_mfma_f32_32x32x8_f16 v[0:15], v[68:69], v[64:65], v[0:15]
 150   ; GCN-NEXT:    v_mul_f32_e32 v64, s4, v16
 151   ; GCN-NEXT:    v_mul_f32_e32 v65, s4, v17
 152   ; GCN-NEXT:    v_mul_f32_e32 v68, s4, v18
 153   ; GCN-NEXT:    v_mul_f32_e32 v69, s4, v19
 154   ; GCN-NEXT:    v_max3_f32 v64, v64, s5, v65
 155   ; GCN-NEXT:    v_mul_f32_e32 v80, s4, v21
 156   ; GCN-NEXT:    v_max3_f32 v64, v64, v68, v69
 157   ; GCN-NEXT:    v_mul_f32_e32 v84, s4, v22
 158   ; GCN-NEXT:    v_mul_f32_e32 v85, s4, v23
 159   ; GCN-NEXT:    v_max3_f32 v64, v64, v74, v80
 160   ; GCN-NEXT:    v_mul_f32_e32 v86, s4, v24
 161   ; GCN-NEXT:    v_mul_f32_e32 v87, s4, v25
 162   ; GCN-NEXT:    v_max3_f32 v64, v64, v84, v85
 163   ; GCN-NEXT:    v_mul_f32_e32 v65, s4, v26
 164   ; GCN-NEXT:    v_mul_f32_e32 v68, s4, v27
 165   ; GCN-NEXT:    v_max3_f32 v64, v64, v86, v87
 166   ; GCN-NEXT:    v_mul_f32_e32 v69, s4, v28
 167   ; GCN-NEXT:    v_mul_f32_e32 v74, s4, v29
 168   ; GCN-NEXT:    v_max3_f32 v64, v64, v65, v68
 169   ; GCN-NEXT:    v_mul_f32_e32 v80, s4, v30
 170   ; GCN-NEXT:    v_mul_f32_e32 v84, s4, v31
 171   ; GCN-NEXT:    v_max3_f32 v64, v64, v69, v74
 172   ; GCN-NEXT:    v_mul_f32_e32 v85, s4, v0
 173   ; GCN-NEXT:    v_mul_f32_e32 v86, s4, v1
 174   ; GCN-NEXT:    v_max3_f32 v64, v64, v80, v84
 175   ; GCN-NEXT:    v_mul_f32_e32 v87, s4, v2
 176   ; GCN-NEXT:    v_mul_f32_e32 v65, s4, v3
 177   ; GCN-NEXT:    v_max3_f32 v64, v64, v85, v86
 178   ; GCN-NEXT:    v_mul_f32_e32 v68, s4, v4
 179   ; GCN-NEXT:    v_mul_f32_e32 v69, s4, v5
 180   ; GCN-NEXT:    v_max3_f32 v64, v64, v87, v65
 181   ; GCN-NEXT:    v_mul_f32_e32 v74, s4, v6
 182   ; GCN-NEXT:    v_mul_f32_e32 v80, s4, v7
 183   ; GCN-NEXT:    v_max3_f32 v64, v64, v68, v69
 184   ; GCN-NEXT:    v_mul_f32_e32 v84, s4, v8
 185   ; GCN-NEXT:    v_mul_f32_e32 v85, s4, v9
 186   ; GCN-NEXT:    v_max3_f32 v64, v64, v74, v80
 187   ; GCN-NEXT:    v_mul_f32_e32 v86, s4, v10
 188   ; GCN-NEXT:    v_mul_f32_e32 v65, s4, v11
 189   ; GCN-NEXT:    v_max3_f32 v64, v64, v84, v85
 190   ; GCN-NEXT:    v_mul_f32_e32 v87, s4, v12
 191   ; GCN-NEXT:    v_mul_f32_e32 v68, s4, v13
 192   ; GCN-NEXT:    v_max3_f32 v64, v64, v86, v65
 193   ; GCN-NEXT:    v_mul_f32_e32 v69, s4, v14
 194   ; GCN-NEXT:    v_mul_f32_e32 v74, s4, v15
 195   ; GCN-NEXT:    v_max3_f32 v64, v64, v87, v68
 196   ; GCN-NEXT:    v_max3_f32 v64, v64, v69, v74
 197   ; GCN-NEXT:    ds_bpermute_b32 v65, v66, v64
 198   ; GCN-NEXT:    v_perm_b32 v68, v75, v73, s3
 199   ; GCN-NEXT:    buffer_wbl2 sc0 sc1
 200   ; GCN-NEXT:    s_waitcnt lgkmcnt(0)
 201   ; GCN-NEXT:    ds_write_b32 v79, v68
 202   ; GCN-NEXT:    ; implicit-def: $vgpr84
 203   ; GCN-NEXT:    v_max_f32_e32 v65, v65, v65
 204   ; GCN-NEXT:    v_max_f32_e32 v70, v64, v65
 205   ; GCN-NEXT:    s_waitcnt lgkmcnt(0)
 206   ; GCN-NEXT:    buffer_load_dwordx2 v[64:65], v82, s[16:19], 0 offen sc0 sc1
 207   ; GCN-NEXT:    s_waitcnt vmcnt(0)
 208   ; GCN-NEXT:    buffer_inv sc0 sc1
 209   ; GCN-NEXT:    buffer_load_dwordx2 v[68:69], v83, s[16:19], 0 offen sc0 sc1
 210   ; GCN-NEXT:    s_waitcnt vmcnt(0)
 211   ; GCN-NEXT:    buffer_inv sc0 sc1
 212   ; GCN-NEXT:    ds_bpermute_b32 v71, v66, v70
 213   ; GCN-NEXT:    ;;#ASMSTART
 214   ; GCN-NEXT:    s_waitcnt vmcnt(8)
 215   ; GCN-NEXT:    ;;#ASMEND
 216   ; GCN-NEXT:    s_waitcnt lgkmcnt(0)
 217   ; GCN-NEXT:    v_cndmask_b32_e64 v70, v71, v70, s[0:1]
 218   ; GCN-NEXT:    v_max_f32_e32 v70, v70, v70
 219   ; GCN-NEXT:    v_max_f32_e32 v72, v81, v70
 220   ; GCN-NEXT:    v_fma_f32 v16, s4, v16, -v72
 221   ; GCN-NEXT:    v_fma_f32 v18, s4, v18, -v72
 222   ; GCN-NEXT:    v_fma_f32 v19, s4, v19, -v72
 223   ; GCN-NEXT:    v_mul_f32_e32 v16, 0x3fb8aa3b, v16
 224   ; GCN-NEXT:    v_mul_f32_e32 v18, 0x3fb8aa3b, v18
 225   ; GCN-NEXT:    v_mul_f32_e32 v19, 0x3fb8aa3b, v19
 226   ; GCN-NEXT:    v_fma_f32 v17, s4, v17, -v72
 227   ; GCN-NEXT:    v_fma_f32 v20, s4, v20, -v72
 228   ; GCN-NEXT:    v_fma_f32 v21, s4, v21, -v72
 229   ; GCN-NEXT:    v_fma_f32 v22, s4, v22, -v72
 230   ; GCN-NEXT:    v_fma_f32 v23, s4, v23, -v72
 231   ; GCN-NEXT:    v_exp_f32_e32 v73, v16
 232   ; GCN-NEXT:    v_exp_f32_e32 v74, v18
 233   ; GCN-NEXT:    v_exp_f32_e32 v75, v19
 234   ; GCN-NEXT:    v_mul_f32_e32 v20, 0x3fb8aa3b, v20
 235   ; GCN-NEXT:    v_mul_f32_e32 v21, 0x3fb8aa3b, v21
 236   ; GCN-NEXT:    v_mul_f32_e32 v22, 0x3fb8aa3b, v22
 237   ; GCN-NEXT:    v_exp_f32_e32 v80, v20
 238   ; GCN-NEXT:    v_cvt_f16_f32_e32 v16, v73
 239   ; GCN-NEXT:    v_fma_f32 v18, s4, v24, -v72
 240   ; GCN-NEXT:    v_exp_f32_e32 v81, v21
 241   ; GCN-NEXT:    v_cvt_f16_f32_e32 v21, v74
 242   ; GCN-NEXT:    v_fma_f32 v20, s4, v25, -v72
 243   ; GCN-NEXT:    v_exp_f32_e32 v82, v22
 244   ; GCN-NEXT:    v_cvt_f16_f32_e32 v22, v75
 245   ; GCN-NEXT:    v_mul_f32_e32 v17, 0x3fb8aa3b, v17
 246   ; GCN-NEXT:    v_mul_f32_e32 v23, 0x3fb8aa3b, v23
 247   ; GCN-NEXT:    v_fma_f32 v26, s4, v26, -v72
 248   ; GCN-NEXT:    v_pack_b32_f16 v71, v21, v22
 249   ; GCN-NEXT:    v_mul_f32_e32 v22, 0x3fb8aa3b, v18
 250   ; GCN-NEXT:    v_sub_f32_e32 v24, v67, v72
 251   ; GCN-NEXT:    v_exp_f32_e32 v83, v23
 252   ; GCN-NEXT:    v_fma_f32 v67, s4, v27, -v72
 253   ; GCN-NEXT:    v_exp_f32_e32 v85, v22
 254   ; GCN-NEXT:    v_exp_f32_e32 v17, v17
 255   ; GCN-NEXT:    v_mul_f32_e32 v24, 0x3fb8aa3b, v24
 256   ; GCN-NEXT:    v_mul_f32_e32 v23, 0x3fb8aa3b, v20
 257   ; GCN-NEXT:    v_cvt_f16_f32_e32 v19, v17
 258   ; GCN-NEXT:    v_fma_f32 v87, s4, v29, -v72
 259   ; GCN-NEXT:    v_exp_f32_e32 v88, v23
 260   ; GCN-NEXT:    v_fma_f32 v0, s4, v0, -v72
 261   ; GCN-NEXT:    v_pack_b32_f16 v70, v16, v19
 262   ; GCN-NEXT:    ds_read_b128 v[18:21], v84
 263   ; GCN-NEXT:    s_waitcnt lgkmcnt(0)
 264   ; GCN-NEXT:    buffer_inv sc0 sc1
 265   ; GCN-NEXT:    v_exp_f32_e32 v16, v24
 266   ; GCN-NEXT:    ds_read_b128 v[22:25], v84 offset:576
 267   ; GCN-NEXT:    s_waitcnt lgkmcnt(0)
 268   ; GCN-NEXT:    buffer_inv sc0 sc1
 269   ; GCN-NEXT:    v_pk_mul_f32 v[48:49], v[48:49], v[16:17] op_sel_hi:[1,0]
 270   ; GCN-NEXT:    v_pk_mul_f32 v[50:51], v[50:51], v[16:17] op_sel_hi:[1,0]
 271   ; GCN-NEXT:    v_pk_mul_f32 v[52:53], v[52:53], v[16:17] op_sel_hi:[1,0]
 272   ; GCN-NEXT:    v_pk_mul_f32 v[54:55], v[54:55], v[16:17] op_sel_hi:[1,0]
 273   ; GCN-NEXT:    v_pk_mul_f32 v[56:57], v[56:57], v[16:17] op_sel_hi:[1,0]
 274   ; GCN-NEXT:    v_pk_mul_f32 v[58:59], v[58:59], v[16:17] op_sel_hi:[1,0]
 275   ; GCN-NEXT:    v_pk_mul_f32 v[60:61], v[60:61], v[16:17] op_sel_hi:[1,0]
 276   ; GCN-NEXT:    v_pk_mul_f32 v[62:63], v[62:63], v[16:17] op_sel_hi:[1,0]
 277   ; GCN-NEXT:    v_pk_mul_f32 v[32:33], v[32:33], v[16:17] op_sel_hi:[1,0]
 278   ; GCN-NEXT:    v_pk_mul_f32 v[34:35], v[34:35], v[16:17] op_sel_hi:[1,0]
 279   ; GCN-NEXT:    v_pk_mul_f32 v[36:37], v[36:37], v[16:17] op_sel_hi:[1,0]
 280   ; GCN-NEXT:    v_pk_mul_f32 v[38:39], v[38:39], v[16:17] op_sel_hi:[1,0]
 281   ; GCN-NEXT:    v_pk_mul_f32 v[40:41], v[40:41], v[16:17] op_sel_hi:[1,0]
 282   ; GCN-NEXT:    v_pk_mul_f32 v[42:43], v[42:43], v[16:17] op_sel_hi:[1,0]
 283   ; GCN-NEXT:    v_pk_mul_f32 v[44:45], v[44:45], v[16:17] op_sel_hi:[1,0]
 284   ; GCN-NEXT:    v_pk_mul_f32 v[46:47], v[46:47], v[16:17] op_sel_hi:[1,0]
 285   ; GCN-NEXT:    v_mfma_f32_32x32x8_f16 v[48:63], v[18:19], v[70:71], v[48:63]
 286   ; GCN-NEXT:    v_add_f32_e32 v18, 0, v73
 287   ; GCN-NEXT:    v_cvt_f16_f32_e32 v89, v83
 288   ; GCN-NEXT:    v_fma_f32 v73, s4, v28, -v72
 289   ; GCN-NEXT:    v_cvt_f16_f32_e32 v19, v80
 290   ; GCN-NEXT:    v_fma_f32 v1, s4, v1, -v72
 291   ; GCN-NEXT:    v_perm_b32 v90, v69, v65, s2
 292   ; GCN-NEXT:    v_mfma_f32_32x32x8_f16 v[32:47], v[22:23], v[70:71], v[32:47]
 293   ; GCN-NEXT:    v_add_f32_e32 v17, v17, v18
 294   ; GCN-NEXT:    v_mul_f32_e32 v18, 0x3fb8aa3b, v26
 295   ; GCN-NEXT:    v_cvt_f16_f32_e32 v86, v81
 296   ; GCN-NEXT:    v_fma_f32 v23, s4, v30, -v72
 297   ; GCN-NEXT:    v_exp_f32_e32 v30, v18
 298   ; GCN-NEXT:    v_cvt_f16_f32_e32 v22, v82
 299   ; GCN-NEXT:    v_fma_f32 v18, s4, v31, -v72
 300   ; GCN-NEXT:    v_perm_b32 v31, v68, v64, s2
 301   ; GCN-NEXT:    v_perm_b32 v64, v68, v64, s3
 302   ; GCN-NEXT:    v_perm_b32 v65, v69, v65, s3
 303   ; GCN-NEXT:    ds_read_b128 v[26:29], v91
 304   ; GCN-NEXT:    s_waitcnt lgkmcnt(0)
 305   ; GCN-NEXT:    buffer_inv sc0 sc1
 306   ; GCN-NEXT:    ds_read_b128 v[68:71], v91 offset:576
 307   ; GCN-NEXT:    s_waitcnt lgkmcnt(0)
 308   ; GCN-NEXT:    buffer_inv sc0 sc1
 309   ; GCN-NEXT:    ;;#ASMSTART
 310   ; GCN-NEXT:    s_waitcnt vmcnt(8)
 311   ; GCN-NEXT:    ;;#ASMEND
 312   ; GCN-NEXT:    buffer_wbl2 sc0 sc1
 313   ; GCN-NEXT:    ds_write_b32 v76, v31
 314   ; GCN-NEXT:    v_mul_f32_e32 v31, 0x3fb8aa3b, v67
 315   ; GCN-NEXT:    v_exp_f32_e32 v31, v31
 316   ; GCN-NEXT:    v_mul_f32_e32 v67, 0x3fb8aa3b, v18
 317   ; GCN-NEXT:    v_pack_b32_f16 v18, v19, v86
 318   ; GCN-NEXT:    v_pack_b32_f16 v19, v22, v89
 319   ; GCN-NEXT:    buffer_wbl2 sc0 sc1
 320   ; GCN-NEXT:    s_waitcnt lgkmcnt(0)
 321   ; GCN-NEXT:    ds_write_b32 v77, v64
 322   ; GCN-NEXT:    buffer_wbl2 sc0 sc1
 323   ; GCN-NEXT:    s_waitcnt lgkmcnt(0)
 324   ; GCN-NEXT:    ds_write_b32 v78, v90
 325   ; GCN-NEXT:    buffer_wbl2 sc0 sc1
 326   ; GCN-NEXT:    s_waitcnt lgkmcnt(0)
 327   ; GCN-NEXT:    ds_write_b32 v79, v65
 328   ; GCN-NEXT:    v_mul_f32_e32 v64, 0x3fb8aa3b, v73
 329   ; GCN-NEXT:    v_mul_f32_e32 v65, 0x3fb8aa3b, v87
 330   ; GCN-NEXT:    v_mfma_f32_32x32x8_f16 v[48:63], v[20:21], v[18:19], v[48:63]
 331   ; GCN-NEXT:    v_add_f32_e32 v17, v74, v17
 332   ; GCN-NEXT:    v_cvt_f16_f32_e32 v20, v85
 333   ; GCN-NEXT:    v_fma_f32 v2, s4, v2, -v72
 334   ; GCN-NEXT:    v_exp_f32_e32 v22, v64
 335   ; GCN-NEXT:    v_cvt_f16_f32_e32 v21, v88
 336   ; GCN-NEXT:    v_exp_f32_e32 v64, v65
 337   ; GCN-NEXT:    v_mul_f32_e32 v23, 0x3fb8aa3b, v23
 338   ; GCN-NEXT:    v_mfma_f32_32x32x8_f16 v[32:47], v[24:25], v[18:19], v[32:47]
 339   ; GCN-NEXT:    v_add_f32_e32 v17, v75, v17
 340   ; GCN-NEXT:    v_cvt_f16_f32_e32 v18, v30
 341   ; GCN-NEXT:    v_fma_f32 v24, s4, v3, -v72
 342   ; GCN-NEXT:    v_exp_f32_e32 v23, v23
 343   ; GCN-NEXT:    v_cvt_f16_f32_e32 v19, v31
 344   ; GCN-NEXT:    v_mul_f32_e32 v3, 0x3fb8aa3b, v0
 345   ; GCN-NEXT:    v_mul_f32_e32 v65, 0x3fb8aa3b, v1
 346   ; GCN-NEXT:    v_pack_b32_f16 v0, v20, v21
 347   ; GCN-NEXT:    v_pack_b32_f16 v1, v18, v19
 348   ; GCN-NEXT:    v_fma_f32 v6, s4, v6, -v72
 349   ; GCN-NEXT:    v_exp_f32_e32 v25, v67
 350   ; GCN-NEXT:    v_mfma_f32_32x32x8_f16 v[48:63], v[26:27], v[0:1], v[48:63]
 351   ; GCN-NEXT:    v_add_f32_e32 v17, v80, v17
 352   ; GCN-NEXT:    v_cvt_f16_f32_e32 v18, v22
 353   ; GCN-NEXT:    v_fma_f32 v26, s4, v4, -v72
 354   ; GCN-NEXT:    v_exp_f32_e32 v27, v3
 355   ; GCN-NEXT:    v_cvt_f16_f32_e32 v4, v64
 356   ; GCN-NEXT:    v_fma_f32 v67, s4, v5, -v72
 357   ; GCN-NEXT:    v_exp_f32_e32 v65, v65
 358   ; GCN-NEXT:    v_mfma_f32_32x32x8_f16 v[32:47], v[68:69], v[0:1], v[32:47]
 359   ; GCN-NEXT:    v_mul_f32_e32 v2, 0x3fb8aa3b, v2
 360   ; GCN-NEXT:    v_add_f32_e32 v17, v81, v17
 361   ; GCN-NEXT:    v_cvt_f16_f32_e32 v5, v23
 362   ; GCN-NEXT:    v_fma_f32 v7, s4, v7, -v72
 363   ; GCN-NEXT:    v_exp_f32_e32 v68, v2
 364   ; GCN-NEXT:    v_cvt_f16_f32_e32 v19, v25
 365   ; GCN-NEXT:    ;;#ASMSTART
 366   ; GCN-NEXT:    s_waitcnt vmcnt(8)
 367   ; GCN-NEXT:    ;;#ASMEND
 368   ; GCN-NEXT:    v_mul_f32_e32 v24, 0x3fb8aa3b, v24
 369   ; GCN-NEXT:    s_waitcnt lgkmcnt(0)
 370   ; GCN-NEXT:    ds_read_b128 v[0:3], v84
 371   ; GCN-NEXT:    s_waitcnt lgkmcnt(0)
 372   ; GCN-NEXT:    buffer_inv sc0 sc1
 373   ; GCN-NEXT:    v_pack_b32_f16 v4, v18, v4
 374   ; GCN-NEXT:    v_pack_b32_f16 v5, v5, v19
 375   ; GCN-NEXT:    v_exp_f32_e32 v24, v24
 376   ; GCN-NEXT:    ds_read_b128 v[18:21], v84 offset:576
 377   ; GCN-NEXT:    s_waitcnt lgkmcnt(0)
 378   ; GCN-NEXT:    buffer_inv sc0 sc1
 379   ; GCN-NEXT:    v_mul_f32_e32 v26, 0x3fb8aa3b, v26
 380   ; GCN-NEXT:    v_mul_f32_e32 v67, 0x3fb8aa3b, v67
 381   ; GCN-NEXT:    v_mfma_f32_32x32x8_f16 v[48:63], v[28:29], v[4:5], v[48:63]
 382   ; GCN-NEXT:    v_add_f32_e32 v17, v82, v17
 383   ; GCN-NEXT:    v_cvt_f16_f32_e32 v28, v27
 384   ; GCN-NEXT:    v_exp_f32_e32 v26, v26
 385   ; GCN-NEXT:    v_cvt_f16_f32_e32 v29, v65
 386   ; GCN-NEXT:    v_fma_f32 v10, s4, v10, -v72
 387   ; GCN-NEXT:    v_exp_f32_e32 v67, v67
 388   ; GCN-NEXT:    v_mul_f32_e32 v6, 0x3fb8aa3b, v6
 389   ; GCN-NEXT:    v_mfma_f32_32x32x8_f16 v[32:47], v[70:71], v[4:5], v[32:47]
 390   ; GCN-NEXT:    v_add_f32_e32 v17, v83, v17
 391   ; GCN-NEXT:    v_cvt_f16_f32_e32 v5, v68
 392   ; GCN-NEXT:    v_exp_f32_e32 v6, v6
 393   ; GCN-NEXT:    v_cvt_f16_f32_e32 v69, v24
 394   ; GCN-NEXT:    v_mul_f32_e32 v7, 0x3fb8aa3b, v7
 395   ; GCN-NEXT:    v_exp_f32_e32 v7, v7
 396   ; GCN-NEXT:    v_pack_b32_f16 v4, v28, v29
 397   ; GCN-NEXT:    v_pack_b32_f16 v5, v5, v69
 398   ; GCN-NEXT:    ; implicit-def: $sgpr2
 399   ; GCN-NEXT:    s_nop 1
 400   ; GCN-NEXT:    v_mfma_f32_32x32x8_f16 v[48:63], v[0:1], v[4:5], v[48:63]
 401   ; GCN-NEXT:    v_add_f32_e32 v0, v85, v17
 402   ; GCN-NEXT:    v_cvt_f16_f32_e32 v17, v26
 403   ; GCN-NEXT:    v_cvt_f16_f32_e32 v28, v67
 404   ; GCN-NEXT:    v_mfma_f32_32x32x8_f16 v[32:47], v[18:19], v[4:5], v[32:47]
 405   ; GCN-NEXT:    v_add_f32_e32 v4, v88, v0
 406   ; GCN-NEXT:    v_mul_f32_e32 v0, 0x3fb8aa3b, v10
 407   ; GCN-NEXT:    v_cvt_f16_f32_e32 v1, v6
 408   ; GCN-NEXT:    v_exp_f32_e32 v10, v0
 409   ; GCN-NEXT:    v_cvt_f16_f32_e32 v0, v7
 410   ; GCN-NEXT:    v_pack_b32_f16 v1, v1, v0
 411   ; GCN-NEXT:    v_pack_b32_f16 v0, v17, v28
 412   ; GCN-NEXT:    s_nop 1
 413   ; GCN-NEXT:    v_mfma_f32_32x32x8_f16 v[48:63], v[2:3], v[0:1], v[48:63]
 414   ; GCN-NEXT:    v_add_f32_e32 v2, v30, v4
 415   ; GCN-NEXT:    v_mfma_f32_32x32x8_f16 v[32:47], v[20:21], v[0:1], v[32:47]
 416   ; GCN-NEXT:    v_add_f32_e32 v0, v31, v2
 417   ; GCN-NEXT:    v_add_f32_e32 v0, v22, v0
 418   ; GCN-NEXT:    v_add_f32_e32 v0, v64, v0
 419   ; GCN-NEXT:    v_add_f32_e32 v0, v23, v0
 420   ; GCN-NEXT:    v_add_f32_e32 v0, v25, v0
 421   ; GCN-NEXT:    v_add_f32_e32 v0, v27, v0
 422   ; GCN-NEXT:    v_fma_f32 v8, s4, v8, -v72
 423   ; GCN-NEXT:    v_add_f32_e32 v0, v65, v0
 424   ; GCN-NEXT:    v_fma_f32 v9, s4, v9, -v72
 425   ; GCN-NEXT:    v_mul_f32_e32 v8, 0x3fb8aa3b, v8
 426   ; GCN-NEXT:    v_add_f32_e32 v0, v68, v0
 427   ; GCN-NEXT:    v_fma_f32 v11, s4, v11, -v72
 428   ; GCN-NEXT:    v_mul_f32_e32 v9, 0x3fb8aa3b, v9
 429   ; GCN-NEXT:    v_fma_f32 v12, s4, v12, -v72
 430   ; GCN-NEXT:    v_fma_f32 v13, s4, v13, -v72
 431   ; GCN-NEXT:    v_exp_f32_e32 v8, v8
 432   ; GCN-NEXT:    v_add_f32_e32 v0, v24, v0
 433   ; GCN-NEXT:    v_fma_f32 v5, s4, v14, -v72
 434   ; GCN-NEXT:    v_exp_f32_e32 v9, v9
 435   ; GCN-NEXT:    v_add_f32_e32 v0, v26, v0
 436   ; GCN-NEXT:    v_add_f32_e32 v0, v67, v0
 437   ; GCN-NEXT:    v_fma_f32 v14, s4, v15, -v72
 438   ; GCN-NEXT:    v_mul_f32_e32 v11, 0x3fb8aa3b, v11
 439   ; GCN-NEXT:    v_mul_f32_e32 v3, 0x3fb8aa3b, v12
 440   ; GCN-NEXT:    v_mul_f32_e32 v1, 0x3fb8aa3b, v5
 441   ; GCN-NEXT:    v_add_f32_e32 v0, v6, v0
 442   ; GCN-NEXT:    v_exp_f32_e32 v11, v11
 443   ; GCN-NEXT:    v_cvt_f16_f32_e32 v4, v8
 444   ; GCN-NEXT:    v_exp_f32_e32 v12, v3
 445   ; GCN-NEXT:    v_mul_f32_e32 v3, 0x3fb8aa3b, v13
 446   ; GCN-NEXT:    v_exp_f32_e32 v17, v1
 447   ; GCN-NEXT:    v_mul_f32_e32 v1, 0x3fb8aa3b, v14
 448   ; GCN-NEXT:    v_add_f32_e32 v0, v7, v0
 449   ; GCN-NEXT:    v_cvt_f16_f32_e32 v13, v9
 450   ; GCN-NEXT:    v_exp_f32_e32 v15, v3
 451   ; GCN-NEXT:    v_exp_f32_e32 v18, v1
 452   ; GCN-NEXT:    v_add_f32_e32 v6, v8, v0
 453   ; GCN-NEXT:    ds_read_b128 v[0:3], v91
 454   ; GCN-NEXT:    s_waitcnt lgkmcnt(0)
 455   ; GCN-NEXT:    buffer_inv sc0 sc1
 456   ; GCN-NEXT:    v_cvt_f16_f32_e32 v5, v10
 457   ; GCN-NEXT:    v_cvt_f16_f32_e32 v14, v11
 458   ; GCN-NEXT:    v_add_f32_e32 v6, v9, v6
 459   ; GCN-NEXT:    v_pack_b32_f16 v8, v4, v13
 460   ; GCN-NEXT:    v_add_f32_e32 v6, v10, v6
 461   ; GCN-NEXT:    v_pack_b32_f16 v9, v5, v14
 462   ; GCN-NEXT:    v_cvt_f16_f32_e32 v7, v18
 463   ; GCN-NEXT:    v_cvt_f16_f32_e32 v10, v15
 464   ; GCN-NEXT:    v_mfma_f32_32x32x8_f16 v[48:63], v[0:1], v[8:9], v[48:63]
 465   ; GCN-NEXT:    v_cvt_f16_f32_e32 v0, v17
 466   ; GCN-NEXT:    v_cvt_f16_f32_e32 v4, v12
 467   ; GCN-NEXT:    v_add_f32_e32 v6, v11, v6
 468   ; GCN-NEXT:    v_add_f32_e32 v6, v12, v6
 469   ; GCN-NEXT:    v_add_f32_e32 v1, v15, v6
 470   ; GCN-NEXT:    v_add_f32_e32 v11, v17, v1
 471   ; GCN-NEXT:    v_pack_b32_f16 v1, v0, v7
 472   ; GCN-NEXT:    v_pack_b32_f16 v0, v4, v10
 473   ; GCN-NEXT:    ds_read_b128 v[4:7], v91 offset:576
 474   ; GCN-NEXT:    s_waitcnt lgkmcnt(0)
 475   ; GCN-NEXT:    buffer_inv sc0 sc1
 476   ; GCN-NEXT:    v_mfma_f32_32x32x8_f16 v[32:47], v[4:5], v[8:9], v[32:47]
 477   ; GCN-NEXT:    ;;#ASMSTART
 478   ; GCN-NEXT:    s_waitcnt vmcnt(8)
 479   ; GCN-NEXT:    ;;#ASMEND
 480   ; GCN-NEXT:    v_mov_b32_e32 v4, 0
 481   ; GCN-NEXT:    v_mfma_f32_32x32x8_f16 v[32:47], v[6:7], v[0:1], v[32:47]
 482   ; GCN-NEXT:    v_mfma_f32_32x32x8_f16 v[48:63], v[2:3], v[0:1], v[48:63]
 483   ; GCN-NEXT:    v_add_f32_e32 v2, v18, v11
 484   ; GCN-NEXT:    ds_bpermute_b32 v3, v66, v2
 485   ; GCN-NEXT:    s_waitcnt lgkmcnt(0)
 486   ; GCN-NEXT:    v_add_f32_e32 v2, v2, v3
 487   ; GCN-NEXT:    ds_bpermute_b32 v3, v66, v2
 488   ; GCN-NEXT:    s_waitcnt lgkmcnt(0)
 489   ; GCN-NEXT:    v_cndmask_b32_e64 v2, v3, v2, s[0:1]
 490   ; GCN-NEXT:    v_fmac_f32_e32 v2, v4, v16
 491   ; GCN-NEXT:    s_endpgm
 492   attributes #0 = {"amdgpu-flat-work-group-size"="256,256"}
 493
 494   !0 = !{i64 2862105}
 495 ...
 496
 497 ---
 498 name:            smallInterleave
 499 tracksRegLiveness: true
 500 machineFunctionInfo:
 501   stackPtrOffsetReg: '$sgpr32'
 502 body:             |
 503  bb.0:
 504   liveins: $vgpr0, $sgpr0_sgpr1, $sgpr2, $sgpr3, $sgpr4
 505   %0:vgpr_32 = V_MOV_B32_e32 0, implicit $exec
 506   %1:vgpr_32 = COPY %0:vgpr_32
 507   %2:vgpr_32 = IMPLICIT_DEF
 508   %3:sreg_32 = IMPLICIT_DEF
 509   %4:vreg_64_align2 = IMPLICIT_DEF
 510   %5:sgpr_128 = IMPLICIT_DEF
 511   %6:vgpr_32 = IMPLICIT_DEF
 512   %7:vgpr_32 = IMPLICIT_DEF
 513   %8:sgpr_128 = IMPLICIT_DEF
 514   %9:vgpr_32 = IMPLICIT_DEF
 515   %10:sgpr_512 = IMPLICIT_DEF
 516   %11:sgpr_32 = IMPLICIT_DEF
 517   %12:sreg_64_xexec = IMPLICIT_DEF
 518   %13:vgpr_32 = IMPLICIT_DEF
 519   %14:sreg_32 = IMPLICIT_DEF
 520   %15:sreg_32 = IMPLICIT_DEF
 521   %16:vgpr_32 = IMPLICIT_DEF
 522   %17:sreg_32 = IMPLICIT_DEF
 523   %18:vgpr_32 = IMPLICIT_DEF
 524   %19:vgpr_32 = IMPLICIT_DEF
 525   %20:vgpr_32 = IMPLICIT_DEF
 526   %21:vgpr_32 = IMPLICIT_DEF
 527   %22:vgpr_32 = IMPLICIT_DEF
 528   %23:vgpr_32 = IMPLICIT_DEF
 529   %24:vgpr_32 = IMPLICIT_DEF
 530   %25:vgpr_32 = IMPLICIT_DEF
 531   %26:sreg_32 = IMPLICIT_DEF
 532   %42:vgpr_32 = IMPLICIT_DEF
 533   %44:vreg_128_align2 = IMPLICIT_DEF
 534   %48:vgpr_32 = IMPLICIT_DEF
 535   %49:vreg_128_align2 = IMPLICIT_DEF
 536   %52:vreg_128_align2 = IMPLICIT_DEF
 537   %55:vreg_128_align2 = IMPLICIT_DEF
 538   %106:vgpr_32 = IMPLICIT_DEF
 539   %29:vgpr_32 = IMPLICIT_DEF
 540   %37:vgpr_32 = IMPLICIT_DEF
 541   %259:vreg_512_align2 = IMPLICIT_DEF
 542   %260:vreg_512_align2 = IMPLICIT_DEF
 543   IGLP_OPT 2
 544   %27:sreg_32 = V_READFIRSTLANE_B32 %2:vgpr_32, implicit $exec
 545   %28:vgpr_32 = V_LSHL_ADD_U32_e64 %27:sreg_32, 4, %29:vgpr_32, implicit $exec
 546   %30:vreg_64_align2, dead %31:sreg_64 = V_MAD_U64_U32_e64 %3:sreg_32, %28:vgpr_32, %4:vreg_64_align2, 0, implicit $exec
 547   %32:vreg_128_align2 = BUFFER_LOAD_DWORDX4_OFFEN %30.sub0:vreg_64_align2, %5:sgpr_128, 0, 0, 0, 0, implicit $exec
 548   %33:sreg_32 = S_LSHL_B32 %27:sreg_32, 7, implicit-def dead $scc
 549   %34:vgpr_32 = V_ADD_LSHL_U32_e64 %6:vgpr_32, %33:sreg_32, 1, implicit $exec
 550   DS_WRITE_B128_gfx9 %34:vgpr_32, %32:vreg_128_align2, 0, 0, implicit $exec
 551   %35:vreg_128_align2 = BUFFER_LOAD_DWORDX4_OFFEN %30.sub0:vreg_64_align2, %5:sgpr_128, 0, 64, 0, 0, implicit $exec
 552   %36:vgpr_32 = V_ADD_U32_e32 %7:vgpr_32, %37:vgpr_32, implicit $exec
 553   %38:vreg_64_align2 = BUFFER_LOAD_DWORDX2_OFFEN %36:vgpr_32, %8:sgpr_128, 0, 0, 0, 0, implicit $exec
 554   %39:vgpr_32 = V_ADD_U32_e32 %9:vgpr_32, %37:vgpr_32, implicit $exec
 555   %40:vreg_64_align2 = BUFFER_LOAD_DWORDX2_OFFEN %39:vgpr_32, %8:sgpr_128, 0, 0, 0, 0, implicit $exec
 556   INLINEASM &"s_waitcnt vmcnt($0)", 57 /* sideeffect mayload maystore isconvergent attdialect */, 13 /* imm */, 8, !0
 557   %41:vreg_128_align2 = DS_READ_B128_gfx9 %42:vgpr_32, 0, 0, implicit $exec
 558   early-clobber %43:vreg_512_align2 = contract V_MFMA_F32_32X32X8F16_vgprcd_e64 %41.sub0_sub1:vreg_128_align2, %44.sub0_sub1:vreg_128_align2, 0, 0, 0, 0, implicit $mode, implicit $exec
 559   %43:vreg_512_align2 = contract V_MFMA_F32_32X32X8F16_mac_vgprcd_e64 %41.sub2_sub3:vreg_128_align2, %44.sub2_sub3:vreg_128_align2, %43:vreg_512_align2, 0, 0, 0, implicit $mode, implicit $exec
 560   %45:vreg_128_align2 = DS_READ_B128_gfx9 %42:vgpr_32, 512, 0, implicit $exec
 561   early-clobber %46:vreg_512_align2 = contract V_MFMA_F32_32X32X8F16_vgprcd_e64 %45.sub0_sub1:vreg_128_align2, %44.sub0_sub1:vreg_128_align2, 0, 0, 0, 0, implicit $mode, implicit $exec
 562   %46:vreg_512_align2 = contract V_MFMA_F32_32X32X8F16_mac_vgprcd_e64 %45.sub2_sub3:vreg_128_align2, %44.sub2_sub3:vreg_128_align2, %46:vreg_512_align2, 0, 0, 0, implicit $mode, implicit $exec
 563   %47:vreg_128_align2 = DS_READ_B128_gfx9 %48:vgpr_32, 0, 0, implicit $exec
 564   %43:vreg_512_align2 = contract V_MFMA_F32_32X32X8F16_mac_vgprcd_e64 %47.sub0_sub1:vreg_128_align2, %49.sub0_sub1:vreg_128_align2, %43:vreg_512_align2, 0, 0, 0, implicit $mode, implicit $exec
 565   %43:vreg_512_align2 = contract V_MFMA_F32_32X32X8F16_mac_vgprcd_e64 %47.sub2_sub3:vreg_128_align2, %49.sub2_sub3:vreg_128_align2, %43:vreg_512_align2, 0, 0, 0, implicit $mode, implicit $exec
 566   %50:vreg_128_align2 = DS_READ_B128_gfx9 %48:vgpr_32, 512, 0, implicit $exec
 567   %46:vreg_512_align2 = contract V_MFMA_F32_32X32X8F16_mac_vgprcd_e64 %50.sub0_sub1:vreg_128_align2, %49.sub0_sub1:vreg_128_align2, %46:vreg_512_align2, 0, 0, 0, implicit $mode, implicit $exec
 568   %46:vreg_512_align2 = contract V_MFMA_F32_32X32X8F16_mac_vgprcd_e64 %50.sub2_sub3:vreg_128_align2, %49.sub2_sub3:vreg_128_align2, %46:vreg_512_align2, 0, 0, 0, implicit $mode, implicit $exec
 569   INLINEASM &"s_waitcnt vmcnt($0)", 57 /* sideeffect mayload maystore isconvergent attdialect */, 13 /* imm */, 8, !0
 570   DS_WRITE_B128_gfx9 %34:vgpr_32, %35:vreg_128_align2, 0, 0, implicit $exec
 571   INLINEASM &"s_waitcnt vmcnt($0)", 57 /* sideeffect mayload maystore isconvergent attdialect */, 13 /* imm */, 8, !0
 572   %51:vreg_128_align2 = DS_READ_B128_gfx9 %42:vgpr_32, 0, 0, implicit $exec
 573   %43:vreg_512_align2 = contract V_MFMA_F32_32X32X8F16_mac_vgprcd_e64 %51.sub0_sub1:vreg_128_align2, %52.sub0_sub1:vreg_128_align2, %43:vreg_512_align2, 0, 0, 0, implicit $mode, implicit $exec
 574   %43:vreg_512_align2 = contract V_MFMA_F32_32X32X8F16_mac_vgprcd_e64 %51.sub2_sub3:vreg_128_align2, %52.sub2_sub3:vreg_128_align2, %43:vreg_512_align2, 0, 0, 0, implicit $mode, implicit $exec
 575   %53:vreg_128_align2 = DS_READ_B128_gfx9 %42:vgpr_32, 512, 0, implicit $exec
 576   %46:vreg_512_align2 = contract V_MFMA_F32_32X32X8F16_mac_vgprcd_e64 %53.sub0_sub1:vreg_128_align2, %52.sub0_sub1:vreg_128_align2, %46:vreg_512_align2, 0, 0, 0, implicit $mode, implicit $exec
 577   %46:vreg_512_align2 = contract V_MFMA_F32_32X32X8F16_mac_vgprcd_e64 %53.sub2_sub3:vreg_128_align2, %52.sub2_sub3:vreg_128_align2, %46:vreg_512_align2, 0, 0, 0, implicit $mode, implicit $exec
 578   %54:vreg_128_align2 = DS_READ_B128_gfx9 %48:vgpr_32, 0, 0, implicit $exec
 579   %43:vreg_512_align2 = contract V_MFMA_F32_32X32X8F16_mac_vgprcd_e64 %54.sub0_sub1:vreg_128_align2, %55.sub0_sub1:vreg_128_align2, %43:vreg_512_align2, 0, 0, 0, implicit $mode, implicit $exec
 580   %43:vreg_512_align2 = contract V_MFMA_F32_32X32X8F16_mac_vgprcd_e64 %54.sub2_sub3:vreg_128_align2, %55.sub2_sub3:vreg_128_align2, %43:vreg_512_align2, 0, 0, 0, implicit $mode, implicit $exec
 581   %56:vreg_128_align2 = DS_READ_B128_gfx9 %48:vgpr_32, 512, 0, implicit $exec
 582   %46:vreg_512_align2 = contract V_MFMA_F32_32X32X8F16_mac_vgprcd_e64 %56.sub0_sub1:vreg_128_align2, %55.sub0_sub1:vreg_128_align2, %46:vreg_512_align2, 0, 0, 0, implicit $mode, implicit $exec
 583   %46:vreg_512_align2 = contract V_MFMA_F32_32X32X8F16_mac_vgprcd_e64 %56.sub2_sub3:vreg_128_align2, %55.sub2_sub3:vreg_128_align2, %46:vreg_512_align2, 0, 0, 0, implicit $mode, implicit $exec
 584   %57:vgpr_32 = contract nofpexcept V_MUL_F32_e32 %10.sub4:sgpr_512, %43.sub0:vreg_512_align2, implicit $mode, implicit $exec
 585   %58:vgpr_32 = contract nofpexcept V_MUL_F32_e32 %10.sub4:sgpr_512, %43.sub1:vreg_512_align2, implicit $mode, implicit $exec
 586   %59:vgpr_32 = contract nofpexcept V_MUL_F32_e32 %10.sub4:sgpr_512, %43.sub2:vreg_512_align2, implicit $mode, implicit $exec
 587   %60:vgpr_32 = contract nofpexcept V_MUL_F32_e32 %10.sub4:sgpr_512, %43.sub3:vreg_512_align2, implicit $mode, implicit $exec
 588   %61:vgpr_32 = contract nofpexcept V_MUL_F32_e32 %10.sub4:sgpr_512, %43.sub4:vreg_512_align2, implicit $mode, implicit $exec
 589   %62:vgpr_32 = contract nofpexcept V_MUL_F32_e32 %10.sub4:sgpr_512, %43.sub5:vreg_512_align2, implicit $mode, implicit $exec
 590   %63:vgpr_32 = contract nofpexcept V_MUL_F32_e32 %10.sub4:sgpr_512, %43.sub6:vreg_512_align2, implicit $mode, implicit $exec
 591   %64:vgpr_32 = contract nofpexcept V_MUL_F32_e32 %10.sub4:sgpr_512, %43.sub7:vreg_512_align2, implicit $mode, implicit $exec
 592   %65:vgpr_32 = contract nofpexcept V_MUL_F32_e32 %10.sub4:sgpr_512, %43.sub8:vreg_512_align2, implicit $mode, implicit $exec
 593   %66:vgpr_32 = contract nofpexcept V_MUL_F32_e32 %10.sub4:sgpr_512, %43.sub9:vreg_512_align2, implicit $mode, implicit $exec
 594   %67:vgpr_32 = contract nofpexcept V_MUL_F32_e32 %10.sub4:sgpr_512, %43.sub10:vreg_512_align2, implicit $mode, implicit $exec
 595   %68:vgpr_32 = contract nofpexcept V_MUL_F32_e32 %10.sub4:sgpr_512, %43.sub11:vreg_512_align2, implicit $mode, implicit $exec
 596   %69:vgpr_32 = contract nofpexcept V_MUL_F32_e32 %10.sub4:sgpr_512, %43.sub12:vreg_512_align2, implicit $mode, implicit $exec
 597   %70:vgpr_32 = contract nofpexcept V_MUL_F32_e32 %10.sub4:sgpr_512, %43.sub13:vreg_512_align2, implicit $mode, implicit $exec
 598   %71:vgpr_32 = contract nofpexcept V_MUL_F32_e32 %10.sub4:sgpr_512, %43.sub14:vreg_512_align2, implicit $mode, implicit $exec
 599   %72:vgpr_32 = contract nofpexcept V_MUL_F32_e32 %10.sub4:sgpr_512, %43.sub15:vreg_512_align2, implicit $mode, implicit $exec
 600   %73:vgpr_32 = contract nofpexcept V_MUL_F32_e32 %10.sub4:sgpr_512, %46.sub0:vreg_512_align2, implicit $mode, implicit $exec
 601   %74:vgpr_32 = contract nofpexcept V_MUL_F32_e32 %10.sub4:sgpr_512, %46.sub1:vreg_512_align2, implicit $mode, implicit $exec
 602   %75:vgpr_32 = contract nofpexcept V_MUL_F32_e32 %10.sub4:sgpr_512, %46.sub2:vreg_512_align2, implicit $mode, implicit $exec
 603   %76:vgpr_32 = contract nofpexcept V_MUL_F32_e32 %10.sub4:sgpr_512, %46.sub3:vreg_512_align2, implicit $mode, implicit $exec
 604   %77:vgpr_32 = contract nofpexcept V_MUL_F32_e32 %10.sub4:sgpr_512, %46.sub4:vreg_512_align2, implicit $mode, implicit $exec
 605   %78:vgpr_32 = contract nofpexcept V_MUL_F32_e32 %10.sub4:sgpr_512, %46.sub5:vreg_512_align2, implicit $mode, implicit $exec
 606   %79:vgpr_32 = contract nofpexcept V_MUL_F32_e32 %10.sub4:sgpr_512, %46.sub6:vreg_512_align2, implicit $mode, implicit $exec
 607   %80:vgpr_32 = contract nofpexcept V_MUL_F32_e32 %10.sub4:sgpr_512, %46.sub7:vreg_512_align2, implicit $mode, implicit $exec
 608   %81:vgpr_32 = contract nofpexcept V_MUL_F32_e32 %10.sub4:sgpr_512, %46.sub8:vreg_512_align2, implicit $mode, implicit $exec
 609   %82:vgpr_32 = contract nofpexcept V_MUL_F32_e32 %10.sub4:sgpr_512, %46.sub9:vreg_512_align2, implicit $mode, implicit $exec
 610   %83:vgpr_32 = contract nofpexcept V_MUL_F32_e32 %10.sub4:sgpr_512, %46.sub10:vreg_512_align2, implicit $mode, implicit $exec
 611   %84:vgpr_32 = contract nofpexcept V_MUL_F32_e32 %10.sub4:sgpr_512, %46.sub11:vreg_512_align2, implicit $mode, implicit $exec
 612   %85:vgpr_32 = contract nofpexcept V_MUL_F32_e32 %10.sub4:sgpr_512, %46.sub12:vreg_512_align2, implicit $mode, implicit $exec
 613   %86:vgpr_32 = contract nofpexcept V_MUL_F32_e32 %10.sub4:sgpr_512, %46.sub13:vreg_512_align2, implicit $mode, implicit $exec
 614   %87:vgpr_32 = contract nofpexcept V_MUL_F32_e32 %10.sub4:sgpr_512, %46.sub14:vreg_512_align2, implicit $mode, implicit $exec
 615   %88:vgpr_32 = contract nofpexcept V_MUL_F32_e32 %10.sub4:sgpr_512, %46.sub15:vreg_512_align2, implicit $mode, implicit $exec
 616   %89:vgpr_32 = V_MAX3_F32_e64 0, %57:vgpr_32, 0, %11:sgpr_32, 0, %58:vgpr_32, 0, 0, implicit $mode, implicit $exec
 617   %90:vgpr_32 = V_MAX3_F32_e64 0, %89:vgpr_32, 0, %59:vgpr_32, 0, %60:vgpr_32, 0, 0, implicit $mode, implicit $exec
 618   %91:vgpr_32 = V_MAX3_F32_e64 0, %90:vgpr_32, 0, %61:vgpr_32, 0, %62:vgpr_32, 0, 0, implicit $mode, implicit $exec
 619   %92:vgpr_32 = V_MAX3_F32_e64 0, %91:vgpr_32, 0, %63:vgpr_32, 0, %64:vgpr_32, 0, 0, implicit $mode, implicit $exec
 620   %93:vgpr_32 = V_MAX3_F32_e64 0, %92:vgpr_32, 0, %65:vgpr_32, 0, %66:vgpr_32, 0, 0, implicit $mode, implicit $exec
 621   %94:vgpr_32 = V_MAX3_F32_e64 0, %93:vgpr_32, 0, %67:vgpr_32, 0, %68:vgpr_32, 0, 0, implicit $mode, implicit $exec
 622   %95:vgpr_32 = V_MAX3_F32_e64 0, %94:vgpr_32, 0, %69:vgpr_32, 0, %70:vgpr_32, 0, 0, implicit $mode, implicit $exec
 623   %96:vgpr_32 = V_MAX3_F32_e64 0, %95:vgpr_32, 0, %71:vgpr_32, 0, %72:vgpr_32, 0, 0, implicit $mode, implicit $exec
 624   %97:vgpr_32 = V_MAX3_F32_e64 0, %96:vgpr_32, 0, %73:vgpr_32, 0, %74:vgpr_32, 0, 0, implicit $mode, implicit $exec
 625   %98:vgpr_32 = V_MAX3_F32_e64 0, %97:vgpr_32, 0, %75:vgpr_32, 0, %76:vgpr_32, 0, 0, implicit $mode, implicit $exec
 626   %99:vgpr_32 = V_MAX3_F32_e64 0, %98:vgpr_32, 0, %77:vgpr_32, 0, %78:vgpr_32, 0, 0, implicit $mode, implicit $exec
 627   %100:vgpr_32 = V_MAX3_F32_e64 0, %99:vgpr_32, 0, %79:vgpr_32, 0, %80:vgpr_32, 0, 0, implicit $mode, implicit $exec
 628   %101:vgpr_32 = V_MAX3_F32_e64 0, %100:vgpr_32, 0, %81:vgpr_32, 0, %82:vgpr_32, 0, 0, implicit $mode, implicit $exec
 629   %102:vgpr_32 = V_MAX3_F32_e64 0, %101:vgpr_32, 0, %83:vgpr_32, 0, %84:vgpr_32, 0, 0, implicit $mode, implicit $exec
 630   %103:vgpr_32 = V_MAX3_F32_e64 0, %102:vgpr_32, 0, %85:vgpr_32, 0, %86:vgpr_32, 0, 0, implicit $mode, implicit $exec
 631   %104:vgpr_32 = V_MAX3_F32_e64 0, %103:vgpr_32, 0, %87:vgpr_32, 0, %88:vgpr_32, 0, 0, implicit $mode, implicit $exec
 632   %105:vgpr_32 = DS_BPERMUTE_B32 %106:vgpr_32, %104:vgpr_32, 0, implicit $exec
 633   %107:vgpr_32 = contract nofpexcept V_MAX_F32_e32 %105:vgpr_32, %105:vgpr_32, implicit $mode, implicit $exec
 634   %108:vgpr_32 = contract nofpexcept V_MAX_F32_e32 %104:vgpr_32, %107:vgpr_32, implicit $mode, implicit $exec
 635   %109:vgpr_32 = DS_BPERMUTE_B32 %106:vgpr_32, %108:vgpr_32, 0, implicit $exec
 636   %110:vgpr_32 = V_CNDMASK_B32_e64 0, %109:vgpr_32, 0, %108:vgpr_32, %12:sreg_64_xexec, implicit $exec
 637   %111:vgpr_32 = contract nofpexcept V_MAX_F32_e32 %110:vgpr_32, %110:vgpr_32, implicit $mode, implicit $exec
 638   %112:vgpr_32 = contract nofpexcept V_MAX_F32_e32 %13:vgpr_32, %13:vgpr_32, implicit $mode, implicit $exec
 639   %113:vgpr_32 = contract nofpexcept V_MAX_F32_e32 %112:vgpr_32, %111:vgpr_32, implicit $mode, implicit $exec
 640   %114:vgpr_32 = contract nofpexcept V_FMA_F32_e64 0, %10.sub4:sgpr_512, 0, %43.sub0:vreg_512_align2, 1, %113:vgpr_32, 0, 0, implicit $mode, implicit $exec
 641   %115:vgpr_32 = afn nofpexcept V_MUL_F32_e32 1069066811, %114:vgpr_32, implicit $mode, implicit $exec
 642   %116:vgpr_32 = afn nofpexcept V_EXP_F32_e32 %115:vgpr_32, implicit $mode, implicit $exec
 643   %117:vgpr_32 = contract nofpexcept V_FMA_F32_e64 0, %10.sub4:sgpr_512, 0, %43.sub1:vreg_512_align2, 1, %113:vgpr_32, 0, 0, implicit $mode, implicit $exec
 644   %118:vgpr_32 = afn nofpexcept V_MUL_F32_e32 1069066811, %117:vgpr_32, implicit $mode, implicit $exec
 645   %119:vgpr_32 = afn nofpexcept V_EXP_F32_e32 %118:vgpr_32, implicit $mode, implicit $exec
 646   %120:vgpr_32 = contract nofpexcept V_FMA_F32_e64 0, %10.sub4:sgpr_512, 0, %43.sub2:vreg_512_align2, 1, %113:vgpr_32, 0, 0, implicit $mode, implicit $exec
 647   %121:vgpr_32 = afn nofpexcept V_MUL_F32_e32 1069066811, %120:vgpr_32, implicit $mode, implicit $exec
 648   %122:vgpr_32 = afn nofpexcept V_EXP_F32_e32 %121:vgpr_32, implicit $mode, implicit $exec
 649   %123:vgpr_32 = contract nofpexcept V_FMA_F32_e64 0, %10.sub4:sgpr_512, 0, %43.sub3:vreg_512_align2, 1, %113:vgpr_32, 0, 0, implicit $mode, implicit $exec
 650   %124:vgpr_32 = afn nofpexcept V_MUL_F32_e32 1069066811, %123:vgpr_32, implicit $mode, implicit $exec
 651   %125:vgpr_32 = afn nofpexcept V_EXP_F32_e32 %124:vgpr_32, implicit $mode, implicit $exec
 652   %126:vgpr_32 = contract nofpexcept V_FMA_F32_e64 0, %10.sub4:sgpr_512, 0, %43.sub4:vreg_512_align2, 1, %113:vgpr_32, 0, 0, implicit $mode, implicit $exec
 653   %127:vgpr_32 = afn nofpexcept V_MUL_F32_e32 1069066811, %126:vgpr_32, implicit $mode, implicit $exec
 654   %128:vgpr_32 = afn nofpexcept V_EXP_F32_e32 %127:vgpr_32, implicit $mode, implicit $exec
 655   %129:vgpr_32 = contract nofpexcept V_FMA_F32_e64 0, %10.sub4:sgpr_512, 0, %43.sub5:vreg_512_align2, 1, %113:vgpr_32, 0, 0, implicit $mode, implicit $exec
 656   %130:vgpr_32 = afn nofpexcept V_MUL_F32_e32 1069066811, %129:vgpr_32, implicit $mode, implicit $exec
 657   %131:vgpr_32 = afn nofpexcept V_EXP_F32_e32 %130:vgpr_32, implicit $mode, implicit $exec
 658   %132:vgpr_32 = contract nofpexcept V_FMA_F32_e64 0, %10.sub4:sgpr_512, 0, %43.sub6:vreg_512_align2, 1, %113:vgpr_32, 0, 0, implicit $mode, implicit $exec
 659   %133:vgpr_32 = afn nofpexcept V_MUL_F32_e32 1069066811, %132:vgpr_32, implicit $mode, implicit $exec
 660   %134:vgpr_32 = afn nofpexcept V_EXP_F32_e32 %133:vgpr_32, implicit $mode, implicit $exec
 661   %135:vgpr_32 = contract nofpexcept V_FMA_F32_e64 0, %10.sub4:sgpr_512, 0, %43.sub7:vreg_512_align2, 1, %113:vgpr_32, 0, 0, implicit $mode, implicit $exec
 662   %136:vgpr_32 = afn nofpexcept V_MUL_F32_e32 1069066811, %135:vgpr_32, implicit $mode, implicit $exec
 663   %137:vgpr_32 = afn nofpexcept V_EXP_F32_e32 %136:vgpr_32, implicit $mode, implicit $exec
 664   %138:vgpr_32 = contract nofpexcept V_FMA_F32_e64 0, %10.sub4:sgpr_512, 0, %43.sub8:vreg_512_align2, 1, %113:vgpr_32, 0, 0, implicit $mode, implicit $exec
 665   %139:vgpr_32 = afn nofpexcept V_MUL_F32_e32 1069066811, %138:vgpr_32, implicit $mode, implicit $exec
 666   %140:vgpr_32 = afn nofpexcept V_EXP_F32_e32 %139:vgpr_32, implicit $mode, implicit $exec
 667   %141:vgpr_32 = contract nofpexcept V_FMA_F32_e64 0, %10.sub4:sgpr_512, 0, %43.sub9:vreg_512_align2, 1, %113:vgpr_32, 0, 0, implicit $mode, implicit $exec
 668   %142:vgpr_32 = afn nofpexcept V_MUL_F32_e32 1069066811, %141:vgpr_32, implicit $mode, implicit $exec
 669   %143:vgpr_32 = afn nofpexcept V_EXP_F32_e32 %142:vgpr_32, implicit $mode, implicit $exec
 670   %144:vgpr_32 = contract nofpexcept V_FMA_F32_e64 0, %10.sub4:sgpr_512, 0, %43.sub10:vreg_512_align2, 1, %113:vgpr_32, 0, 0, implicit $mode, implicit $exec
 671   %145:vgpr_32 = afn nofpexcept V_MUL_F32_e32 1069066811, %144:vgpr_32, implicit $mode, implicit $exec
 672   %146:vgpr_32 = afn nofpexcept V_EXP_F32_e32 %145:vgpr_32, implicit $mode, implicit $exec
 673   %147:vgpr_32 = contract nofpexcept V_FMA_F32_e64 0, %10.sub4:sgpr_512, 0, %43.sub11:vreg_512_align2, 1, %113:vgpr_32, 0, 0, implicit $mode, implicit $exec
 674   %148:vgpr_32 = afn nofpexcept V_MUL_F32_e32 1069066811, %147:vgpr_32, implicit $mode, implicit $exec
 675   %149:vgpr_32 = afn nofpexcept V_EXP_F32_e32 %148:vgpr_32, implicit $mode, implicit $exec
 676   %150:vgpr_32 = contract nofpexcept V_FMA_F32_e64 0, %10.sub4:sgpr_512, 0, %43.sub12:vreg_512_align2, 1, %113:vgpr_32, 0, 0, implicit $mode, implicit $exec
 677   %151:vgpr_32 = afn nofpexcept V_MUL_F32_e32 1069066811, %150:vgpr_32, implicit $mode, implicit $exec
 678   %152:vgpr_32 = afn nofpexcept V_EXP_F32_e32 %151:vgpr_32, implicit $mode, implicit $exec
 679   %153:vgpr_32 = contract nofpexcept V_FMA_F32_e64 0, %10.sub4:sgpr_512, 0, %43.sub13:vreg_512_align2, 1, %113:vgpr_32, 0, 0, implicit $mode, implicit $exec
 680   %154:vgpr_32 = afn nofpexcept V_MUL_F32_e32 1069066811, %153:vgpr_32, implicit $mode, implicit $exec
 681   %155:vgpr_32 = afn nofpexcept V_EXP_F32_e32 %154:vgpr_32, implicit $mode, implicit $exec
 682   %156:vgpr_32 = contract nofpexcept V_FMA_F32_e64 0, %10.sub4:sgpr_512, 0, %43.sub14:vreg_512_align2, 1, %113:vgpr_32, 0, 0, implicit $mode, implicit $exec
 683   %157:vgpr_32 = afn nofpexcept V_MUL_F32_e32 1069066811, %156:vgpr_32, implicit $mode, implicit $exec
 684   %158:vgpr_32 = afn nofpexcept V_EXP_F32_e32 %157:vgpr_32, implicit $mode, implicit $exec
 685   %159:vgpr_32 = contract nofpexcept V_FMA_F32_e64 0, %10.sub4:sgpr_512, 0, %43.sub15:vreg_512_align2, 1, %113:vgpr_32, 0, 0, implicit $mode, implicit $exec
 686   %160:vgpr_32 = afn nofpexcept V_MUL_F32_e32 1069066811, %159:vgpr_32, implicit $mode, implicit $exec
 687   %161:vgpr_32 = afn nofpexcept V_EXP_F32_e32 %160:vgpr_32, implicit $mode, implicit $exec
 688   %162:vgpr_32 = contract nofpexcept V_FMA_F32_e64 0, %10.sub4:sgpr_512, 0, %46.sub0:vreg_512_align2, 1, %113:vgpr_32, 0, 0, implicit $mode, implicit $exec
 689   %163:vgpr_32 = afn nofpexcept V_MUL_F32_e32 1069066811, %162:vgpr_32, implicit $mode, implicit $exec
 690   %164:vgpr_32 = afn nofpexcept V_EXP_F32_e32 %163:vgpr_32, implicit $mode, implicit $exec
 691   %165:vgpr_32 = contract nofpexcept V_FMA_F32_e64 0, %10.sub4:sgpr_512, 0, %46.sub1:vreg_512_align2, 1, %113:vgpr_32, 0, 0, implicit $mode, implicit $exec
 692   %166:vgpr_32 = afn nofpexcept V_MUL_F32_e32 1069066811, %165:vgpr_32, implicit $mode, implicit $exec
 693   %167:vgpr_32 = afn nofpexcept V_EXP_F32_e32 %166:vgpr_32, implicit $mode, implicit $exec
 694   %168:vgpr_32 = contract nofpexcept V_FMA_F32_e64 0, %10.sub4:sgpr_512, 0, %46.sub2:vreg_512_align2, 1, %113:vgpr_32, 0, 0, implicit $mode, implicit $exec
 695   %169:vgpr_32 = afn nofpexcept V_MUL_F32_e32 1069066811, %168:vgpr_32, implicit $mode, implicit $exec
 696   %170:vgpr_32 = afn nofpexcept V_EXP_F32_e32 %169:vgpr_32, implicit $mode, implicit $exec
 697   %171:vgpr_32 = contract nofpexcept V_FMA_F32_e64 0, %10.sub4:sgpr_512, 0, %46.sub3:vreg_512_align2, 1, %113:vgpr_32, 0, 0, implicit $mode, implicit $exec
 698   %172:vgpr_32 = afn nofpexcept V_MUL_F32_e32 1069066811, %171:vgpr_32, implicit $mode, implicit $exec
 699   %173:vgpr_32 = afn nofpexcept V_EXP_F32_e32 %172:vgpr_32, implicit $mode, implicit $exec
 700   %174:vgpr_32 = contract nofpexcept V_FMA_F32_e64 0, %10.sub4:sgpr_512, 0, %46.sub4:vreg_512_align2, 1, %113:vgpr_32, 0, 0, implicit $mode, implicit $exec
 701   %175:vgpr_32 = afn nofpexcept V_MUL_F32_e32 1069066811, %174:vgpr_32, implicit $mode, implicit $exec
 702   %176:vgpr_32 = afn nofpexcept V_EXP_F32_e32 %175:vgpr_32, implicit $mode, implicit $exec
 703   %177:vgpr_32 = contract nofpexcept V_FMA_F32_e64 0, %10.sub4:sgpr_512, 0, %46.sub5:vreg_512_align2, 1, %113:vgpr_32, 0, 0, implicit $mode, implicit $exec
 704   %178:vgpr_32 = afn nofpexcept V_MUL_F32_e32 1069066811, %177:vgpr_32, implicit $mode, implicit $exec
 705   %179:vgpr_32 = afn nofpexcept V_EXP_F32_e32 %178:vgpr_32, implicit $mode, implicit $exec
 706   %180:vgpr_32 = contract nofpexcept V_FMA_F32_e64 0, %10.sub4:sgpr_512, 0, %46.sub6:vreg_512_align2, 1, %113:vgpr_32, 0, 0, implicit $mode, implicit $exec
 707   %181:vgpr_32 = afn nofpexcept V_MUL_F32_e32 1069066811, %180:vgpr_32, implicit $mode, implicit $exec
 708   %182:vgpr_32 = afn nofpexcept V_EXP_F32_e32 %181:vgpr_32, implicit $mode, implicit $exec
 709   %183:vgpr_32 = contract nofpexcept V_FMA_F32_e64 0, %10.sub4:sgpr_512, 0, %46.sub7:vreg_512_align2, 1, %113:vgpr_32, 0, 0, implicit $mode, implicit $exec
 710   %184:vgpr_32 = afn nofpexcept V_MUL_F32_e32 1069066811, %183:vgpr_32, implicit $mode, implicit $exec
 711   %185:vgpr_32 = afn nofpexcept V_EXP_F32_e32 %184:vgpr_32, implicit $mode, implicit $exec
 712   %186:vgpr_32 = contract nofpexcept V_FMA_F32_e64 0, %10.sub4:sgpr_512, 0, %46.sub8:vreg_512_align2, 1, %113:vgpr_32, 0, 0, implicit $mode, implicit $exec
 713   %187:vgpr_32 = afn nofpexcept V_MUL_F32_e32 1069066811, %186:vgpr_32, implicit $mode, implicit $exec
 714   %188:vgpr_32 = afn nofpexcept V_EXP_F32_e32 %187:vgpr_32, implicit $mode, implicit $exec
 715   %189:vgpr_32 = contract nofpexcept V_FMA_F32_e64 0, %10.sub4:sgpr_512, 0, %46.sub9:vreg_512_align2, 1, %113:vgpr_32, 0, 0, implicit $mode, implicit $exec
 716   %190:vgpr_32 = afn nofpexcept V_MUL_F32_e32 1069066811, %189:vgpr_32, implicit $mode, implicit $exec
 717   %191:vgpr_32 = afn nofpexcept V_EXP_F32_e32 %190:vgpr_32, implicit $mode, implicit $exec
 718   %192:vgpr_32 = contract nofpexcept V_FMA_F32_e64 0, %10.sub4:sgpr_512, 0, %46.sub10:vreg_512_align2, 1, %113:vgpr_32, 0, 0, implicit $mode, implicit $exec
 719   %193:vgpr_32 = afn nofpexcept V_MUL_F32_e32 1069066811, %192:vgpr_32, implicit $mode, implicit $exec
 720   %194:vgpr_32 = afn nofpexcept V_EXP_F32_e32 %193:vgpr_32, implicit $mode, implicit $exec
 721   %195:vgpr_32 = contract nofpexcept V_FMA_F32_e64 0, %10.sub4:sgpr_512, 0, %46.sub11:vreg_512_align2, 1, %113:vgpr_32, 0, 0, implicit $mode, implicit $exec
 722   %196:vgpr_32 = afn nofpexcept V_MUL_F32_e32 1069066811, %195:vgpr_32, implicit $mode, implicit $exec
 723   %197:vgpr_32 = afn nofpexcept V_EXP_F32_e32 %196:vgpr_32, implicit $mode, implicit $exec
 724   %198:vgpr_32 = contract nofpexcept V_FMA_F32_e64 0, %10.sub4:sgpr_512, 0, %46.sub12:vreg_512_align2, 1, %113:vgpr_32, 0, 0, implicit $mode, implicit $exec
 725   %199:vgpr_32 = afn nofpexcept V_MUL_F32_e32 1069066811, %198:vgpr_32, implicit $mode, implicit $exec
 726   %200:vgpr_32 = afn nofpexcept V_EXP_F32_e32 %199:vgpr_32, implicit $mode, implicit $exec
 727   %201:vgpr_32 = contract nofpexcept V_FMA_F32_e64 0, %10.sub4:sgpr_512, 0, %46.sub13:vreg_512_align2, 1, %113:vgpr_32, 0, 0, implicit $mode, implicit $exec
 728   %202:vgpr_32 = afn nofpexcept V_MUL_F32_e32 1069066811, %201:vgpr_32, implicit $mode, implicit $exec
 729   %203:vgpr_32 = afn nofpexcept V_EXP_F32_e32 %202:vgpr_32, implicit $mode, implicit $exec
 730   %204:vgpr_32 = contract nofpexcept V_FMA_F32_e64 0, %10.sub4:sgpr_512, 0, %46.sub14:vreg_512_align2, 1, %113:vgpr_32, 0, 0, implicit $mode, implicit $exec
 731   %205:vgpr_32 = afn nofpexcept V_MUL_F32_e32 1069066811, %204:vgpr_32, implicit $mode, implicit $exec
 732   %206:vgpr_32 = afn nofpexcept V_EXP_F32_e32 %205:vgpr_32, implicit $mode, implicit $exec
 733   %207:vgpr_32 = contract nofpexcept V_FMA_F32_e64 0, %10.sub4:sgpr_512, 0, %46.sub15:vreg_512_align2, 1, %113:vgpr_32, 0, 0, implicit $mode, implicit $exec
 734   %208:vgpr_32 = afn nofpexcept V_MUL_F32_e32 1069066811, %207:vgpr_32, implicit $mode, implicit $exec
 735   %209:vgpr_32 = afn nofpexcept V_EXP_F32_e32 %208:vgpr_32, implicit $mode, implicit $exec
 736   %210:vgpr_32 = contract nofpexcept V_ADD_F32_e32 0, %116:vgpr_32, implicit $mode, implicit $exec
 737   %211:vgpr_32 = contract nofpexcept V_ADD_F32_e32 %119:vgpr_32, %210:vgpr_32, implicit $mode, implicit $exec
 738   %212:vgpr_32 = contract nofpexcept V_ADD_F32_e32 %122:vgpr_32, %211:vgpr_32, implicit $mode, implicit $exec
 739   %213:vgpr_32 = contract nofpexcept V_ADD_F32_e32 %125:vgpr_32, %212:vgpr_32, implicit $mode, implicit $exec
 740   %214:vgpr_32 = contract nofpexcept V_ADD_F32_e32 %128:vgpr_32, %213:vgpr_32, implicit $mode, implicit $exec
 741   %215:vgpr_32 = contract nofpexcept V_ADD_F32_e32 %131:vgpr_32, %214:vgpr_32, implicit $mode, implicit $exec
 742   %216:vgpr_32 = contract nofpexcept V_ADD_F32_e32 %134:vgpr_32, %215:vgpr_32, implicit $mode, implicit $exec
 743   %217:vgpr_32 = contract nofpexcept V_ADD_F32_e32 %137:vgpr_32, %216:vgpr_32, implicit $mode, implicit $exec
 744   %218:vgpr_32 = contract nofpexcept V_ADD_F32_e32 %140:vgpr_32, %217:vgpr_32, implicit $mode, implicit $exec
 745   %219:vgpr_32 = contract nofpexcept V_ADD_F32_e32 %143:vgpr_32, %218:vgpr_32, implicit $mode, implicit $exec
 746   %220:vgpr_32 = contract nofpexcept V_ADD_F32_e32 %146:vgpr_32, %219:vgpr_32, implicit $mode, implicit $exec
 747   %221:vgpr_32 = contract nofpexcept V_ADD_F32_e32 %149:vgpr_32, %220:vgpr_32, implicit $mode, implicit $exec
 748   %222:vgpr_32 = contract nofpexcept V_ADD_F32_e32 %152:vgpr_32, %221:vgpr_32, implicit $mode, implicit $exec
 749   %223:vgpr_32 = contract nofpexcept V_ADD_F32_e32 %155:vgpr_32, %222:vgpr_32, implicit $mode, implicit $exec
 750   %224:vgpr_32 = contract nofpexcept V_ADD_F32_e32 %158:vgpr_32, %223:vgpr_32, implicit $mode, implicit $exec
 751   %225:vgpr_32 = contract nofpexcept V_ADD_F32_e32 %161:vgpr_32, %224:vgpr_32, implicit $mode, implicit $exec
 752   %226:vgpr_32 = contract nofpexcept V_ADD_F32_e32 %164:vgpr_32, %225:vgpr_32, implicit $mode, implicit $exec
 753   %227:vgpr_32 = contract nofpexcept V_ADD_F32_e32 %167:vgpr_32, %226:vgpr_32, implicit $mode, implicit $exec
 754   %228:vgpr_32 = contract nofpexcept V_ADD_F32_e32 %170:vgpr_32, %227:vgpr_32, implicit $mode, implicit $exec
 755   %229:vgpr_32 = contract nofpexcept V_ADD_F32_e32 %173:vgpr_32, %228:vgpr_32, implicit $mode, implicit $exec
 756   %230:vgpr_32 = contract nofpexcept V_ADD_F32_e32 %176:vgpr_32, %229:vgpr_32, implicit $mode, implicit $exec
 757   %231:vgpr_32 = contract nofpexcept V_ADD_F32_e32 %179:vgpr_32, %230:vgpr_32, implicit $mode, implicit $exec
 758   %232:vgpr_32 = contract nofpexcept V_ADD_F32_e32 %182:vgpr_32, %231:vgpr_32, implicit $mode, implicit $exec
 759   %233:vgpr_32 = contract nofpexcept V_ADD_F32_e32 %185:vgpr_32, %232:vgpr_32, implicit $mode, implicit $exec
 760   %234:vgpr_32 = contract nofpexcept V_ADD_F32_e32 %188:vgpr_32, %233:vgpr_32, implicit $mode, implicit $exec
 761   %235:vgpr_32 = contract nofpexcept V_ADD_F32_e32 %191:vgpr_32, %234:vgpr_32, implicit $mode, implicit $exec
 762   %236:vgpr_32 = contract nofpexcept V_ADD_F32_e32 %194:vgpr_32, %235:vgpr_32, implicit $mode, implicit $exec
 763   %237:vgpr_32 = contract nofpexcept V_ADD_F32_e32 %197:vgpr_32, %236:vgpr_32, implicit $mode, implicit $exec
 764   %238:vgpr_32 = contract nofpexcept V_ADD_F32_e32 %200:vgpr_32, %237:vgpr_32, implicit $mode, implicit $exec
 765   %239:vgpr_32 = contract nofpexcept V_ADD_F32_e32 %203:vgpr_32, %238:vgpr_32, implicit $mode, implicit $exec
 766   %240:vgpr_32 = contract nofpexcept V_ADD_F32_e32 %206:vgpr_32, %239:vgpr_32, implicit $mode, implicit $exec
 767   %241:vgpr_32 = contract nofpexcept V_ADD_F32_e32 %209:vgpr_32, %240:vgpr_32, implicit $mode, implicit $exec
 768   %242:vgpr_32 = DS_BPERMUTE_B32 %106:vgpr_32, %241:vgpr_32, 0, implicit $exec
 769   %243:vgpr_32 = contract nofpexcept V_ADD_F32_e32 %241:vgpr_32, %242:vgpr_32, implicit $mode, implicit $exec
 770   %244:vgpr_32 = DS_BPERMUTE_B32 %106:vgpr_32, %243:vgpr_32, 0, implicit $exec
 771   %0:vgpr_32 = V_CNDMASK_B32_e64 0, %244:vgpr_32, 0, %243:vgpr_32, %12:sreg_64_xexec, implicit $exec
 772   %245:vgpr_32 = contract nofpexcept V_SUB_F32_e32 %13:vgpr_32, %113:vgpr_32, implicit $mode, implicit $exec
 773   %246:vgpr_32 = afn nofpexcept V_MUL_F32_e32 1069066811, %245:vgpr_32, implicit $mode, implicit $exec
 774   undef %247.sub0:vreg_64_align2 = afn nofpexcept V_EXP_F32_e32 %246:vgpr_32, implicit $mode, implicit $exec
 775   INLINEASM &"s_waitcnt vmcnt($0)", 57 /* sideeffect mayload maystore isconvergent attdialect */, 13 /* imm */, 8, !0
 776   %248:vgpr_32 = V_PERM_B32_e64 %40.sub0:vreg_64_align2, %38.sub0:vreg_64_align2, %14:sreg_32, implicit $exec
 777   %249:vgpr_32 = V_PERM_B32_e64 %40.sub0:vreg_64_align2, %38.sub0:vreg_64_align2, %15:sreg_32, implicit $exec
 778   %250:vgpr_32 = V_PERM_B32_e64 %40.sub1:vreg_64_align2, %38.sub1:vreg_64_align2, %14:sreg_32, implicit $exec
 779   %251:vgpr_32 = V_PERM_B32_e64 %40.sub1:vreg_64_align2, %38.sub1:vreg_64_align2, %15:sreg_32, implicit $exec
 780   %252:vgpr_32 = V_ADD_U32_e32 %27:sreg_32, %16:vgpr_32, implicit $exec
 781   %253:vgpr_32 = V_AND_B32_e32 536870911, %252:vgpr_32, implicit $exec
 782   %254:vgpr_32 = nsw V_MUL_LO_U32_e64 %253:vgpr_32, %17:sreg_32, implicit $exec
 783   %255:vgpr_32 = V_ADD_LSHL_U32_e64 %18:vgpr_32, %254:vgpr_32, 1, implicit $exec
 784   DS_WRITE_B32_gfx9 %255:vgpr_32, %248:vgpr_32, 0, 0, implicit $exec
 785   %256:vgpr_32 = V_LSHL_ADD_U32_e64 %19:vgpr_32, 1, %255:vgpr_32, implicit $exec
 786   DS_WRITE_B32_gfx9 %256:vgpr_32, %249:vgpr_32, 0, 0, implicit $exec
 787   %257:vgpr_32 = V_LSHL_ADD_U32_e64 %20:vgpr_32, 1, %256:vgpr_32, implicit $exec
 788   DS_WRITE_B32_gfx9 %257:vgpr_32, %250:vgpr_32, 0, 0, implicit $exec
 789   %258:vgpr_32 = V_LSHL_ADD_U32_e64 %21:vgpr_32, 1, %257:vgpr_32, implicit $exec
 790   DS_WRITE_B32_gfx9 %258:vgpr_32, %251:vgpr_32, 0, 0, implicit $exec
 791   %0:vgpr_32 = contract nofpexcept V_FMAC_F32_e32 %1:vgpr_32, %247.sub0:vreg_64_align2, %0:vgpr_32, implicit $mode, implicit $exec
 792   %259.sub0_sub1:vreg_512_align2 = contract nofpexcept V_PK_MUL_F32 8, %259.sub0_sub1:vreg_512_align2, 0, %247:vreg_64_align2, 0, 0, 0, 0, 0, implicit $mode, implicit $exec
 793   %259.sub2_sub3:vreg_512_align2 = contract nofpexcept V_PK_MUL_F32 8, %259.sub2_sub3:vreg_512_align2, 0, %247:vreg_64_align2, 0, 0, 0, 0, 0, implicit $mode, implicit $exec
 794   %259.sub4_sub5:vreg_512_align2 = contract nofpexcept V_PK_MUL_F32 8, %259.sub4_sub5:vreg_512_align2, 0, %247:vreg_64_align2, 0, 0, 0, 0, 0, implicit $mode, implicit $exec
 795   %259.sub6_sub7:vreg_512_align2 = contract nofpexcept V_PK_MUL_F32 8, %259.sub6_sub7:vreg_512_align2, 0, %247:vreg_64_align2, 0, 0, 0, 0, 0, implicit $mode, implicit $exec
 796   %259.sub8_sub9:vreg_512_align2 = contract nofpexcept V_PK_MUL_F32 8, %259.sub8_sub9:vreg_512_align2, 0, %247:vreg_64_align2, 0, 0, 0, 0, 0, implicit $mode, implicit $exec
 797   %259.sub10_sub11:vreg_512_align2 = contract nofpexcept V_PK_MUL_F32 8, %259.sub10_sub11:vreg_512_align2, 0, %247:vreg_64_align2, 0, 0, 0, 0, 0, implicit $mode, implicit $exec
 798   %259.sub12_sub13:vreg_512_align2 = contract nofpexcept V_PK_MUL_F32 8, %259.sub12_sub13:vreg_512_align2, 0, %247:vreg_64_align2, 0, 0, 0, 0, 0, implicit $mode, implicit $exec
 799   %259.sub14_sub15:vreg_512_align2 = contract nofpexcept V_PK_MUL_F32 8, %259.sub14_sub15:vreg_512_align2, 0, %247:vreg_64_align2, 0, 0, 0, 0, 0, implicit $mode, implicit $exec
 800   %260.sub0_sub1:vreg_512_align2 = contract nofpexcept V_PK_MUL_F32 8, %260.sub0_sub1:vreg_512_align2, 0, %247:vreg_64_align2, 0, 0, 0, 0, 0, implicit $mode, implicit $exec
 801   %260.sub2_sub3:vreg_512_align2 = contract nofpexcept V_PK_MUL_F32 8, %260.sub2_sub3:vreg_512_align2, 0, %247:vreg_64_align2, 0, 0, 0, 0, 0, implicit $mode, implicit $exec
 802   %260.sub4_sub5:vreg_512_align2 = contract nofpexcept V_PK_MUL_F32 8, %260.sub4_sub5:vreg_512_align2, 0, %247:vreg_64_align2, 0, 0, 0, 0, 0, implicit $mode, implicit $exec
 803   %260.sub6_sub7:vreg_512_align2 = contract nofpexcept V_PK_MUL_F32 8, %260.sub6_sub7:vreg_512_align2, 0, %247:vreg_64_align2, 0, 0, 0, 0, 0, implicit $mode, implicit $exec
 804   %260.sub8_sub9:vreg_512_align2 = contract nofpexcept V_PK_MUL_F32 8, %260.sub8_sub9:vreg_512_align2, 0, %247:vreg_64_align2, 0, 0, 0, 0, 0, implicit $mode, implicit $exec
 805   %260.sub10_sub11:vreg_512_align2 = contract nofpexcept V_PK_MUL_F32 8, %260.sub10_sub11:vreg_512_align2, 0, %247:vreg_64_align2, 0, 0, 0, 0, 0, implicit $mode, implicit $exec
 806   %260.sub12_sub13:vreg_512_align2 = contract nofpexcept V_PK_MUL_F32 8, %260.sub12_sub13:vreg_512_align2, 0, %247:vreg_64_align2, 0, 0, 0, 0, 0, implicit $mode, implicit $exec
 807   %260.sub14_sub15:vreg_512_align2 = contract nofpexcept V_PK_MUL_F32 8, %260.sub14_sub15:vreg_512_align2, 0, %247:vreg_64_align2, 0, 0, 0, 0, 0, implicit $mode, implicit $exec
 808   %261:vgpr_32 = nofpexcept V_CVT_F16_F32_e32 %116:vgpr_32, implicit $mode, implicit $exec
 809   %262:vgpr_32 = nofpexcept V_CVT_F16_F32_e32 %119:vgpr_32, implicit $mode, implicit $exec
 810   %263:vgpr_32 = nofpexcept V_CVT_F16_F32_e32 %122:vgpr_32, implicit $mode, implicit $exec
 811   %264:vgpr_32 = nofpexcept V_CVT_F16_F32_e32 %128:vgpr_32, implicit $mode, implicit $exec
 812   %265:vgpr_32 = nofpexcept V_CVT_F16_F32_e32 %131:vgpr_32, implicit $mode, implicit $exec
 813   %266:vgpr_32 = nofpexcept V_CVT_F16_F32_e32 %134:vgpr_32, implicit $mode, implicit $exec
 814   %267:vgpr_32 = nofpexcept V_CVT_F16_F32_e32 %140:vgpr_32, implicit $mode, implicit $exec
 815   %268:vgpr_32 = nofpexcept V_CVT_F16_F32_e32 %143:vgpr_32, implicit $mode, implicit $exec
 816   %269:vgpr_32 = nofpexcept V_CVT_F16_F32_e32 %146:vgpr_32, implicit $mode, implicit $exec
 817   %270:vgpr_32 = nofpexcept V_CVT_F16_F32_e32 %152:vgpr_32, implicit $mode, implicit $exec
 818   %271:vgpr_32 = nofpexcept V_CVT_F16_F32_e32 %155:vgpr_32, implicit $mode, implicit $exec
 819   %272:vgpr_32 = nofpexcept V_CVT_F16_F32_e32 %158:vgpr_32, implicit $mode, implicit $exec
 820   %273:vgpr_32 = V_ADD_U32_e32 %22:vgpr_32, %37:vgpr_32, implicit $exec
 821   %274:vreg_64_align2 = BUFFER_LOAD_DWORDX2_OFFEN %273:vgpr_32, %8:sgpr_128, 0, 0, 0, 0, implicit $exec
 822   %275:vgpr_32 = V_ADD_U32_e32 %23:vgpr_32, %37:vgpr_32, implicit $exec
 823   %276:vreg_64_align2 = BUFFER_LOAD_DWORDX2_OFFEN %275:vgpr_32, %8:sgpr_128, 0, 0, 0, 0, implicit $exec
 824   INLINEASM &"s_waitcnt vmcnt($0)", 57 /* sideeffect mayload maystore isconvergent attdialect */, 13 /* imm */, 8, !0
 825   %277:vreg_128_align2 = DS_READ_B128_gfx9 %24:vgpr_32, 0, 0, implicit $exec
 826   %278:vreg_128_align2 = DS_READ_B128_gfx9 %24:vgpr_32, 576, 0, implicit $exec
 827   %279:vreg_128_align2 = DS_READ_B128_gfx9 %25:vgpr_32, 0, 0, implicit $exec
 828   %280:vreg_128_align2 = DS_READ_B128_gfx9 %25:vgpr_32, 576, 0, implicit $exec
 829   INLINEASM &"s_waitcnt vmcnt($0)", 57 /* sideeffect mayload maystore isconvergent attdialect */, 13 /* imm */, 8, !0
 830   %281:vgpr_32 = V_PERM_B32_e64 %276.sub0:vreg_64_align2, %274.sub0:vreg_64_align2, %14:sreg_32, implicit $exec
 831   %282:vgpr_32 = V_PERM_B32_e64 %276.sub0:vreg_64_align2, %274.sub0:vreg_64_align2, %15:sreg_32, implicit $exec
 832   %283:vgpr_32 = V_PERM_B32_e64 %276.sub1:vreg_64_align2, %274.sub1:vreg_64_align2, %14:sreg_32, implicit $exec
 833   %284:vgpr_32 = V_PERM_B32_e64 %276.sub1:vreg_64_align2, %274.sub1:vreg_64_align2, %15:sreg_32, implicit $exec
 834   DS_WRITE_B32_gfx9 %255:vgpr_32, %281:vgpr_32, 0, 0, implicit $exec
 835   DS_WRITE_B32_gfx9 %256:vgpr_32, %282:vgpr_32, 0, 0, implicit $exec
 836   DS_WRITE_B32_gfx9 %257:vgpr_32, %283:vgpr_32, 0, 0, implicit $exec
 837   DS_WRITE_B32_gfx9 %258:vgpr_32, %284:vgpr_32, 0, 0, implicit $exec
 838   %285:vgpr_32 = nofpexcept V_CVT_F16_F32_e32 %125:vgpr_32, implicit $mode, implicit $exec
 839   %286:vgpr_32 = nofpexcept V_CVT_F16_F32_e32 %137:vgpr_32, implicit $mode, implicit $exec
 840   %287:vgpr_32 = nofpexcept V_CVT_F16_F32_e32 %149:vgpr_32, implicit $mode, implicit $exec
 841   %288:vgpr_32 = nofpexcept V_CVT_F16_F32_e32 %161:vgpr_32, implicit $mode, implicit $exec
 842   undef %289.sub1:vreg_64_align2 = nofpexcept V_PACK_B32_F16_e64 0, %263:vgpr_32, 0, %285:vgpr_32, 0, 0, implicit $mode, implicit $exec
 843   %289.sub0:vreg_64_align2 = nofpexcept V_PACK_B32_F16_e64 0, %261:vgpr_32, 0, %262:vgpr_32, 0, 0, implicit $mode, implicit $exec
 844   undef %290.sub1:vreg_64_align2 = nofpexcept V_PACK_B32_F16_e64 0, %266:vgpr_32, 0, %286:vgpr_32, 0, 0, implicit $mode, implicit $exec
 845   %290.sub0:vreg_64_align2 = nofpexcept V_PACK_B32_F16_e64 0, %264:vgpr_32, 0, %265:vgpr_32, 0, 0, implicit $mode, implicit $exec
 846   undef %291.sub1:vreg_64_align2 = nofpexcept V_PACK_B32_F16_e64 0, %269:vgpr_32, 0, %287:vgpr_32, 0, 0, implicit $mode, implicit $exec
 847   %291.sub0:vreg_64_align2 = nofpexcept V_PACK_B32_F16_e64 0, %267:vgpr_32, 0, %268:vgpr_32, 0, 0, implicit $mode, implicit $exec
 848   undef %292.sub1:vreg_64_align2 = nofpexcept V_PACK_B32_F16_e64 0, %272:vgpr_32, 0, %288:vgpr_32, 0, 0, implicit $mode, implicit $exec
 849   %292.sub0:vreg_64_align2 = nofpexcept V_PACK_B32_F16_e64 0, %270:vgpr_32, 0, %271:vgpr_32, 0, 0, implicit $mode, implicit $exec
 850   %259:vreg_512_align2 = contract V_MFMA_F32_32X32X8F16_mac_vgprcd_e64 %277.sub0_sub1:vreg_128_align2, %289:vreg_64_align2, %259:vreg_512_align2, 0, 0, 0, implicit $mode, implicit $exec
 851   %259:vreg_512_align2 = contract V_MFMA_F32_32X32X8F16_mac_vgprcd_e64 %277.sub2_sub3:vreg_128_align2, %290:vreg_64_align2, %259:vreg_512_align2, 0, 0, 0, implicit $mode, implicit $exec
 852   %260:vreg_512_align2 = contract V_MFMA_F32_32X32X8F16_mac_vgprcd_e64 %278.sub0_sub1:vreg_128_align2, %289:vreg_64_align2, %260:vreg_512_align2, 0, 0, 0, implicit $mode, implicit $exec
 853   %260:vreg_512_align2 = contract V_MFMA_F32_32X32X8F16_mac_vgprcd_e64 %278.sub2_sub3:vreg_128_align2, %290:vreg_64_align2, %260:vreg_512_align2, 0, 0, 0, implicit $mode, implicit $exec
 854   %259:vreg_512_align2 = contract V_MFMA_F32_32X32X8F16_mac_vgprcd_e64 %279.sub0_sub1:vreg_128_align2, %291:vreg_64_align2, %259:vreg_512_align2, 0, 0, 0, implicit $mode, implicit $exec
 855   %259:vreg_512_align2 = contract V_MFMA_F32_32X32X8F16_mac_vgprcd_e64 %279.sub2_sub3:vreg_128_align2, %292:vreg_64_align2, %259:vreg_512_align2, 0, 0, 0, implicit $mode, implicit $exec
 856   %260:vreg_512_align2 = contract V_MFMA_F32_32X32X8F16_mac_vgprcd_e64 %280.sub0_sub1:vreg_128_align2, %291:vreg_64_align2, %260:vreg_512_align2, 0, 0, 0, implicit $mode, implicit $exec
 857   %260:vreg_512_align2 = contract V_MFMA_F32_32X32X8F16_mac_vgprcd_e64 %280.sub2_sub3:vreg_128_align2, %292:vreg_64_align2, %260:vreg_512_align2, 0, 0, 0, implicit $mode, implicit $exec
 858   %293:vgpr_32 = nofpexcept V_CVT_F16_F32_e32 %209:vgpr_32, implicit $mode, implicit $exec
 859   %294:vgpr_32 = nofpexcept V_CVT_F16_F32_e32 %203:vgpr_32, implicit $mode, implicit $exec
 860   %295:vgpr_32 = nofpexcept V_CVT_F16_F32_e32 %197:vgpr_32, implicit $mode, implicit $exec
 861   %296:vgpr_32 = nofpexcept V_CVT_F16_F32_e32 %191:vgpr_32, implicit $mode, implicit $exec
 862   %297:vgpr_32 = nofpexcept V_CVT_F16_F32_e32 %185:vgpr_32, implicit $mode, implicit $exec
 863   %298:vgpr_32 = nofpexcept V_CVT_F16_F32_e32 %179:vgpr_32, implicit $mode, implicit $exec
 864   %299:vgpr_32 = nofpexcept V_CVT_F16_F32_e32 %173:vgpr_32, implicit $mode, implicit $exec
 865   %300:vgpr_32 = nofpexcept V_CVT_F16_F32_e32 %167:vgpr_32, implicit $mode, implicit $exec
 866   %301:vgpr_32 = nofpexcept V_CVT_F16_F32_e32 %206:vgpr_32, implicit $mode, implicit $exec
 867   %302:vgpr_32 = nofpexcept V_CVT_F16_F32_e32 %200:vgpr_32, implicit $mode, implicit $exec
 868   %303:vgpr_32 = nofpexcept V_CVT_F16_F32_e32 %194:vgpr_32, implicit $mode, implicit $exec
 869   %304:vgpr_32 = nofpexcept V_CVT_F16_F32_e32 %188:vgpr_32, implicit $mode, implicit $exec
 870   %305:vgpr_32 = nofpexcept V_CVT_F16_F32_e32 %182:vgpr_32, implicit $mode, implicit $exec
 871   %306:vgpr_32 = nofpexcept V_CVT_F16_F32_e32 %176:vgpr_32, implicit $mode, implicit $exec
 872   %307:vgpr_32 = nofpexcept V_CVT_F16_F32_e32 %170:vgpr_32, implicit $mode, implicit $exec
 873   %308:vgpr_32 = nofpexcept V_CVT_F16_F32_e32 %164:vgpr_32, implicit $mode, implicit $exec
 874   INLINEASM &"s_waitcnt vmcnt($0)", 57 /* sideeffect mayload maystore isconvergent attdialect */, 13 /* imm */, 8, !0
 875   undef %309.sub1:vreg_64_align2 = nofpexcept V_PACK_B32_F16_e64 0, %307:vgpr_32, 0, %299:vgpr_32, 0, 0, implicit $mode, implicit $exec
 876   %309.sub0:vreg_64_align2 = nofpexcept V_PACK_B32_F16_e64 0, %308:vgpr_32, 0, %300:vgpr_32, 0, 0, implicit $mode, implicit $exec
 877   undef %310.sub1:vreg_64_align2 = nofpexcept V_PACK_B32_F16_e64 0, %305:vgpr_32, 0, %297:vgpr_32, 0, 0, implicit $mode, implicit $exec
 878   %310.sub0:vreg_64_align2 = nofpexcept V_PACK_B32_F16_e64 0, %306:vgpr_32, 0, %298:vgpr_32, 0, 0, implicit $mode, implicit $exec
 879   undef %311.sub1:vreg_64_align2 = nofpexcept V_PACK_B32_F16_e64 0, %303:vgpr_32, 0, %295:vgpr_32, 0, 0, implicit $mode, implicit $exec
 880   %311.sub0:vreg_64_align2 = nofpexcept V_PACK_B32_F16_e64 0, %304:vgpr_32, 0, %296:vgpr_32, 0, 0, implicit $mode, implicit $exec
 881   undef %312.sub1:vreg_64_align2 = nofpexcept V_PACK_B32_F16_e64 0, %301:vgpr_32, 0, %293:vgpr_32, 0, 0, implicit $mode, implicit $exec
 882   %312.sub0:vreg_64_align2 = nofpexcept V_PACK_B32_F16_e64 0, %302:vgpr_32, 0, %294:vgpr_32, 0, 0, implicit $mode, implicit $exec
 883   %313:vreg_128_align2 = DS_READ_B128_gfx9 %24:vgpr_32, 0, 0, implicit $exec
 884   %259:vreg_512_align2 = contract V_MFMA_F32_32X32X8F16_mac_vgprcd_e64 %313.sub0_sub1:vreg_128_align2, %309:vreg_64_align2, %259:vreg_512_align2, 0, 0, 0, implicit $mode, implicit $exec
 885   %259:vreg_512_align2 = contract V_MFMA_F32_32X32X8F16_mac_vgprcd_e64 %313.sub2_sub3:vreg_128_align2, %310:vreg_64_align2, %259:vreg_512_align2, 0, 0, 0, implicit $mode, implicit $exec
 886   %314:vreg_128_align2 = DS_READ_B128_gfx9 %24:vgpr_32, 576, 0, implicit $exec
 887   %260:vreg_512_align2 = contract V_MFMA_F32_32X32X8F16_mac_vgprcd_e64 %314.sub0_sub1:vreg_128_align2, %309:vreg_64_align2, %260:vreg_512_align2, 0, 0, 0, implicit $mode, implicit $exec
 888   %260:vreg_512_align2 = contract V_MFMA_F32_32X32X8F16_mac_vgprcd_e64 %314.sub2_sub3:vreg_128_align2, %310:vreg_64_align2, %260:vreg_512_align2, 0, 0, 0, implicit $mode, implicit $exec
 889   %315:vreg_128_align2 = DS_READ_B128_gfx9 %25:vgpr_32, 0, 0, implicit $exec
 890   %259:vreg_512_align2 = contract V_MFMA_F32_32X32X8F16_mac_vgprcd_e64 %315.sub0_sub1:vreg_128_align2, %311:vreg_64_align2, %259:vreg_512_align2, 0, 0, 0, implicit $mode, implicit $exec
 891   %259:vreg_512_align2 = contract V_MFMA_F32_32X32X8F16_mac_vgprcd_e64 %315.sub2_sub3:vreg_128_align2, %312:vreg_64_align2, %259:vreg_512_align2, 0, 0, 0, implicit $mode, implicit $exec
 892   %316:vreg_128_align2 = DS_READ_B128_gfx9 %25:vgpr_32, 576, 0, implicit $exec
 893   %260:vreg_512_align2 = contract V_MFMA_F32_32X32X8F16_mac_vgprcd_e64 %316.sub0_sub1:vreg_128_align2, %311:vreg_64_align2, %260:vreg_512_align2, 0, 0, 0, implicit $mode, implicit $exec
 894   %260:vreg_512_align2 = contract V_MFMA_F32_32X32X8F16_mac_vgprcd_e64 %316.sub2_sub3:vreg_128_align2, %312:vreg_64_align2, %260:vreg_512_align2, 0, 0, 0, implicit $mode, implicit $exec
 895   INLINEASM &"s_waitcnt vmcnt($0)", 57 /* sideeffect mayload maystore isconvergent attdialect */, 13 /* imm */, 8, !0
 896   %37:vgpr_32 = V_ADD_U32_e32 %26:sreg_32, %37:vgpr_32, implicit $exec
 897   %29:vgpr_32 = nuw V_ADD_U32_e32 64, %29:vgpr_32, implicit $exec
 898   S_ENDPGM 0
 899 ...
 900