llvm/test/MC/Disassembler/AMDGPU/gfx940_mai.txt

   1 # RUN: llvm-mc -triple=amdgcn -mcpu=gfx940 -show-encoding -disassemble %s | FileCheck -check-prefix=GFX940 %s
   2
   3 # GFX940: v_accvgpr_write_b32 a10, s20 ; encoding: [0x0a,0x40,0xd9,0xd3,0x14,0x00,0x00,0x18]
   4 0x0a,0x40,0xd9,0xd3,0x14,0x00,0x00,0x18
   5
   6 # GFX940: v_mfma_i32_32x32x16_i8 v[0:15], v[2:3], v[4:5], v[0:15] ; encoding: [0x00,0x00,0xd6,0xd3,0x02,0x09,0x02,0x04]
   7 0x00,0x00,0xd6,0xd3,0x02,0x09,0x02,0x04
   8
   9 # GFX940: v_mfma_i32_32x32x16_i8 a[0:15], v[2:3], v[4:5], a[0:15] ; encoding: [0x00,0x80,0xd6,0xd3,0x02,0x09,0x02,0x04]
  10 0x00,0x80,0xd6,0xd3,0x02,0x09,0x02,0x04
  11
  12 # GFX940: v_mfma_i32_32x32x16_i8 a[0:15], v[2:3], v[4:5], a[0:15] blgp:5 ; encoding: [0x00,0x80,0xd6,0xd3,0x02,0x09,0x02,0xa4]
  13 0x00,0x80,0xd6,0xd3,0x02,0x09,0x02,0xa4
  14
  15 # GFX940: v_mfma_i32_16x16x32_i8 v[0:3], v[2:3], v[4:5], v[0:3] ; encoding: [0x00,0x00,0xd7,0xd3,0x02,0x09,0x02,0x04]
  16 0x00,0x00,0xd7,0xd3,0x02,0x09,0x02,0x04
  17
  18 # GFX940: v_mfma_i32_16x16x32_i8 a[0:3], v[2:3], v[4:5], a[0:3] ; encoding: [0x00,0x80,0xd7,0xd3,0x02,0x09,0x02,0x04]
  19 0x00,0x80,0xd7,0xd3,0x02,0x09,0x02,0x04
  20
  21 # GFX940: v_mfma_i32_16x16x32_i8 a[0:3], v[2:3], v[4:5], a[0:3] blgp:5 ; encoding: [0x00,0x80,0xd7,0xd3,0x02,0x09,0x02,0xa4]
  22 0x00,0x80,0xd7,0xd3,0x02,0x09,0x02,0xa4
  23
  24 # GFX940: v_mfma_f32_32x32x4_2b_bf16 v[0:31], v[2:3], v[4:5], v[2:33] ; encoding: [0x00,0x00,0xdd,0xd3,0x02,0x09,0x0a,0x04]
  25 0x00,0x00,0xdd,0xd3,0x02,0x09,0x0a,0x04
  26
  27 # GFX940: v_mfma_f32_32x32x4_2b_bf16 a[0:31], v[2:3], v[4:5], a[2:33] ; encoding: [0x00,0x80,0xdd,0xd3,0x02,0x09,0x0a,0x04]
  28 0x00,0x80,0xdd,0xd3,0x02,0x09,0x0a,0x04
  29
  30 # GFX940: v_mfma_f32_16x16x4_4b_bf16 v[0:15], v[2:3], v[4:5], v[2:17] ; encoding: [0x00,0x00,0xde,0xd3,0x02,0x09,0x0a,0x04]
  31 0x00,0x00,0xde,0xd3,0x02,0x09,0x0a,0x04
  32
  33 # GFX940: v_mfma_f32_16x16x4_4b_bf16 a[0:15], v[2:3], v[4:5], a[2:17] ; encoding: [0x00,0x80,0xde,0xd3,0x02,0x09,0x0a,0x04]
  34 0x00,0x80,0xde,0xd3,0x02,0x09,0x0a,0x04
  35
  36 # GFX940: v_mfma_f32_4x4x4_16b_bf16 v[0:3], v[2:3], v[4:5], v[2:5] ; encoding: [0x00,0x00,0xdf,0xd3,0x02,0x09,0x0a,0x04]
  37 0x00,0x00,0xdf,0xd3,0x02,0x09,0x0a,0x04
  38
  39 # GFX940: v_mfma_f32_4x4x4_16b_bf16 a[0:3], v[2:3], v[4:5], a[2:5] ; encoding: [0x00,0x80,0xdf,0xd3,0x02,0x09,0x0a,0x04]
  40 0x00,0x80,0xdf,0xd3,0x02,0x09,0x0a,0x04
  41
  42 # GFX940: v_mfma_f32_32x32x8_bf16 v[0:15], v[2:3], v[4:5], v[2:17] ; encoding: [0x00,0x00,0xe0,0xd3,0x02,0x09,0x0a,0x04]
  43 0x00,0x00,0xe0,0xd3,0x02,0x09,0x0a,0x04
  44
  45 # GFX940: v_mfma_f32_32x32x8_bf16 a[0:15], v[2:3], v[4:5], a[2:17] ; encoding: [0x00,0x80,0xe0,0xd3,0x02,0x09,0x0a,0x04]
  46 0x00,0x80,0xe0,0xd3,0x02,0x09,0x0a,0x04
  47
  48 # GFX940: v_mfma_f32_16x16x16_bf16 v[0:3], v[2:3], v[4:5], v[2:5] ; encoding: [0x00,0x00,0xe1,0xd3,0x02,0x09,0x0a,0x04]
  49 0x00,0x00,0xe1,0xd3,0x02,0x09,0x0a,0x04
  50
  51 # GFX940: v_mfma_f32_16x16x16_bf16 a[0:3], v[2:3], v[4:5], a[2:5] ; encoding: [0x00,0x80,0xe1,0xd3,0x02,0x09,0x0a,0x04]
  52 0x00,0x80,0xe1,0xd3,0x02,0x09,0x0a,0x04
  53
  54 # GFX940: v_mfma_f32_16x16x8_xf32 a[0:3], v[2:3], v[4:5], a[2:5] ; encoding: [0x00,0x80,0xbe,0xd3,0x02,0x09,0x0a,0x04]
  55 0x00,0x80,0xbe,0xd3,0x02,0x09,0x0a,0x04
  56
  57 # GFX940: v_mfma_f32_16x16x8_xf32 v[0:3], v[2:3], v[4:5], v[2:5] ; encoding: [0x00,0x00,0xbe,0xd3,0x02,0x09,0x0a,0x04]
  58 0x00,0x00,0xbe,0xd3,0x02,0x09,0x0a,0x04
  59
  60 # GFX940: v_mfma_f32_32x32x4_xf32 v[0:15], v[2:3], v[4:5], v[2:17] ; encoding: [0x00,0x00,0xbf,0xd3,0x02,0x09,0x0a,0x04]
  61 0x00,0x00,0xbf,0xd3,0x02,0x09,0x0a,0x04
  62
  63 # GFX940: v_mfma_f32_32x32x4_xf32 a[0:15], v[2:3], v[4:5], a[2:17] ; encoding: [0x00,0x80,0xbf,0xd3,0x02,0x09,0x0a,0x04]
  64 0x00,0x80,0xbf,0xd3,0x02,0x09,0x0a,0x04
  65
  66 # GFX940: v_mfma_f32_16x16x32_bf8_bf8 v[0:3], v[2:3], v[4:5], v[0:3] ; encoding: [0x00,0x00,0xf0,0xd3,0x02,0x09,0x02,0x04]
  67 0x00,0x00,0xf0,0xd3,0x02,0x09,0x02,0x04
  68
  69 # GFX940: v_mfma_f32_16x16x32_bf8_bf8 a[0:3], v[2:3], v[4:5], a[0:3] ; encoding: [0x00,0x80,0xf0,0xd3,0x02,0x09,0x02,0x04]
  70 0x00,0x80,0xf0,0xd3,0x02,0x09,0x02,0x04
  71
  72 # GFX940: v_mfma_f32_16x16x32_bf8_bf8 a[0:3], v[2:3], v[4:5], a[0:3] blgp:5 ; encoding: [0x00,0x80,0xf0,0xd3,0x02,0x09,0x02,0xa4]
  73 0x00,0x80,0xf0,0xd3,0x02,0x09,0x02,0xa4
  74
  75 # GFX940: v_mfma_f32_16x16x32_bf8_fp8 v[0:3], v[2:3], v[4:5], v[0:3] ; encoding: [0x00,0x00,0xf1,0xd3,0x02,0x09,0x02,0x04]
  76 0x00,0x00,0xf1,0xd3,0x02,0x09,0x02,0x04
  77
  78 # GFX940: v_mfma_f32_16x16x32_bf8_fp8 a[0:3], v[2:3], v[4:5], a[0:3] ; encoding: [0x00,0x80,0xf1,0xd3,0x02,0x09,0x02,0x04]
  79 0x00,0x80,0xf1,0xd3,0x02,0x09,0x02,0x04
  80
  81 # GFX940: v_mfma_f32_16x16x32_bf8_fp8 a[0:3], v[2:3], v[4:5], a[0:3] blgp:5 ; encoding: [0x00,0x80,0xf1,0xd3,0x02,0x09,0x02,0xa4]
  82 0x00,0x80,0xf1,0xd3,0x02,0x09,0x02,0xa4
  83
  84 # GFX940: v_mfma_f32_16x16x32_fp8_bf8 v[0:3], v[2:3], v[4:5], v[0:3] ; encoding: [0x00,0x00,0xf2,0xd3,0x02,0x09,0x02,0x04]
  85 0x00,0x00,0xf2,0xd3,0x02,0x09,0x02,0x04
  86
  87 # GFX940: v_mfma_f32_16x16x32_fp8_bf8 a[0:3], v[2:3], v[4:5], a[0:3] ; encoding: [0x00,0x80,0xf2,0xd3,0x02,0x09,0x02,0x04]
  88 0x00,0x80,0xf2,0xd3,0x02,0x09,0x02,0x04
  89
  90 # GFX940: v_mfma_f32_16x16x32_fp8_bf8 a[0:3], v[2:3], v[4:5], a[0:3] blgp:5 ; encoding: [0x00,0x80,0xf2,0xd3,0x02,0x09,0x02,0xa4]
  91 0x00,0x80,0xf2,0xd3,0x02,0x09,0x02,0xa4
  92
  93 # GFX940: v_mfma_f32_16x16x32_fp8_fp8 v[0:3], v[2:3], v[4:5], v[0:3] ; encoding: [0x00,0x00,0xf3,0xd3,0x02,0x09,0x02,0x04]
  94 0x00,0x00,0xf3,0xd3,0x02,0x09,0x02,0x04
  95
  96 # GFX940: v_mfma_f32_16x16x32_fp8_fp8 a[0:3], v[2:3], v[4:5], a[0:3] ; encoding: [0x00,0x80,0xf3,0xd3,0x02,0x09,0x02,0x04]
  97 0x00,0x80,0xf3,0xd3,0x02,0x09,0x02,0x04
  98
  99 # GFX940: v_mfma_f32_16x16x32_fp8_fp8 a[0:3], v[2:3], v[4:5], a[0:3] blgp:5 ; encoding: [0x00,0x80,0xf3,0xd3,0x02,0x09,0x02,0xa4]
 100 0x00,0x80,0xf3,0xd3,0x02,0x09,0x02,0xa4
 101
 102 # GFX940: v_mfma_f32_32x32x16_bf8_bf8 v[0:15], v[2:3], v[4:5], v[0:15] ; encoding: [0x00,0x00,0xf4,0xd3,0x02,0x09,0x02,0x04]
 103 0x00,0x00,0xf4,0xd3,0x02,0x09,0x02,0x04
 104
 105 # GFX940: v_mfma_f32_32x32x16_bf8_bf8 a[0:15], v[2:3], v[4:5], a[0:15] ; encoding: [0x00,0x80,0xf4,0xd3,0x02,0x09,0x02,0x04]
 106 0x00,0x80,0xf4,0xd3,0x02,0x09,0x02,0x04
 107
 108 # GFX940: v_mfma_f32_32x32x16_bf8_bf8 a[0:15], v[2:3], v[4:5], a[0:15] blgp:5 ; encoding: [0x00,0x80,0xf4,0xd3,0x02,0x09,0x02,0xa4]
 109 0x00,0x80,0xf4,0xd3,0x02,0x09,0x02,0xa4
 110
 111 # GFX940: v_mfma_f32_32x32x16_bf8_fp8 v[0:15], v[2:3], v[4:5], v[0:15] ; encoding: [0x00,0x00,0xf5,0xd3,0x02,0x09,0x02,0x04]
 112 0x00,0x00,0xf5,0xd3,0x02,0x09,0x02,0x04
 113
 114 # GFX940: v_mfma_f32_32x32x16_bf8_fp8 a[0:15], v[2:3], v[4:5], a[0:15] ; encoding: [0x00,0x80,0xf5,0xd3,0x02,0x09,0x02,0x04]
 115 0x00,0x80,0xf5,0xd3,0x02,0x09,0x02,0x04
 116
 117 # GFX940: v_mfma_f32_32x32x16_bf8_fp8 a[0:15], v[2:3], v[4:5], a[0:15] blgp:5 ; encoding: [0x00,0x80,0xf5,0xd3,0x02,0x09,0x02,0xa4]
 118 0x00,0x80,0xf5,0xd3,0x02,0x09,0x02,0xa4
 119
 120 # GFX940: v_mfma_f32_32x32x16_fp8_bf8 v[0:15], v[2:3], v[4:5], v[0:15] ; encoding: [0x00,0x00,0xf6,0xd3,0x02,0x09,0x02,0x04]
 121 0x00,0x00,0xf6,0xd3,0x02,0x09,0x02,0x04
 122
 123 # GFX940: v_mfma_f32_32x32x16_fp8_bf8 a[0:15], v[2:3], v[4:5], a[0:15] ; encoding: [0x00,0x80,0xf6,0xd3,0x02,0x09,0x02,0x04]
 124 0x00,0x80,0xf6,0xd3,0x02,0x09,0x02,0x04
 125
 126 # GFX940: v_mfma_f32_32x32x16_fp8_bf8 a[0:15], v[2:3], v[4:5], a[0:15] blgp:5 ; encoding: [0x00,0x80,0xf6,0xd3,0x02,0x09,0x02,0xa4]
 127 0x00,0x80,0xf6,0xd3,0x02,0x09,0x02,0xa4
 128
 129 # GFX940: v_mfma_f32_32x32x16_fp8_fp8 v[0:15], v[2:3], v[4:5], v[0:15] ; encoding: [0x00,0x00,0xf7,0xd3,0x02,0x09,0x02,0x04]
 130 0x00,0x00,0xf7,0xd3,0x02,0x09,0x02,0x04
 131
 132 # GFX940: v_mfma_f32_32x32x16_fp8_fp8 a[0:15], v[2:3], v[4:5], a[0:15] ; encoding: [0x00,0x80,0xf7,0xd3,0x02,0x09,0x02,0x04]
 133 0x00,0x80,0xf7,0xd3,0x02,0x09,0x02,0x04
 134
 135 # GFX940: v_mfma_f32_32x32x16_fp8_fp8 a[0:15], v[2:3], v[4:5], a[0:15] blgp:5 ; encoding: [0x00,0x80,0xf7,0xd3,0x02,0x09,0x02,0xa4]
 136 0x00,0x80,0xf7,0xd3,0x02,0x09,0x02,0xa4
 137
 138 # GFX940: v_smfmac_f32_16x16x32_f16 v[10:13], a[2:3], v[4:7], v0 cbsz:3 abid:1 ; encoding: [0x0a,0x0b,0xe2,0xd3,0x02,0x09,0x02,0x0c]
 139 0x0a,0x0b,0xe2,0xd3,0x02,0x09,0x02,0x0c
 140
 141 # GFX940: v_smfmac_f32_16x16x32_f16 a[10:13], v[2:3], a[4:7], v1 ; encoding: [0x0a,0x80,0xe2,0xd3,0x02,0x09,0x06,0x14]
 142 0x0a,0x80,0xe2,0xd3,0x02,0x09,0x06,0x14
 143
 144 # GFX940: v_smfmac_f32_16x16x32_f16 v[252:255], v[2:3], v[4:7], v3 ; encoding: [0xfc,0x00,0xe2,0xd3,0x02,0x09,0x0e,0x04]
 145 0xfc,0x00,0xe2,0xd3,0x02,0x09,0x0e,0x04
 146
 147 # GFX940: v_smfmac_f32_16x16x32_f16 a[252:255], v[2:3], v[4:7], v3 ; encoding: [0xfc,0x80,0xe2,0xd3,0x02,0x09,0x0e,0x04]
 148 0xfc,0x80,0xe2,0xd3,0x02,0x09,0x0e,0x04
 149
 150 # GFX940: v_smfmac_f32_16x16x32_f16 v[10:13], v[254:255], v[4:7], v3 ; encoding: [0x0a,0x00,0xe2,0xd3,0xfe,0x09,0x0e,0x04]
 151 0x0a,0x00,0xe2,0xd3,0xfe,0x09,0x0e,0x04
 152
 153 # GFX940: v_smfmac_f32_16x16x32_f16 v[10:13], a[254:255], v[4:7], v3 ; encoding: [0x0a,0x00,0xe2,0xd3,0xfe,0x09,0x0e,0x0c]
 154 0x0a,0x00,0xe2,0xd3,0xfe,0x09,0x0e,0x0c
 155
 156 # GFX940: v_smfmac_f32_16x16x32_f16 v[10:13], v[2:3], v[252:255], v3 ; encoding: [0x0a,0x00,0xe2,0xd3,0x02,0xf9,0x0f,0x04]
 157 0x0a,0x00,0xe2,0xd3,0x02,0xf9,0x0f,0x04
 158
 159 # GFX940: v_smfmac_f32_16x16x32_f16 v[10:13], v[2:3], a[252:255], v3 ; encoding: [0x0a,0x00,0xe2,0xd3,0x02,0xf9,0x0f,0x14]
 160 0x0a,0x00,0xe2,0xd3,0x02,0xf9,0x0f,0x14
 161
 162 # GFX940: v_smfmac_f32_16x16x32_f16 v[10:13], v[2:3], v[4:7], v255 ; encoding: [0x0a,0x00,0xe2,0xd3,0x02,0x09,0xfe,0x07]
 163 0x0a,0x00,0xe2,0xd3,0x02,0x09,0xfe,0x07
 164
 165 # GFX940: v_smfmac_f32_16x16x32_f16 v[10:13], v[2:3], v[4:7], v3 cbsz:2 ; encoding: [0x0a,0x02,0xe2,0xd3,0x02,0x09,0x0e,0x04]
 166 0x0a,0x02,0xe2,0xd3,0x02,0x09,0x0e,0x04
 167
 168 # GFX940: v_smfmac_f32_16x16x32_f16 v[10:13], v[2:3], v[4:7], v3 cbsz:7 ; encoding: [0x0a,0x07,0xe2,0xd3,0x02,0x09,0x0e,0x04]
 169 0x0a,0x07,0xe2,0xd3,0x02,0x09,0x0e,0x04
 170
 171 # GFX940: v_smfmac_f32_16x16x32_f16 v[10:13], v[2:3], v[4:7], v3 abid:1 ; encoding: [0x0a,0x08,0xe2,0xd3,0x02,0x09,0x0e,0x04]
 172 0x0a,0x08,0xe2,0xd3,0x02,0x09,0x0e,0x04
 173
 174 # GFX940: v_smfmac_f32_16x16x32_f16 v[10:13], v[2:3], v[4:7], v3 abid:7 ; encoding: [0x0a,0x38,0xe2,0xd3,0x02,0x09,0x0e,0x04]
 175 0x0a,0x38,0xe2,0xd3,0x02,0x09,0x0e,0x04
 176
 177 # GFX940: v_smfmac_f32_16x16x32_f16 v[10:13], v[2:3], v[4:7], v3 abid:15 ; encoding: [0x0a,0x78,0xe2,0xd3,0x02,0x09,0x0e,0x04]
 178 0x0a,0x78,0xe2,0xd3,0x02,0x09,0x0e,0x04
 179
 180 # GFX940: v_smfmac_f32_32x32x16_f16 v[10:25], a[2:3], v[4:7], v2 cbsz:3 abid:1 ; encoding: [0x0a,0x0b,0xe4,0xd3,0x02,0x09,0x0a,0x0c]
 181 0x0a,0x0b,0xe4,0xd3,0x02,0x09,0x0a,0x0c
 182
 183 # GFX940: v_smfmac_f32_32x32x16_f16 a[10:25], v[2:3], a[4:7], v3 ; encoding: [0x0a,0x80,0xe4,0xd3,0x02,0x09,0x0e,0x14]
 184 0x0a,0x80,0xe4,0xd3,0x02,0x09,0x0e,0x14
 185
 186 # GFX940: v_smfmac_f32_32x32x16_f16 v[240:255], v[2:3], v[4:7], v3 ; encoding: [0xf0,0x00,0xe4,0xd3,0x02,0x09,0x0e,0x04]
 187 0xf0,0x00,0xe4,0xd3,0x02,0x09,0x0e,0x04
 188
 189 # GFX940: v_smfmac_f32_32x32x16_f16 a[240:255], v[2:3], v[4:7], v3 ; encoding: [0xf0,0x80,0xe4,0xd3,0x02,0x09,0x0e,0x04]
 190 0xf0,0x80,0xe4,0xd3,0x02,0x09,0x0e,0x04
 191
 192 # GFX940: v_smfmac_f32_32x32x16_f16 v[10:25], v[254:255], v[4:7], v3 ; encoding: [0x0a,0x00,0xe4,0xd3,0xfe,0x09,0x0e,0x04]
 193 0x0a,0x00,0xe4,0xd3,0xfe,0x09,0x0e,0x04
 194
 195 # GFX940: v_smfmac_f32_32x32x16_f16 v[10:25], a[254:255], v[4:7], v3 ; encoding: [0x0a,0x00,0xe4,0xd3,0xfe,0x09,0x0e,0x0c]
 196 0x0a,0x00,0xe4,0xd3,0xfe,0x09,0x0e,0x0c
 197
 198 # GFX940: v_smfmac_f32_32x32x16_f16 v[10:25], v[2:3], v[252:255], v3 ; encoding: [0x0a,0x00,0xe4,0xd3,0x02,0xf9,0x0f,0x04]
 199 0x0a,0x00,0xe4,0xd3,0x02,0xf9,0x0f,0x04
 200
 201 # GFX940: v_smfmac_f32_32x32x16_f16 v[10:25], v[2:3], a[252:255], v3 ; encoding: [0x0a,0x00,0xe4,0xd3,0x02,0xf9,0x0f,0x14]
 202 0x0a,0x00,0xe4,0xd3,0x02,0xf9,0x0f,0x14
 203
 204 # GFX940: v_smfmac_f32_32x32x16_f16 v[10:25], v[2:3], v[4:7], v255 ; encoding: [0x0a,0x00,0xe4,0xd3,0x02,0x09,0xfe,0x07]
 205 0x0a,0x00,0xe4,0xd3,0x02,0x09,0xfe,0x07
 206
 207 # GFX940: v_smfmac_f32_32x32x16_f16 v[10:25], v[2:3], v[4:7], v3 cbsz:2 ; encoding: [0x0a,0x02,0xe4,0xd3,0x02,0x09,0x0e,0x04]
 208 0x0a,0x02,0xe4,0xd3,0x02,0x09,0x0e,0x04
 209
 210 # GFX940: v_smfmac_f32_32x32x16_f16 v[10:25], v[2:3], v[4:7], v3 cbsz:7 ; encoding: [0x0a,0x07,0xe4,0xd3,0x02,0x09,0x0e,0x04]
 211 0x0a,0x07,0xe4,0xd3,0x02,0x09,0x0e,0x04
 212
 213 # GFX940: v_smfmac_f32_32x32x16_f16 v[10:25], v[2:3], v[4:7], v3 abid:1 ; encoding: [0x0a,0x08,0xe4,0xd3,0x02,0x09,0x0e,0x04]
 214 0x0a,0x08,0xe4,0xd3,0x02,0x09,0x0e,0x04
 215
 216 # GFX940: v_smfmac_f32_32x32x16_f16 v[10:25], v[2:3], v[4:7], v3 abid:7 ; encoding: [0x0a,0x38,0xe4,0xd3,0x02,0x09,0x0e,0x04]
 217 0x0a,0x38,0xe4,0xd3,0x02,0x09,0x0e,0x04
 218
 219 # GFX940: v_smfmac_f32_32x32x16_f16 v[10:25], v[2:3], v[4:7], v3 abid:15 ; encoding: [0x0a,0x78,0xe4,0xd3,0x02,0x09,0x0e,0x04]
 220 0x0a,0x78,0xe4,0xd3,0x02,0x09,0x0e,0x04
 221
 222 # GFX940: v_smfmac_f32_16x16x32_bf16 v[10:13], a[2:3], v[4:7], v4 cbsz:3 abid:1 ; encoding: [0x0a,0x0b,0xe6,0xd3,0x02,0x09,0x12,0x0c]
 223 0x0a,0x0b,0xe6,0xd3,0x02,0x09,0x12,0x0c
 224
 225 # GFX940: v_smfmac_f32_16x16x32_bf16 a[10:13], v[2:3], a[4:7], v5 ; encoding: [0x0a,0x80,0xe6,0xd3,0x02,0x09,0x16,0x14]
 226 0x0a,0x80,0xe6,0xd3,0x02,0x09,0x16,0x14
 227
 228 # GFX940: v_smfmac_f32_16x16x32_bf16 v[252:255], v[2:3], v[4:7], v3 ; encoding: [0xfc,0x00,0xe6,0xd3,0x02,0x09,0x0e,0x04]
 229 0xfc,0x00,0xe6,0xd3,0x02,0x09,0x0e,0x04
 230
 231 # GFX940: v_smfmac_f32_16x16x32_bf16 a[252:255], v[2:3], v[4:7], v3 ; encoding: [0xfc,0x80,0xe6,0xd3,0x02,0x09,0x0e,0x04]
 232 0xfc,0x80,0xe6,0xd3,0x02,0x09,0x0e,0x04
 233
 234 # GFX940: v_smfmac_f32_16x16x32_bf16 v[10:13], v[254:255], v[4:7], v3 ; encoding: [0x0a,0x00,0xe6,0xd3,0xfe,0x09,0x0e,0x04]
 235 0x0a,0x00,0xe6,0xd3,0xfe,0x09,0x0e,0x04
 236
 237 # GFX940: v_smfmac_f32_16x16x32_bf16 v[10:13], a[254:255], v[4:7], v3 ; encoding: [0x0a,0x00,0xe6,0xd3,0xfe,0x09,0x0e,0x0c]
 238 0x0a,0x00,0xe6,0xd3,0xfe,0x09,0x0e,0x0c
 239
 240 # GFX940: v_smfmac_f32_16x16x32_bf16 v[10:13], v[2:3], v[252:255], v3 ; encoding: [0x0a,0x00,0xe6,0xd3,0x02,0xf9,0x0f,0x04]
 241 0x0a,0x00,0xe6,0xd3,0x02,0xf9,0x0f,0x04
 242
 243 # GFX940: v_smfmac_f32_16x16x32_bf16 v[10:13], v[2:3], a[252:255], v3 ; encoding: [0x0a,0x00,0xe6,0xd3,0x02,0xf9,0x0f,0x14]
 244 0x0a,0x00,0xe6,0xd3,0x02,0xf9,0x0f,0x14
 245
 246 # GFX940: v_smfmac_f32_16x16x32_bf16 v[10:13], v[2:3], v[4:7], v255 ; encoding: [0x0a,0x00,0xe6,0xd3,0x02,0x09,0xfe,0x07]
 247 0x0a,0x00,0xe6,0xd3,0x02,0x09,0xfe,0x07
 248
 249 # GFX940: v_smfmac_f32_16x16x32_bf16 v[10:13], v[2:3], v[4:7], v3 cbsz:2 ; encoding: [0x0a,0x02,0xe6,0xd3,0x02,0x09,0x0e,0x04]
 250 0x0a,0x02,0xe6,0xd3,0x02,0x09,0x0e,0x04
 251
 252 # GFX940: v_smfmac_f32_16x16x32_bf16 v[10:13], v[2:3], v[4:7], v3 cbsz:7 ; encoding: [0x0a,0x07,0xe6,0xd3,0x02,0x09,0x0e,0x04]
 253 0x0a,0x07,0xe6,0xd3,0x02,0x09,0x0e,0x04
 254
 255 # GFX940: v_smfmac_f32_16x16x32_bf16 v[10:13], v[2:3], v[4:7], v3 abid:1 ; encoding: [0x0a,0x08,0xe6,0xd3,0x02,0x09,0x0e,0x04]
 256 0x0a,0x08,0xe6,0xd3,0x02,0x09,0x0e,0x04
 257
 258 # GFX940: v_smfmac_f32_16x16x32_bf16 v[10:13], v[2:3], v[4:7], v3 abid:7 ; encoding: [0x0a,0x38,0xe6,0xd3,0x02,0x09,0x0e,0x04]
 259 0x0a,0x38,0xe6,0xd3,0x02,0x09,0x0e,0x04
 260
 261 # GFX940: v_smfmac_f32_16x16x32_bf16 v[10:13], v[2:3], v[4:7], v3 abid:15 ; encoding: [0x0a,0x78,0xe6,0xd3,0x02,0x09,0x0e,0x04]
 262 0x0a,0x78,0xe6,0xd3,0x02,0x09,0x0e,0x04
 263
 264 # GFX940: v_smfmac_f32_32x32x16_bf16 v[10:25], a[2:3], v[4:7], v6 cbsz:3 abid:1 ; encoding: [0x0a,0x0b,0xe8,0xd3,0x02,0x09,0x1a,0x0c]
 265 0x0a,0x0b,0xe8,0xd3,0x02,0x09,0x1a,0x0c
 266
 267 # GFX940: v_smfmac_f32_32x32x16_bf16 a[10:25], v[2:3], a[4:7], v7 ; encoding: [0x0a,0x80,0xe8,0xd3,0x02,0x09,0x1e,0x14]
 268 0x0a,0x80,0xe8,0xd3,0x02,0x09,0x1e,0x14
 269
 270 # GFX940: v_smfmac_f32_32x32x16_bf16 v[10:25], a[2:3], v[4:7], v8 cbsz:3 abid:1 ; encoding: [0x0a,0x0b,0xe8,0xd3,0x02,0x09,0x22,0x0c]
 271 0x0a,0x0b,0xe8,0xd3,0x02,0x09,0x22,0x0c
 272
 273 # GFX940: v_smfmac_f32_32x32x16_bf16 a[10:25], v[2:3], a[4:7], v9 ; encoding: [0x0a,0x80,0xe8,0xd3,0x02,0x09,0x26,0x14]
 274 0x0a,0x80,0xe8,0xd3,0x02,0x09,0x26,0x14
 275
 276 # GFX940: v_smfmac_f32_32x32x16_bf16 v[240:255], v[2:3], v[4:7], v3 ; encoding: [0xf0,0x00,0xe8,0xd3,0x02,0x09,0x0e,0x04]
 277 0xf0,0x00,0xe8,0xd3,0x02,0x09,0x0e,0x04
 278
 279 # GFX940: v_smfmac_f32_32x32x16_bf16 a[240:255], v[2:3], v[4:7], v3 ; encoding: [0xf0,0x80,0xe8,0xd3,0x02,0x09,0x0e,0x04]
 280 0xf0,0x80,0xe8,0xd3,0x02,0x09,0x0e,0x04
 281
 282 # GFX940: v_smfmac_f32_32x32x16_bf16 v[10:25], v[254:255], v[4:7], v3 ; encoding: [0x0a,0x00,0xe8,0xd3,0xfe,0x09,0x0e,0x04]
 283 0x0a,0x00,0xe8,0xd3,0xfe,0x09,0x0e,0x04
 284
 285 # GFX940: v_smfmac_f32_32x32x16_bf16 v[10:25], a[254:255], v[4:7], v3 ; encoding: [0x0a,0x00,0xe8,0xd3,0xfe,0x09,0x0e,0x0c]
 286 0x0a,0x00,0xe8,0xd3,0xfe,0x09,0x0e,0x0c
 287
 288 # GFX940: v_smfmac_f32_32x32x16_bf16 v[10:25], v[2:3], v[252:255], v3 ; encoding: [0x0a,0x00,0xe8,0xd3,0x02,0xf9,0x0f,0x04]
 289 0x0a,0x00,0xe8,0xd3,0x02,0xf9,0x0f,0x04
 290
 291 # GFX940: v_smfmac_f32_32x32x16_bf16 v[10:25], v[2:3], a[252:255], v3 ; encoding: [0x0a,0x00,0xe8,0xd3,0x02,0xf9,0x0f,0x14]
 292 0x0a,0x00,0xe8,0xd3,0x02,0xf9,0x0f,0x14
 293
 294 # GFX940: v_smfmac_f32_32x32x16_bf16 v[10:25], v[2:3], v[4:7], v255 ; encoding: [0x0a,0x00,0xe8,0xd3,0x02,0x09,0xfe,0x07]
 295 0x0a,0x00,0xe8,0xd3,0x02,0x09,0xfe,0x07
 296
 297 # GFX940: v_smfmac_f32_32x32x16_bf16 v[10:25], v[2:3], v[4:7], v3 cbsz:2 ; encoding: [0x0a,0x02,0xe8,0xd3,0x02,0x09,0x0e,0x04]
 298 0x0a,0x02,0xe8,0xd3,0x02,0x09,0x0e,0x04
 299
 300 # GFX940: v_smfmac_f32_32x32x16_bf16 v[10:25], v[2:3], v[4:7], v3 cbsz:7 ; encoding: [0x0a,0x07,0xe8,0xd3,0x02,0x09,0x0e,0x04]
 301 0x0a,0x07,0xe8,0xd3,0x02,0x09,0x0e,0x04
 302
 303 # GFX940: v_smfmac_f32_32x32x16_bf16 v[10:25], v[2:3], v[4:7], v3 abid:1 ; encoding: [0x0a,0x08,0xe8,0xd3,0x02,0x09,0x0e,0x04]
 304 0x0a,0x08,0xe8,0xd3,0x02,0x09,0x0e,0x04
 305
 306 # GFX940: v_smfmac_f32_32x32x16_bf16 v[10:25], v[2:3], v[4:7], v3 abid:7 ; encoding: [0x0a,0x38,0xe8,0xd3,0x02,0x09,0x0e,0x04]
 307 0x0a,0x38,0xe8,0xd3,0x02,0x09,0x0e,0x04
 308
 309 # GFX940: v_smfmac_f32_32x32x16_bf16 v[10:25], v[2:3], v[4:7], v3 abid:15 ; encoding: [0x0a,0x78,0xe8,0xd3,0x02,0x09,0x0e,0x04]
 310 0x0a,0x78,0xe8,0xd3,0x02,0x09,0x0e,0x04
 311
 312 # GFX940: v_smfmac_i32_16x16x64_i8 v[10:13], a[2:3], v[4:7], v10 cbsz:3 abid:1 ; encoding: [0x0a,0x0b,0xea,0xd3,0x02,0x09,0x2a,0x0c]
 313 0x0a,0x0b,0xea,0xd3,0x02,0x09,0x2a,0x0c
 314
 315 # GFX940: v_smfmac_i32_16x16x64_i8 a[10:13], v[2:3], a[4:7], v11 ; encoding: [0x0a,0x80,0xea,0xd3,0x02,0x09,0x2e,0x14]
 316 0x0a,0x80,0xea,0xd3,0x02,0x09,0x2e,0x14
 317
 318 # GFX940: v_smfmac_i32_16x16x64_i8 v[252:255], v[2:3], v[4:7], v3 ; encoding: [0xfc,0x00,0xea,0xd3,0x02,0x09,0x0e,0x04]
 319 0xfc,0x00,0xea,0xd3,0x02,0x09,0x0e,0x04
 320
 321 # GFX940: v_smfmac_i32_16x16x64_i8 a[252:255], v[2:3], v[4:7], v3 ; encoding: [0xfc,0x80,0xea,0xd3,0x02,0x09,0x0e,0x04]
 322 0xfc,0x80,0xea,0xd3,0x02,0x09,0x0e,0x04
 323
 324 # GFX940: v_smfmac_i32_16x16x64_i8 v[10:13], v[254:255], v[4:7], v3 ; encoding: [0x0a,0x00,0xea,0xd3,0xfe,0x09,0x0e,0x04]
 325 0x0a,0x00,0xea,0xd3,0xfe,0x09,0x0e,0x04
 326
 327 # GFX940: v_smfmac_i32_16x16x64_i8 v[10:13], a[254:255], v[4:7], v3 ; encoding: [0x0a,0x00,0xea,0xd3,0xfe,0x09,0x0e,0x0c]
 328 0x0a,0x00,0xea,0xd3,0xfe,0x09,0x0e,0x0c
 329
 330 # GFX940: v_smfmac_i32_16x16x64_i8 v[10:13], v[2:3], v[252:255], v3 ; encoding: [0x0a,0x00,0xea,0xd3,0x02,0xf9,0x0f,0x04]
 331 0x0a,0x00,0xea,0xd3,0x02,0xf9,0x0f,0x04
 332
 333 # GFX940: v_smfmac_i32_16x16x64_i8 v[10:13], v[2:3], a[252:255], v3 ; encoding: [0x0a,0x00,0xea,0xd3,0x02,0xf9,0x0f,0x14]
 334 0x0a,0x00,0xea,0xd3,0x02,0xf9,0x0f,0x14
 335
 336 # GFX940: v_smfmac_i32_16x16x64_i8 v[10:13], v[2:3], v[4:7], v255 ; encoding: [0x0a,0x00,0xea,0xd3,0x02,0x09,0xfe,0x07]
 337 0x0a,0x00,0xea,0xd3,0x02,0x09,0xfe,0x07
 338
 339 # GFX940: v_smfmac_i32_16x16x64_i8 v[10:13], v[2:3], v[4:7], v3 cbsz:2 ; encoding: [0x0a,0x02,0xea,0xd3,0x02,0x09,0x0e,0x04]
 340 0x0a,0x02,0xea,0xd3,0x02,0x09,0x0e,0x04
 341
 342 # GFX940: v_smfmac_i32_16x16x64_i8 v[10:13], v[2:3], v[4:7], v3 cbsz:7 ; encoding: [0x0a,0x07,0xea,0xd3,0x02,0x09,0x0e,0x04]
 343 0x0a,0x07,0xea,0xd3,0x02,0x09,0x0e,0x04
 344
 345 # GFX940: v_smfmac_i32_16x16x64_i8 v[10:13], v[2:3], v[4:7], v3 abid:1 ; encoding: [0x0a,0x08,0xea,0xd3,0x02,0x09,0x0e,0x04]
 346 0x0a,0x08,0xea,0xd3,0x02,0x09,0x0e,0x04
 347
 348 # GFX940: v_smfmac_i32_16x16x64_i8 v[10:13], v[2:3], v[4:7], v3 abid:7 ; encoding: [0x0a,0x38,0xea,0xd3,0x02,0x09,0x0e,0x04]
 349 0x0a,0x38,0xea,0xd3,0x02,0x09,0x0e,0x04
 350
 351 # GFX940: v_smfmac_i32_16x16x64_i8 v[10:13], v[2:3], v[4:7], v3 abid:15 ; encoding: [0x0a,0x78,0xea,0xd3,0x02,0x09,0x0e,0x04]
 352 0x0a,0x78,0xea,0xd3,0x02,0x09,0x0e,0x04
 353
 354 # GFX940: v_smfmac_i32_32x32x32_i8 v[10:25], a[2:3], v[4:7], v12 cbsz:3 abid:1 ; encoding: [0x0a,0x0b,0xec,0xd3,0x02,0x09,0x32,0x0c]
 355 0x0a,0x0b,0xec,0xd3,0x02,0x09,0x32,0x0c
 356
 357 # GFX940: v_smfmac_i32_32x32x32_i8 a[10:25], v[2:3], a[4:7], v13 ; encoding: [0x0a,0x80,0xec,0xd3,0x02,0x09,0x36,0x14]
 358 0x0a,0x80,0xec,0xd3,0x02,0x09,0x36,0x14
 359
 360 # GFX940: v_smfmac_i32_32x32x32_i8 v[240:255], v[2:3], v[4:7], v3 ; encoding: [0xf0,0x00,0xec,0xd3,0x02,0x09,0x0e,0x04]
 361 0xf0,0x00,0xec,0xd3,0x02,0x09,0x0e,0x04
 362
 363 # GFX940: v_smfmac_i32_32x32x32_i8 a[240:255], v[2:3], v[4:7], v3 ; encoding: [0xf0,0x80,0xec,0xd3,0x02,0x09,0x0e,0x04]
 364 0xf0,0x80,0xec,0xd3,0x02,0x09,0x0e,0x04
 365
 366 # GFX940: v_smfmac_i32_32x32x32_i8 v[10:25], v[254:255], v[4:7], v3 ; encoding: [0x0a,0x00,0xec,0xd3,0xfe,0x09,0x0e,0x04]
 367 0x0a,0x00,0xec,0xd3,0xfe,0x09,0x0e,0x04
 368
 369 # GFX940: v_smfmac_i32_32x32x32_i8 v[10:25], a[254:255], v[4:7], v3 ; encoding: [0x0a,0x00,0xec,0xd3,0xfe,0x09,0x0e,0x0c]
 370 0x0a,0x00,0xec,0xd3,0xfe,0x09,0x0e,0x0c
 371
 372 # GFX940: v_smfmac_i32_32x32x32_i8 v[10:25], v[2:3], v[252:255], v3 ; encoding: [0x0a,0x00,0xec,0xd3,0x02,0xf9,0x0f,0x04]
 373 0x0a,0x00,0xec,0xd3,0x02,0xf9,0x0f,0x04
 374
 375 # GFX940: v_smfmac_i32_32x32x32_i8 v[10:25], v[2:3], a[252:255], v3 ; encoding: [0x0a,0x00,0xec,0xd3,0x02,0xf9,0x0f,0x14]
 376 0x0a,0x00,0xec,0xd3,0x02,0xf9,0x0f,0x14
 377
 378 # GFX940: v_smfmac_i32_32x32x32_i8 v[10:25], v[2:3], v[4:7], v255 ; encoding: [0x0a,0x00,0xec,0xd3,0x02,0x09,0xfe,0x07]
 379 0x0a,0x00,0xec,0xd3,0x02,0x09,0xfe,0x07
 380
 381 # GFX940: v_smfmac_i32_32x32x32_i8 v[10:25], v[2:3], v[4:7], v3 cbsz:2 ; encoding: [0x0a,0x02,0xec,0xd3,0x02,0x09,0x0e,0x04]
 382 0x0a,0x02,0xec,0xd3,0x02,0x09,0x0e,0x04
 383
 384 # GFX940: v_smfmac_i32_32x32x32_i8 v[10:25], v[2:3], v[4:7], v3 cbsz:7 ; encoding: [0x0a,0x07,0xec,0xd3,0x02,0x09,0x0e,0x04]
 385 0x0a,0x07,0xec,0xd3,0x02,0x09,0x0e,0x04
 386
 387 # GFX940: v_smfmac_i32_32x32x32_i8 v[10:25], v[2:3], v[4:7], v3 abid:1 ; encoding: [0x0a,0x08,0xec,0xd3,0x02,0x09,0x0e,0x04]
 388 0x0a,0x08,0xec,0xd3,0x02,0x09,0x0e,0x04
 389
 390 # GFX940: v_smfmac_i32_32x32x32_i8 v[10:25], v[2:3], v[4:7], v3 abid:7 ; encoding: [0x0a,0x38,0xec,0xd3,0x02,0x09,0x0e,0x04]
 391 0x0a,0x38,0xec,0xd3,0x02,0x09,0x0e,0x04
 392
 393 # GFX940: v_smfmac_i32_32x32x32_i8 v[10:25], v[2:3], v[4:7], v3 abid:15 ; encoding: [0x0a,0x78,0xec,0xd3,0x02,0x09,0x0e,0x04]
 394 0x0a,0x78,0xec,0xd3,0x02,0x09,0x0e,0x04
 395
 396 # GFX940: v_smfmac_f32_16x16x64_bf8_bf8 v[0:3], a[2:3], v[4:7], v1 cbsz:3 abid:1 ; encoding: [0x00,0x0b,0xf8,0xd3,0x02,0x09,0x06,0x0c]
 397 0x00,0x0b,0xf8,0xd3,0x02,0x09,0x06,0x0c
 398
 399 # GFX940: v_smfmac_f32_16x16x64_bf8_bf8 a[0:3], v[2:3], a[4:7], v1 ; encoding: [0x00,0x80,0xf8,0xd3,0x02,0x09,0x06,0x14]
 400 0x00,0x80,0xf8,0xd3,0x02,0x09,0x06,0x14
 401
 402 # GFX940: v_smfmac_f32_16x16x64_bf8_fp8 v[0:3], a[2:3], v[4:7], v1 cbsz:3 abid:1 ; encoding: [0x00,0x0b,0xf9,0xd3,0x02,0x09,0x06,0x0c]
 403 0x00,0x0b,0xf9,0xd3,0x02,0x09,0x06,0x0c
 404
 405 # GFX940: v_smfmac_f32_16x16x64_bf8_fp8 a[0:3], v[2:3], a[4:7], v1 ; encoding: [0x00,0x80,0xf9,0xd3,0x02,0x09,0x06,0x14]
 406 0x00,0x80,0xf9,0xd3,0x02,0x09,0x06,0x14
 407
 408 # GFX940: v_smfmac_f32_16x16x64_fp8_bf8 v[0:3], a[2:3], v[4:7], v1 cbsz:3 abid:1 ; encoding: [0x00,0x0b,0xfa,0xd3,0x02,0x09,0x06,0x0c]
 409 0x00,0x0b,0xfa,0xd3,0x02,0x09,0x06,0x0c
 410
 411 # GFX940: v_smfmac_f32_16x16x64_fp8_bf8 a[0:3], v[2:3], a[4:7], v1 ; encoding: [0x00,0x80,0xfa,0xd3,0x02,0x09,0x06,0x14]
 412 0x00,0x80,0xfa,0xd3,0x02,0x09,0x06,0x14
 413
 414 # GFX940: v_smfmac_f32_16x16x64_fp8_fp8 v[0:3], a[2:3], v[4:7], v1 cbsz:3 abid:1 ; encoding: [0x00,0x0b,0xfb,0xd3,0x02,0x09,0x06,0x0c]
 415 0x00,0x0b,0xfb,0xd3,0x02,0x09,0x06,0x0c
 416
 417 # GFX940: v_smfmac_f32_16x16x64_fp8_fp8 a[0:3], v[2:3], a[4:7], v1 ; encoding: [0x00,0x80,0xfb,0xd3,0x02,0x09,0x06,0x14]
 418 0x00,0x80,0xfb,0xd3,0x02,0x09,0x06,0x14
 419
 420 # GFX940: v_smfmac_f32_32x32x32_bf8_bf8 v[0:15], a[2:3], v[4:7], v1 cbsz:3 abid:1 ; encoding: [0x00,0x0b,0xfc,0xd3,0x02,0x09,0x06,0x0c]
 421 0x00,0x0b,0xfc,0xd3,0x02,0x09,0x06,0x0c
 422
 423 # GFX940: v_smfmac_f32_32x32x32_bf8_bf8 a[0:15], v[2:3], a[4:7], v1 ; encoding: [0x00,0x80,0xfc,0xd3,0x02,0x09,0x06,0x14]
 424 0x00,0x80,0xfc,0xd3,0x02,0x09,0x06,0x14
 425
 426 # GFX940: v_smfmac_f32_32x32x32_bf8_fp8 v[0:15], a[2:3], v[4:7], v1 cbsz:3 abid:1 ; encoding: [0x00,0x0b,0xfd,0xd3,0x02,0x09,0x06,0x0c]
 427 0x00,0x0b,0xfd,0xd3,0x02,0x09,0x06,0x0c
 428
 429 # GFX940: v_smfmac_f32_32x32x32_bf8_fp8 a[0:15], v[2:3], a[4:7], v1 ; encoding: [0x00,0x80,0xfd,0xd3,0x02,0x09,0x06,0x14]
 430 0x00,0x80,0xfd,0xd3,0x02,0x09,0x06,0x14
 431
 432 # GFX940: v_smfmac_f32_32x32x32_fp8_bf8 v[0:15], a[2:3], v[4:7], v1 cbsz:3 abid:1 ; encoding: [0x00,0x0b,0xfe,0xd3,0x02,0x09,0x06,0x0c]
 433 0x00,0x0b,0xfe,0xd3,0x02,0x09,0x06,0x0c
 434
 435 # GFX940: v_smfmac_f32_32x32x32_fp8_bf8 a[0:15], v[2:3], a[4:7], v1 ; encoding: [0x00,0x80,0xfe,0xd3,0x02,0x09,0x06,0x14]
 436 0x00,0x80,0xfe,0xd3,0x02,0x09,0x06,0x14
 437
 438 # GFX940: v_smfmac_f32_32x32x32_fp8_fp8 v[0:15], a[2:3], v[4:7], v1 cbsz:3 abid:1 ; encoding: [0x00,0x0b,0xff,0xd3,0x02,0x09,0x06,0x0c]
 439 0x00,0x0b,0xff,0xd3,0x02,0x09,0x06,0x0c
 440
 441 # GFX940: v_smfmac_f32_32x32x32_fp8_fp8 a[0:15], v[2:3], a[4:7], v1 ; encoding: [0x00,0x80,0xff,0xd3,0x02,0x09,0x06,0x14]
 442 0x00,0x80,0xff,0xd3,0x02,0x09,0x06,0x14
 443
 444 # GFX940: v_mfma_f32_16x16x16_f16 v[10:13], v[2:3], v[4:5], v[6:9] ; encoding: [0x0a,0x00,0xcd,0xd3,0x02,0x09,0x1a,0x04]
 445 0x0a,0x00,0xcd,0xd3,0x02,0x09,0x1a,0x04
 446
 447 # GFX940: v_mfma_f32_16x16x16_f16 v[252:255], a[254:255], v[254:255], v[252:255] ; encoding: [0xfc,0x00,0xcd,0xd3,0xfe,0xfd,0xf3,0x0f]
 448 0xfc,0x00,0xcd,0xd3,0xfe,0xfd,0xf3,0x0f
 449
 450 # GFX940: v_mfma_f32_16x16x16_f16 v[252:255], v[254:255], a[254:255], v[252:255] cbsz:2 abid:7 blgp:3 ; encoding: [0xfc,0x3a,0xcd,0xd3,0xfe,0xfd,0xf3,0x77]
 451 0xfc,0x3a,0xcd,0xd3,0xfe,0xfd,0xf3,0x77
 452
 453 # GFX940: v_mfma_f32_16x16x16_f16 a[252:255], a[254:255], a[254:255], a[252:255] cbsz:7 abid:15 blgp:7 ; encoding: [0xfc,0xff,0xcd,0xd3,0xfe,0xfd,0xf3,0xff]
 454 0xfc,0xff,0xcd,0xd3,0xfe,0xfd,0xf3,0xff
 455
 456 # GFX940: v_mfma_f32_16x16x1_4b_f32 v[240:255], v1, v2, v[240:255] ; encoding: [0xf0,0x00,0xc1,0xd3,0x01,0x05,0xc2,0x07]
 457 0xf0,0x00,0xc1,0xd3,0x01,0x05,0xc2,0x07
 458
 459 # GFX940: v_mfma_f32_16x16x1_4b_f32 v[240:255], a1, v2, v[240:255] ; encoding: [0xf0,0x00,0xc1,0xd3,0x01,0x05,0xc2,0x0f]
 460 0xf0,0x00,0xc1,0xd3,0x01,0x05,0xc2,0x0f
 461
 462 # GFX940: v_mfma_f32_16x16x1_4b_f32 v[240:255], v1, a2, v[240:255] cbsz:2 abid:7 blgp:3 ; encoding: [0xf0,0x3a,0xc1,0xd3,0x01,0x05,0xc2,0x77]
 463 0xf0,0x3a,0xc1,0xd3,0x01,0x05,0xc2,0x77
 464
 465 # GFX940: v_mfma_f32_16x16x1_4b_f32 a[240:255], a255, a255, a[240:255] cbsz:7 abid:15 blgp:7 ; encoding: [0xf0,0xff,0xc1,0xd3,0xff,0xff,0xc3,0xff]
 466 0xf0,0xff,0xc1,0xd3,0xff,0xff,0xc3,0xff
 467
 468 # GFX940: v_mfma_f32_16x16x4_4b_f16 v[240:255], v[2:3], v[4:5], v[240:255] ; encoding: [0xf0,0x00,0xc9,0xd3,0x02,0x09,0xc2,0x07]
 469 0xf0,0x00,0xc9,0xd3,0x02,0x09,0xc2,0x07
 470
 471 # GFX940: v_mfma_f32_16x16x4_4b_f16 v[240:255], a[2:3], v[4:5], v[240:255] ; encoding: [0xf0,0x00,0xc9,0xd3,0x02,0x09,0xc2,0x0f]
 472 0xf0,0x00,0xc9,0xd3,0x02,0x09,0xc2,0x0f
 473
 474 # GFX940: v_mfma_f32_16x16x4_4b_f16 v[240:255], v[2:3], a[4:5], v[240:255] cbsz:2 abid:7 blgp:3 ; encoding: [0xf0,0x3a,0xc9,0xd3,0x02,0x09,0xc2,0x77]
 475 0xf0,0x3a,0xc9,0xd3,0x02,0x09,0xc2,0x77
 476
 477 # GFX940: v_mfma_f32_16x16x4_4b_f16 a[240:255], a[254:255], a[254:255], a[240:255] cbsz:7 abid:15 blgp:7 ; encoding: [0xf0,0xff,0xc9,0xd3,0xfe,0xfd,0xc3,0xff]
 478 0xf0,0xff,0xc9,0xd3,0xfe,0xfd,0xc3,0xff
 479
 480 # GFX940: v_mfma_f32_16x16x4_f32 v[10:13], v1, v2, v[6:9] ; encoding: [0x0a,0x00,0xc5,0xd3,0x01,0x05,0x1a,0x04]
 481 0x0a,0x00,0xc5,0xd3,0x01,0x05,0x1a,0x04
 482
 483 # GFX940: v_mfma_f32_16x16x4_f32 v[252:255], a255, v255, v[252:255] ; encoding: [0xfc,0x00,0xc5,0xd3,0xff,0xff,0xf3,0x0f]
 484 0xfc,0x00,0xc5,0xd3,0xff,0xff,0xf3,0x0f
 485
 486 # GFX940: v_mfma_f32_16x16x4_f32 a[10:13], v1, a2, a[6:9] cbsz:2 abid:7 blgp:3 ; encoding: [0x0a,0xba,0xc5,0xd3,0x01,0x05,0x1a,0x74]
 487 0x0a,0xba,0xc5,0xd3,0x01,0x05,0x1a,0x74
 488
 489 # GFX940: v_mfma_f32_16x16x4_f32 a[252:255], a255, a255, a[252:255] cbsz:7 abid:15 blgp:7 ; encoding: [0xfc,0xff,0xc5,0xd3,0xff,0xff,0xf3,0xff]
 490 0xfc,0xff,0xc5,0xd3,0xff,0xff,0xf3,0xff
 491
 492 # GFX940: v_mfma_f32_32x32x1_2b_f32 v[224:255], v1, v2, v[224:255] ; encoding: [0xe0,0x00,0xc0,0xd3,0x01,0x05,0x82,0x07]
 493 0xe0,0x00,0xc0,0xd3,0x01,0x05,0x82,0x07
 494
 495 # GFX940: v_mfma_f32_32x32x1_2b_f32 v[224:255], a1, v2, v[224:255] ; encoding: [0xe0,0x00,0xc0,0xd3,0x01,0x05,0x82,0x0f]
 496 0xe0,0x00,0xc0,0xd3,0x01,0x05,0x82,0x0f
 497
 498 # GFX940: v_mfma_f32_32x32x1_2b_f32 v[224:255], v1, a2, v[224:255] cbsz:2 abid:7 blgp:3 ; encoding: [0xe0,0x3a,0xc0,0xd3,0x01,0x05,0x82,0x77]
 499 0xe0,0x3a,0xc0,0xd3,0x01,0x05,0x82,0x77
 500
 501 # GFX940: v_mfma_f32_32x32x1_2b_f32 a[224:255], a255, a255, a[224:255] cbsz:7 abid:15 blgp:7 ; encoding: [0xe0,0xff,0xc0,0xd3,0xff,0xff,0x83,0xff]
 502 0xe0,0xff,0xc0,0xd3,0xff,0xff,0x83,0xff
 503
 504 # GFX940: v_mfma_f32_32x32x2_f32 v[240:255], v1, v2, v[240:255] ; encoding: [0xf0,0x00,0xc4,0xd3,0x01,0x05,0xc2,0x07]
 505 0xf0,0x00,0xc4,0xd3,0x01,0x05,0xc2,0x07
 506
 507 # GFX940: v_mfma_f32_32x32x2_f32 v[240:255], a1, v2, v[240:255] ; encoding: [0xf0,0x00,0xc4,0xd3,0x01,0x05,0xc2,0x0f]
 508 0xf0,0x00,0xc4,0xd3,0x01,0x05,0xc2,0x0f
 509
 510 # GFX940: v_mfma_f32_32x32x2_f32 v[240:255], v1, a2, v[240:255] cbsz:2 abid:7 blgp:3 ; encoding: [0xf0,0x3a,0xc4,0xd3,0x01,0x05,0xc2,0x77]
 511 0xf0,0x3a,0xc4,0xd3,0x01,0x05,0xc2,0x77
 512
 513 # GFX940: v_mfma_f32_32x32x2_f32 a[240:255], a255, a255, a[240:255] cbsz:7 abid:15 blgp:7 ; encoding: [0xf0,0xff,0xc4,0xd3,0xff,0xff,0xc3,0xff]
 514 0xf0,0xff,0xc4,0xd3,0xff,0xff,0xc3,0xff
 515
 516 # GFX940: v_mfma_f32_32x32x4_2b_f16 v[224:255], v[2:3], v[4:5], v[224:255] ; encoding: [0xe0,0x00,0xc8,0xd3,0x02,0x09,0x82,0x07]
 517 0xe0,0x00,0xc8,0xd3,0x02,0x09,0x82,0x07
 518
 519 # GFX940: v_mfma_f32_32x32x4_2b_f16 v[224:255], a[2:3], v[4:5], v[224:255] ; encoding: [0xe0,0x00,0xc8,0xd3,0x02,0x09,0x82,0x0f]
 520 0xe0,0x00,0xc8,0xd3,0x02,0x09,0x82,0x0f
 521
 522 # GFX940: v_mfma_f32_32x32x4_2b_f16 v[224:255], v[2:3], a[4:5], v[224:255] cbsz:2 abid:7 blgp:3 ; encoding: [0xe0,0x3a,0xc8,0xd3,0x02,0x09,0x82,0x77]
 523 0xe0,0x3a,0xc8,0xd3,0x02,0x09,0x82,0x77
 524
 525 # GFX940: v_mfma_f32_32x32x4_2b_f16 a[224:255], a[254:255], a[254:255], a[224:255] cbsz:7 abid:15 blgp:7 ; encoding: [0xe0,0xff,0xc8,0xd3,0xfe,0xfd,0x83,0xff]
 526 0xe0,0xff,0xc8,0xd3,0xfe,0xfd,0x83,0xff
 527
 528 # GFX940: v_mfma_f32_32x32x8_f16 v[240:255], v[2:3], v[4:5], v[240:255] ; encoding: [0xf0,0x00,0xcc,0xd3,0x02,0x09,0xc2,0x07]
 529 0xf0,0x00,0xcc,0xd3,0x02,0x09,0xc2,0x07
 530
 531 # GFX940: v_mfma_f32_32x32x8_f16 v[240:255], a[2:3], v[4:5], v[240:255] ; encoding: [0xf0,0x00,0xcc,0xd3,0x02,0x09,0xc2,0x0f]
 532 0xf0,0x00,0xcc,0xd3,0x02,0x09,0xc2,0x0f
 533
 534 # GFX940: v_mfma_f32_32x32x8_f16 v[240:255], v[2:3], a[4:5], v[240:255] cbsz:2 abid:7 blgp:3 ; encoding: [0xf0,0x3a,0xcc,0xd3,0x02,0x09,0xc2,0x77]
 535 0xf0,0x3a,0xcc,0xd3,0x02,0x09,0xc2,0x77
 536
 537 # GFX940: v_mfma_f32_32x32x8_f16 a[240:255], a[254:255], a[254:255], a[240:255] cbsz:7 abid:15 blgp:7 ; encoding: [0xf0,0xff,0xcc,0xd3,0xfe,0xfd,0xc3,0xff]
 538 0xf0,0xff,0xcc,0xd3,0xfe,0xfd,0xc3,0xff
 539
 540 # GFX940: v_mfma_f32_4x4x1_16b_f32 v[10:13], v1, v2, v[6:9] ; encoding: [0x0a,0x00,0xc2,0xd3,0x01,0x05,0x1a,0x04]
 541 0x0a,0x00,0xc2,0xd3,0x01,0x05,0x1a,0x04
 542
 543 # GFX940: v_mfma_f32_4x4x1_16b_f32 v[252:255], a255, v255, v[252:255] ; encoding: [0xfc,0x00,0xc2,0xd3,0xff,0xff,0xf3,0x0f]
 544 0xfc,0x00,0xc2,0xd3,0xff,0xff,0xf3,0x0f
 545
 546 # GFX940: v_mfma_f32_4x4x1_16b_f32 a[10:13], v1, a2, a[6:9] cbsz:2 abid:7 blgp:3 ; encoding: [0x0a,0xba,0xc2,0xd3,0x01,0x05,0x1a,0x74]
 547 0x0a,0xba,0xc2,0xd3,0x01,0x05,0x1a,0x74
 548
 549 # GFX940: v_mfma_f32_4x4x1_16b_f32 a[252:255], a255, a255, a[252:255] cbsz:7 abid:15 blgp:7 ; encoding: [0xfc,0xff,0xc2,0xd3,0xff,0xff,0xf3,0xff]
 550 0xfc,0xff,0xc2,0xd3,0xff,0xff,0xf3,0xff
 551
 552 # GFX940: v_mfma_f32_4x4x4_16b_f16 v[10:13], v[2:3], v[4:5], v[6:9] ; encoding: [0x0a,0x00,0xca,0xd3,0x02,0x09,0x1a,0x04]
 553 0x0a,0x00,0xca,0xd3,0x02,0x09,0x1a,0x04
 554
 555 # GFX940: v_mfma_f32_4x4x4_16b_f16 v[252:255], a[254:255], v[254:255], v[252:255] ; encoding: [0xfc,0x00,0xca,0xd3,0xfe,0xfd,0xf3,0x0f]
 556 0xfc,0x00,0xca,0xd3,0xfe,0xfd,0xf3,0x0f
 557
 558 # GFX940: v_mfma_f32_4x4x4_16b_f16 a[10:13], v[2:3], a[4:5], a[6:9] cbsz:2 abid:7 blgp:3 ; encoding: [0x0a,0xba,0xca,0xd3,0x02,0x09,0x1a,0x74]
 559 0x0a,0xba,0xca,0xd3,0x02,0x09,0x1a,0x74
 560
 561 # GFX940: v_mfma_f32_4x4x4_16b_f16 a[252:255], a[254:255], a[254:255], a[252:255] cbsz:7 abid:15 blgp:7 ; encoding: [0xfc,0xff,0xca,0xd3,0xfe,0xfd,0xf3,0xff]
 562 0xfc,0xff,0xca,0xd3,0xfe,0xfd,0xf3,0xff
 563
 564 # GFX940: v_mfma_i32_16x16x4_4b_i8 v[240:255], a1, a2, v[240:255] ; encoding: [0xf0,0x00,0xd1,0xd3,0x01,0x05,0xc2,0x1f]
 565 0xf0,0x00,0xd1,0xd3,0x01,0x05,0xc2,0x1f
 566
 567 # GFX940: v_mfma_i32_16x16x4_4b_i8 v[240:255], v1, a2, v[240:255] ; encoding: [0xf0,0x00,0xd1,0xd3,0x01,0x05,0xc2,0x17]
 568 0xf0,0x00,0xd1,0xd3,0x01,0x05,0xc2,0x17
 569
 570 # GFX940: v_mfma_i32_16x16x4_4b_i8 v[240:255], a1, v2, v[240:255] cbsz:2 abid:7 blgp:3 ; encoding: [0xf0,0x3a,0xd1,0xd3,0x01,0x05,0xc2,0x6f]
 571 0xf0,0x3a,0xd1,0xd3,0x01,0x05,0xc2,0x6f
 572
 573 # GFX940: v_mfma_i32_16x16x4_4b_i8 a[240:255], a255, a255, a[240:255] cbsz:7 abid:15 blgp:7 ; encoding: [0xf0,0xff,0xd1,0xd3,0xff,0xff,0xc3,0xff]
 574 0xf0,0xff,0xd1,0xd3,0xff,0xff,0xc3,0xff
 575
 576 # GFX940: v_mfma_i32_32x32x4_2b_i8 v[224:255], v1, v2, v[224:255] ; encoding: [0xe0,0x00,0xd0,0xd3,0x01,0x05,0x82,0x07]
 577 0xe0,0x00,0xd0,0xd3,0x01,0x05,0x82,0x07
 578
 579 # GFX940: v_mfma_i32_32x32x4_2b_i8 v[224:255], v1, a2, v[224:255] ; encoding: [0xe0,0x00,0xd0,0xd3,0x01,0x05,0x82,0x17]
 580 0xe0,0x00,0xd0,0xd3,0x01,0x05,0x82,0x17
 581
 582 # GFX940: v_mfma_i32_32x32x4_2b_i8 v[224:255], a1, v2, v[224:255] cbsz:2 abid:7 blgp:3 ; encoding: [0xe0,0x3a,0xd0,0xd3,0x01,0x05,0x82,0x6f]
 583 0xe0,0x3a,0xd0,0xd3,0x01,0x05,0x82,0x6f
 584
 585 # GFX940: v_mfma_i32_32x32x4_2b_i8 a[224:255], a255, a255, a[224:255] cbsz:7 abid:15 blgp:7 ; encoding: [0xe0,0xff,0xd0,0xd3,0xff,0xff,0x83,0xff]
 586 0xe0,0xff,0xd0,0xd3,0xff,0xff,0x83,0xff
 587
 588 # GFX940: v_mfma_i32_4x4x4_16b_i8 v[10:13], v1, v2, v[6:9] ; encoding: [0x0a,0x00,0xd2,0xd3,0x01,0x05,0x1a,0x04]
 589 0x0a,0x00,0xd2,0xd3,0x01,0x05,0x1a,0x04
 590
 591 # GFX940: v_mfma_i32_4x4x4_16b_i8 v[252:255], a255, v255, v[252:255] ; encoding: [0xfc,0x00,0xd2,0xd3,0xff,0xff,0xf3,0x0f]
 592 0xfc,0x00,0xd2,0xd3,0xff,0xff,0xf3,0x0f
 593
 594 # GFX940: v_mfma_i32_4x4x4_16b_i8 a[10:13], v1, a2, a[6:9] cbsz:2 abid:7 blgp:3 ; encoding: [0x0a,0xba,0xd2,0xd3,0x01,0x05,0x1a,0x74]
 595 0x0a,0xba,0xd2,0xd3,0x01,0x05,0x1a,0x74
 596
 597 # GFX940: v_mfma_i32_4x4x4_16b_i8 a[252:255], a255, a255, a[252:255] cbsz:7 abid:15 blgp:7 ; encoding: [0xfc,0xff,0xd2,0xd3,0xff,0xff,0xf3,0xff]
 598 0xfc,0xff,0xd2,0xd3,0xff,0xff,0xf3,0xff