test/MC/AMDGPU/mai.s

   1 // RUN: llvm-mc -arch=amdgcn -mcpu=gfx908 -show-encoding %s | FileCheck -check-prefix=GFX908 %s
   2
   3 v_accvgpr_read_b32 v2, a0
   4 // GFX908: v_accvgpr_read_b32 v2, a0       ; encoding: [0x02,0x00,0xd8,0xd3,0x00,0x01,0x00,0x08]
   5
   6 v_accvgpr_read_b32 v2, a1
   7 // GFX908: v_accvgpr_read_b32 v2, a1       ; encoding: [0x02,0x00,0xd8,0xd3,0x01,0x01,0x00,0x08]
   8
   9 v_accvgpr_read_b32 v2, a255
  10 // GFX908: v_accvgpr_read_b32 v2, a255     ; encoding: [0x02,0x00,0xd8,0xd3,0xff,0x01,0x00,0x08]
  11
  12 v_accvgpr_read v2, a10
  13 // GFX908: v_accvgpr_read_b32 v2, a10      ; encoding: [0x02,0x00,0xd8,0xd3,0x0a,0x01,0x00,0x08]
  14
  15 v_accvgpr_write_b32 a2, -2.0
  16 // GFX908: v_accvgpr_write_b32 a2, -2.0    ; encoding: [0x02,0x00,0xd9,0xd3,0xf5,0x00,0x00,0x00]
  17
  18 v_accvgpr_write_b32 a2, -2
  19 // GFX908: v_accvgpr_write_b32 a2, -2      ; encoding: [0x02,0x00,0xd9,0xd3,0xc2,0x00,0x00,0x00]
  20
  21 v_accvgpr_write_b32 a2, v1
  22 // GFX908: v_accvgpr_write_b32 a2, v1      ; encoding: [0x02,0x00,0xd9,0xd3,0x01,0x01,0x00,0x00]
  23
  24 v_accvgpr_write a2, v255
  25 // GFX908: v_accvgpr_write_b32 a2, v255    ; encoding: [0x02,0x00,0xd9,0xd3,0xff,0x01,0x00,0x00]
  26
  27 v_mfma_f32_32x32x1f32 a[0:31], v0, v1, a[1:32]
  28 // GFX908: v_mfma_f32_32x32x1f32 a[0:31], v0, v1, a[1:32] ; encoding: [0x00,0x00,0xc0,0xd3,0x00,0x03,0x06,0x04]
  29
  30 v_mfma_f32_32x32x1f32 a[0:31], v0, v1, a[1:32] cbsz:3 abid:2 blgp:7
  31 // GFX908: v_mfma_f32_32x32x1f32 a[0:31], v0, v1, a[1:32] cbsz:3 abid:2 blgp:7 ; encoding: [0x00,0x13,0xc0,0xd3,0x00,0x03,0x06,0xe4]
  32
  33 v_mfma_f32_32x32x1f32 a[0:31], v0, a1, a[1:32]
  34 // GFX908: v_mfma_f32_32x32x1f32 a[0:31], v0, a1, a[1:32] ; encoding: [0x00,0x00,0xc0,0xd3,0x00,0x03,0x06,0x14]
  35
  36 v_mfma_f32_32x32x1f32 a[0:31], v0, a1, a[1:32] cbsz:3 abid:2 blgp:7
  37 // GFX908: v_mfma_f32_32x32x1f32 a[0:31], v0, a1, a[1:32] cbsz:3 abid:2 blgp:7 ; encoding: [0x00,0x13,0xc0,0xd3,0x00,0x03,0x06,0xf4]
  38
  39 v_mfma_f32_32x32x1f32 a[0:31], a0, v1, a[1:32]
  40 // GFX908: v_mfma_f32_32x32x1f32 a[0:31], a0, v1, a[1:32] ; encoding: [0x00,0x00,0xc0,0xd3,0x00,0x03,0x06,0x0c]
  41
  42 v_mfma_f32_32x32x1f32 a[0:31], a0, v1, a[1:32] cbsz:3 abid:2 blgp:7
  43 // GFX908: v_mfma_f32_32x32x1f32 a[0:31], a0, v1, a[1:32] cbsz:3 abid:2 blgp:7 ; encoding: [0x00,0x13,0xc0,0xd3,0x00,0x03,0x06,0xec]
  44
  45 v_mfma_f32_32x32x1f32 a[0:31], a0, a1, a[1:32]
  46 // GFX908: v_mfma_f32_32x32x1f32 a[0:31], a0, a1, a[1:32] ; encoding: [0x00,0x00,0xc0,0xd3,0x00,0x03,0x06,0x1c]
  47
  48 v_mfma_f32_32x32x1f32 a[0:31], a0, a1, a[1:32] cbsz:3 abid:2 blgp:7
  49 // GFX908: v_mfma_f32_32x32x1f32 a[0:31], a0, a1, a[1:32] cbsz:3 abid:2 blgp:7 ; encoding: [0x00,0x13,0xc0,0xd3,0x00,0x03,0x06,0xfc]
  50
  51 v_mfma_f32_16x16x1f32 a[0:15], v0, v1, a[1:16]
  52 // GFX908: v_mfma_f32_16x16x1f32 a[0:15], v0, v1, a[1:16] ; encoding: [0x00,0x00,0xc1,0xd3,0x00,0x03,0x06,0x04]
  53
  54 v_mfma_f32_16x16x1f32 a[0:15], v0, v1, a[1:16] cbsz:3 abid:2 blgp:7
  55 // GFX908: v_mfma_f32_16x16x1f32 a[0:15], v0, v1, a[1:16] cbsz:3 abid:2 blgp:7 ; encoding: [0x00,0x13,0xc1,0xd3,0x00,0x03,0x06,0xe4]
  56
  57 v_mfma_f32_16x16x1f32 a[0:15], v0, a1, a[1:16]
  58 // GFX908: v_mfma_f32_16x16x1f32 a[0:15], v0, a1, a[1:16] ; encoding: [0x00,0x00,0xc1,0xd3,0x00,0x03,0x06,0x14]
  59
  60 v_mfma_f32_16x16x1f32 a[0:15], v0, a1, a[1:16] cbsz:3 abid:2 blgp:7
  61 // GFX908: v_mfma_f32_16x16x1f32 a[0:15], v0, a1, a[1:16] cbsz:3 abid:2 blgp:7 ; encoding: [0x00,0x13,0xc1,0xd3,0x00,0x03,0x06,0xf4]
  62
  63 v_mfma_f32_16x16x1f32 a[0:15], a0, v1, a[1:16]
  64 // GFX908: v_mfma_f32_16x16x1f32 a[0:15], a0, v1, a[1:16] ; encoding: [0x00,0x00,0xc1,0xd3,0x00,0x03,0x06,0x0c]
  65
  66 v_mfma_f32_16x16x1f32 a[0:15], a0, v1, a[1:16] cbsz:3 abid:2 blgp:7
  67 // GFX908: v_mfma_f32_16x16x1f32 a[0:15], a0, v1, a[1:16] cbsz:3 abid:2 blgp:7 ; encoding: [0x00,0x13,0xc1,0xd3,0x00,0x03,0x06,0xec]
  68
  69 v_mfma_f32_16x16x1f32 a[0:15], a0, a1, a[1:16]
  70 // GFX908: v_mfma_f32_16x16x1f32 a[0:15], a0, a1, a[1:16] ; encoding: [0x00,0x00,0xc1,0xd3,0x00,0x03,0x06,0x1c]
  71
  72 v_mfma_f32_16x16x1f32 a[0:15], a0, a1, a[1:16] cbsz:3 abid:2 blgp:7
  73 // GFX908: v_mfma_f32_16x16x1f32 a[0:15], a0, a1, a[1:16] cbsz:3 abid:2 blgp:7 ; encoding: [0x00,0x13,0xc1,0xd3,0x00,0x03,0x06,0xfc]
  74
  75 v_mfma_f32_4x4x1f32 a[0:3], v0, v1, a[1:4]
  76 // GFX908: v_mfma_f32_4x4x1f32 a[0:3], v0, v1, a[1:4] ; encoding: [0x00,0x00,0xc2,0xd3,0x00,0x03,0x06,0x04]
  77
  78 v_mfma_f32_4x4x1f32 a[0:3], v0, v1, a[1:4] cbsz:3 abid:2 blgp:7
  79 // GFX908: v_mfma_f32_4x4x1f32 a[0:3], v0, v1, a[1:4] cbsz:3 abid:2 blgp:7 ; encoding: [0x00,0x13,0xc2,0xd3,0x00,0x03,0x06,0xe4]
  80
  81 v_mfma_f32_4x4x1f32 a[0:3], v0, a1, a[1:4]
  82 // GFX908: v_mfma_f32_4x4x1f32 a[0:3], v0, a1, a[1:4] ; encoding: [0x00,0x00,0xc2,0xd3,0x00,0x03,0x06,0x14]
  83
  84 v_mfma_f32_4x4x1f32 a[0:3], v0, a1, a[1:4] cbsz:3 abid:2 blgp:7
  85 // GFX908: v_mfma_f32_4x4x1f32 a[0:3], v0, a1, a[1:4] cbsz:3 abid:2 blgp:7 ; encoding: [0x00,0x13,0xc2,0xd3,0x00,0x03,0x06,0xf4]
  86
  87 v_mfma_f32_4x4x1f32 a[0:3], a0, v1, a[1:4]
  88 // GFX908: v_mfma_f32_4x4x1f32 a[0:3], a0, v1, a[1:4] ; encoding: [0x00,0x00,0xc2,0xd3,0x00,0x03,0x06,0x0c]
  89
  90 v_mfma_f32_4x4x1f32 a[0:3], a0, v1, a[1:4] cbsz:3 abid:2 blgp:7
  91 // GFX908: v_mfma_f32_4x4x1f32 a[0:3], a0, v1, a[1:4] cbsz:3 abid:2 blgp:7 ; encoding: [0x00,0x13,0xc2,0xd3,0x00,0x03,0x06,0xec]
  92
  93 v_mfma_f32_4x4x1f32 a[0:3], a0, a1, a[1:4]
  94 // GFX908: v_mfma_f32_4x4x1f32 a[0:3], a0, a1, a[1:4] ; encoding: [0x00,0x00,0xc2,0xd3,0x00,0x03,0x06,0x1c]
  95
  96 v_mfma_f32_4x4x1f32 a[0:3], a0, a1, a[1:4] cbsz:3 abid:2 blgp:7
  97 // GFX908: v_mfma_f32_4x4x1f32 a[0:3], a0, a1, a[1:4] cbsz:3 abid:2 blgp:7 ; encoding: [0x00,0x13,0xc2,0xd3,0x00,0x03,0x06,0xfc]
  98
  99 v_mfma_f32_32x32x2f32 a[0:15], v0, v1, a[1:16]
 100 // GFX908: v_mfma_f32_32x32x2f32 a[0:15], v0, v1, a[1:16] ; encoding: [0x00,0x00,0xc4,0xd3,0x00,0x03,0x06,0x04]
 101
 102 v_mfma_f32_32x32x2f32 a[0:15], v0, v1, a[1:16] cbsz:3 abid:2 blgp:7
 103 // GFX908: v_mfma_f32_32x32x2f32 a[0:15], v0, v1, a[1:16] cbsz:3 abid:2 blgp:7 ; encoding: [0x00,0x13,0xc4,0xd3,0x00,0x03,0x06,0xe4]
 104
 105 v_mfma_f32_32x32x2f32 a[0:15], v0, a1, a[1:16]
 106 // GFX908: v_mfma_f32_32x32x2f32 a[0:15], v0, a1, a[1:16] ; encoding: [0x00,0x00,0xc4,0xd3,0x00,0x03,0x06,0x14]
 107
 108 v_mfma_f32_32x32x2f32 a[0:15], v0, a1, a[1:16] cbsz:3 abid:2 blgp:7
 109 // GFX908: v_mfma_f32_32x32x2f32 a[0:15], v0, a1, a[1:16] cbsz:3 abid:2 blgp:7 ; encoding: [0x00,0x13,0xc4,0xd3,0x00,0x03,0x06,0xf4]
 110
 111 v_mfma_f32_32x32x2f32 a[0:15], a0, v1, a[1:16]
 112 // GFX908: v_mfma_f32_32x32x2f32 a[0:15], a0, v1, a[1:16] ; encoding: [0x00,0x00,0xc4,0xd3,0x00,0x03,0x06,0x0c]
 113
 114 v_mfma_f32_32x32x2f32 a[0:15], a0, v1, a[1:16] cbsz:3 abid:2 blgp:7
 115 // GFX908: v_mfma_f32_32x32x2f32 a[0:15], a0, v1, a[1:16] cbsz:3 abid:2 blgp:7 ; encoding: [0x00,0x13,0xc4,0xd3,0x00,0x03,0x06,0xec]
 116
 117 v_mfma_f32_32x32x2f32 a[0:15], a0, a1, a[1:16]
 118 // GFX908: v_mfma_f32_32x32x2f32 a[0:15], a0, a1, a[1:16] ; encoding: [0x00,0x00,0xc4,0xd3,0x00,0x03,0x06,0x1c]
 119
 120 v_mfma_f32_32x32x2f32 a[0:15], a0, a1, a[1:16] cbsz:3 abid:2 blgp:7
 121 // GFX908: v_mfma_f32_32x32x2f32 a[0:15], a0, a1, a[1:16] cbsz:3 abid:2 blgp:7 ; encoding: [0x00,0x13,0xc4,0xd3,0x00,0x03,0x06,0xfc]
 122
 123 v_mfma_f32_16x16x4f32 a[0:3], v0, v1, a[1:4]
 124 // GFX908: v_mfma_f32_16x16x4f32 a[0:3], v0, v1, a[1:4] ; encoding: [0x00,0x00,0xc5,0xd3,0x00,0x03,0x06,0x04]
 125
 126 v_mfma_f32_16x16x4f32 a[0:3], v0, v1, a[1:4] cbsz:3 abid:2 blgp:7
 127 // GFX908: v_mfma_f32_16x16x4f32 a[0:3], v0, v1, a[1:4] cbsz:3 abid:2 blgp:7 ; encoding: [0x00,0x13,0xc5,0xd3,0x00,0x03,0x06,0xe4]
 128
 129 v_mfma_f32_16x16x4f32 a[0:3], v0, a1, a[1:4]
 130 // GFX908: v_mfma_f32_16x16x4f32 a[0:3], v0, a1, a[1:4] ; encoding: [0x00,0x00,0xc5,0xd3,0x00,0x03,0x06,0x14]
 131
 132 v_mfma_f32_16x16x4f32 a[0:3], v0, a1, a[1:4] cbsz:3 abid:2 blgp:7
 133 // GFX908: v_mfma_f32_16x16x4f32 a[0:3], v0, a1, a[1:4] cbsz:3 abid:2 blgp:7 ; encoding: [0x00,0x13,0xc5,0xd3,0x00,0x03,0x06,0xf4]
 134
 135 v_mfma_f32_16x16x4f32 a[0:3], a0, v1, a[1:4]
 136 // GFX908: v_mfma_f32_16x16x4f32 a[0:3], a0, v1, a[1:4] ; encoding: [0x00,0x00,0xc5,0xd3,0x00,0x03,0x06,0x0c]
 137
 138 v_mfma_f32_16x16x4f32 a[0:3], a0, v1, a[1:4] cbsz:3 abid:2 blgp:7
 139 // GFX908: v_mfma_f32_16x16x4f32 a[0:3], a0, v1, a[1:4] cbsz:3 abid:2 blgp:7 ; encoding: [0x00,0x13,0xc5,0xd3,0x00,0x03,0x06,0xec]
 140
 141 v_mfma_f32_16x16x4f32 a[0:3], a0, a1, a[1:4]
 142 // GFX908: v_mfma_f32_16x16x4f32 a[0:3], a0, a1, a[1:4] ; encoding: [0x00,0x00,0xc5,0xd3,0x00,0x03,0x06,0x1c]
 143
 144 v_mfma_f32_16x16x4f32 a[0:3], a0, a1, a[1:4] cbsz:3 abid:2 blgp:7
 145 // GFX908: v_mfma_f32_16x16x4f32 a[0:3], a0, a1, a[1:4] cbsz:3 abid:2 blgp:7 ; encoding: [0x00,0x13,0xc5,0xd3,0x00,0x03,0x06,0xfc]
 146
 147 v_mfma_f32_32x32x4f16 a[0:31], v[0:1], v[1:2], a[1:32]
 148 // GFX908: v_mfma_f32_32x32x4f16 a[0:31], v[0:1], v[1:2], a[1:32] ; encoding: [0x00,0x00,0xc8,0xd3,0x00,0x03,0x06,0x04]
 149
 150 v_mfma_f32_32x32x4f16 a[0:31], v[0:1], v[1:2], a[1:32] cbsz:3 abid:2 blgp:7
 151 // GFX908: v_mfma_f32_32x32x4f16 a[0:31], v[0:1], v[1:2], a[1:32] cbsz:3 abid:2 blgp:7 ; encoding: [0x00,0x13,0xc8,0xd3,0x00,0x03,0x06,0xe4]
 152
 153 v_mfma_f32_32x32x4f16 a[0:31], v[0:1], a[1:2], a[1:32]
 154 // GFX908: v_mfma_f32_32x32x4f16 a[0:31], v[0:1], a[1:2], a[1:32] ; encoding: [0x00,0x00,0xc8,0xd3,0x00,0x03,0x06,0x14]
 155
 156 v_mfma_f32_32x32x4f16 a[0:31], v[0:1], a[1:2], a[1:32] cbsz:3 abid:2 blgp:7
 157 // GFX908: v_mfma_f32_32x32x4f16 a[0:31], v[0:1], a[1:2], a[1:32] cbsz:3 abid:2 blgp:7 ; encoding: [0x00,0x13,0xc8,0xd3,0x00,0x03,0x06,0xf4]
 158
 159 v_mfma_f32_32x32x4f16 a[0:31], a[0:1], v[1:2], a[1:32]
 160 // GFX908: v_mfma_f32_32x32x4f16 a[0:31], a[0:1], v[1:2], a[1:32] ; encoding: [0x00,0x00,0xc8,0xd3,0x00,0x03,0x06,0x0c]
 161
 162 v_mfma_f32_32x32x4f16 a[0:31], a[0:1], v[1:2], a[1:32] cbsz:3 abid:2 blgp:7
 163 // GFX908: v_mfma_f32_32x32x4f16 a[0:31], a[0:1], v[1:2], a[1:32] cbsz:3 abid:2 blgp:7 ; encoding: [0x00,0x13,0xc8,0xd3,0x00,0x03,0x06,0xec]
 164
 165 v_mfma_f32_32x32x4f16 a[0:31], a[0:1], a[1:2], a[1:32]
 166 // GFX908: v_mfma_f32_32x32x4f16 a[0:31], a[0:1], a[1:2], a[1:32] ; encoding: [0x00,0x00,0xc8,0xd3,0x00,0x03,0x06,0x1c]
 167
 168 v_mfma_f32_32x32x4f16 a[0:31], a[0:1], a[1:2], a[1:32] cbsz:3 abid:2 blgp:7
 169 // GFX908: v_mfma_f32_32x32x4f16 a[0:31], a[0:1], a[1:2], a[1:32] cbsz:3 abid:2 blgp:7 ; encoding: [0x00,0x13,0xc8,0xd3,0x00,0x03,0x06,0xfc]
 170
 171 v_mfma_f32_16x16x4f16 a[0:15], v[0:1], v[1:2], a[1:16]
 172 // GFX908: v_mfma_f32_16x16x4f16 a[0:15], v[0:1], v[1:2], a[1:16] ; encoding: [0x00,0x00,0xc9,0xd3,0x00,0x03,0x06,0x04]
 173
 174 v_mfma_f32_16x16x4f16 a[0:15], v[0:1], v[1:2], a[1:16] cbsz:3 abid:2 blgp:7
 175 // GFX908: v_mfma_f32_16x16x4f16 a[0:15], v[0:1], v[1:2], a[1:16] cbsz:3 abid:2 blgp:7 ; encoding: [0x00,0x13,0xc9,0xd3,0x00,0x03,0x06,0xe4]
 176
 177 v_mfma_f32_16x16x4f16 a[0:15], v[0:1], a[1:2], a[1:16]
 178 // GFX908: v_mfma_f32_16x16x4f16 a[0:15], v[0:1], a[1:2], a[1:16] ; encoding: [0x00,0x00,0xc9,0xd3,0x00,0x03,0x06,0x14]
 179
 180 v_mfma_f32_16x16x4f16 a[0:15], v[0:1], a[1:2], a[1:16] cbsz:3 abid:2 blgp:7
 181 // GFX908: v_mfma_f32_16x16x4f16 a[0:15], v[0:1], a[1:2], a[1:16] cbsz:3 abid:2 blgp:7 ; encoding: [0x00,0x13,0xc9,0xd3,0x00,0x03,0x06,0xf4]
 182
 183 v_mfma_f32_16x16x4f16 a[0:15], a[0:1], v[1:2], a[1:16]
 184 // GFX908: v_mfma_f32_16x16x4f16 a[0:15], a[0:1], v[1:2], a[1:16] ; encoding: [0x00,0x00,0xc9,0xd3,0x00,0x03,0x06,0x0c]
 185
 186 v_mfma_f32_16x16x4f16 a[0:15], a[0:1], v[1:2], a[1:16] cbsz:3 abid:2 blgp:7
 187 // GFX908: v_mfma_f32_16x16x4f16 a[0:15], a[0:1], v[1:2], a[1:16] cbsz:3 abid:2 blgp:7 ; encoding: [0x00,0x13,0xc9,0xd3,0x00,0x03,0x06,0xec]
 188
 189 v_mfma_f32_16x16x4f16 a[0:15], a[0:1], a[1:2], a[1:16]
 190 // GFX908: v_mfma_f32_16x16x4f16 a[0:15], a[0:1], a[1:2], a[1:16] ; encoding: [0x00,0x00,0xc9,0xd3,0x00,0x03,0x06,0x1c]
 191
 192 v_mfma_f32_16x16x4f16 a[0:15], a[0:1], a[1:2], a[1:16] cbsz:3 abid:2 blgp:7
 193 // GFX908: v_mfma_f32_16x16x4f16 a[0:15], a[0:1], a[1:2], a[1:16] cbsz:3 abid:2 blgp:7 ; encoding: [0x00,0x13,0xc9,0xd3,0x00,0x03,0x06,0xfc]
 194
 195 v_mfma_f32_4x4x4f16 a[0:3], v[0:1], v[1:2], a[1:4]
 196 // GFX908: v_mfma_f32_4x4x4f16 a[0:3], v[0:1], v[1:2], a[1:4] ; encoding: [0x00,0x00,0xca,0xd3,0x00,0x03,0x06,0x04]
 197
 198 v_mfma_f32_4x4x4f16 a[0:3], v[0:1], v[1:2], a[1:4] cbsz:3 abid:2 blgp:7
 199 // GFX908: v_mfma_f32_4x4x4f16 a[0:3], v[0:1], v[1:2], a[1:4] cbsz:3 abid:2 blgp:7 ; encoding: [0x00,0x13,0xca,0xd3,0x00,0x03,0x06,0xe4]
 200
 201 v_mfma_f32_4x4x4f16 a[0:3], v[0:1], a[1:2], a[1:4]
 202 // GFX908: v_mfma_f32_4x4x4f16 a[0:3], v[0:1], a[1:2], a[1:4] ; encoding: [0x00,0x00,0xca,0xd3,0x00,0x03,0x06,0x14]
 203
 204 v_mfma_f32_4x4x4f16 a[0:3], v[0:1], a[1:2], a[1:4] cbsz:3 abid:2 blgp:7
 205 // GFX908: v_mfma_f32_4x4x4f16 a[0:3], v[0:1], a[1:2], a[1:4] cbsz:3 abid:2 blgp:7 ; encoding: [0x00,0x13,0xca,0xd3,0x00,0x03,0x06,0xf4]
 206
 207 v_mfma_f32_4x4x4f16 a[0:3], a[0:1], v[1:2], a[1:4]
 208 // GFX908: v_mfma_f32_4x4x4f16 a[0:3], a[0:1], v[1:2], a[1:4] ; encoding: [0x00,0x00,0xca,0xd3,0x00,0x03,0x06,0x0c]
 209
 210 v_mfma_f32_4x4x4f16 a[0:3], a[0:1], v[1:2], a[1:4] cbsz:3 abid:2 blgp:7
 211 // GFX908: v_mfma_f32_4x4x4f16 a[0:3], a[0:1], v[1:2], a[1:4] cbsz:3 abid:2 blgp:7 ; encoding: [0x00,0x13,0xca,0xd3,0x00,0x03,0x06,0xec]
 212
 213 v_mfma_f32_4x4x4f16 a[0:3], a[0:1], a[1:2], a[1:4]
 214 // GFX908: v_mfma_f32_4x4x4f16 a[0:3], a[0:1], a[1:2], a[1:4] ; encoding: [0x00,0x00,0xca,0xd3,0x00,0x03,0x06,0x1c]
 215
 216 v_mfma_f32_4x4x4f16 a[0:3], a[0:1], a[1:2], a[1:4] cbsz:3 abid:2 blgp:7
 217 // GFX908: v_mfma_f32_4x4x4f16 a[0:3], a[0:1], a[1:2], a[1:4] cbsz:3 abid:2 blgp:7 ; encoding: [0x00,0x13,0xca,0xd3,0x00,0x03,0x06,0xfc]
 218
 219 v_mfma_f32_32x32x8f16 a[0:15], v[0:1], v[1:2], a[1:16]
 220 // GFX908: v_mfma_f32_32x32x8f16 a[0:15], v[0:1], v[1:2], a[1:16] ; encoding: [0x00,0x00,0xcc,0xd3,0x00,0x03,0x06,0x04]
 221
 222 v_mfma_f32_32x32x8f16 a[0:15], v[0:1], v[1:2], a[1:16] cbsz:3 abid:2 blgp:7
 223 // GFX908: v_mfma_f32_32x32x8f16 a[0:15], v[0:1], v[1:2], a[1:16] cbsz:3 abid:2 blgp:7 ; encoding: [0x00,0x13,0xcc,0xd3,0x00,0x03,0x06,0xe4]
 224
 225 v_mfma_f32_32x32x8f16 a[0:15], v[0:1], a[1:2], a[1:16]
 226 // GFX908: v_mfma_f32_32x32x8f16 a[0:15], v[0:1], a[1:2], a[1:16] ; encoding: [0x00,0x00,0xcc,0xd3,0x00,0x03,0x06,0x14]
 227
 228 v_mfma_f32_32x32x8f16 a[0:15], v[0:1], a[1:2], a[1:16] cbsz:3 abid:2 blgp:7
 229 // GFX908: v_mfma_f32_32x32x8f16 a[0:15], v[0:1], a[1:2], a[1:16] cbsz:3 abid:2 blgp:7 ; encoding: [0x00,0x13,0xcc,0xd3,0x00,0x03,0x06,0xf4]
 230
 231 v_mfma_f32_32x32x8f16 a[0:15], a[0:1], v[1:2], a[1:16]
 232 // GFX908: v_mfma_f32_32x32x8f16 a[0:15], a[0:1], v[1:2], a[1:16] ; encoding: [0x00,0x00,0xcc,0xd3,0x00,0x03,0x06,0x0c]
 233
 234 v_mfma_f32_32x32x8f16 a[0:15], a[0:1], v[1:2], a[1:16] cbsz:3 abid:2 blgp:7
 235 // GFX908: v_mfma_f32_32x32x8f16 a[0:15], a[0:1], v[1:2], a[1:16] cbsz:3 abid:2 blgp:7 ; encoding: [0x00,0x13,0xcc,0xd3,0x00,0x03,0x06,0xec]
 236
 237 v_mfma_f32_32x32x8f16 a[0:15], a[0:1], a[1:2], a[1:16]
 238 // GFX908: v_mfma_f32_32x32x8f16 a[0:15], a[0:1], a[1:2], a[1:16] ; encoding: [0x00,0x00,0xcc,0xd3,0x00,0x03,0x06,0x1c]
 239
 240 v_mfma_f32_32x32x8f16 a[0:15], a[0:1], a[1:2], a[1:16] cbsz:3 abid:2 blgp:7
 241 // GFX908: v_mfma_f32_32x32x8f16 a[0:15], a[0:1], a[1:2], a[1:16] cbsz:3 abid:2 blgp:7 ; encoding: [0x00,0x13,0xcc,0xd3,0x00,0x03,0x06,0xfc]
 242
 243 v_mfma_f32_16x16x16f16 a[0:3], v[0:1], v[1:2], a[1:4]
 244 // GFX908: v_mfma_f32_16x16x16f16 a[0:3], v[0:1], v[1:2], a[1:4] ; encoding: [0x00,0x00,0xcd,0xd3,0x00,0x03,0x06,0x04]
 245
 246 v_mfma_f32_16x16x16f16 a[0:3], v[0:1], v[1:2], a[1:4] cbsz:3 abid:2 blgp:7
 247 // GFX908: v_mfma_f32_16x16x16f16 a[0:3], v[0:1], v[1:2], a[1:4] cbsz:3 abid:2 blgp:7 ; encoding: [0x00,0x13,0xcd,0xd3,0x00,0x03,0x06,0xe4]
 248
 249 v_mfma_f32_16x16x16f16 a[0:3], v[0:1], a[1:2], a[1:4]
 250 // GFX908: v_mfma_f32_16x16x16f16 a[0:3], v[0:1], a[1:2], a[1:4] ; encoding: [0x00,0x00,0xcd,0xd3,0x00,0x03,0x06,0x14]
 251
 252 v_mfma_f32_16x16x16f16 a[0:3], v[0:1], a[1:2], a[1:4] cbsz:3 abid:2 blgp:7
 253 // GFX908: v_mfma_f32_16x16x16f16 a[0:3], v[0:1], a[1:2], a[1:4] cbsz:3 abid:2 blgp:7 ; encoding: [0x00,0x13,0xcd,0xd3,0x00,0x03,0x06,0xf4]
 254
 255 v_mfma_f32_16x16x16f16 a[0:3], a[0:1], v[1:2], a[1:4]
 256 // GFX908: v_mfma_f32_16x16x16f16 a[0:3], a[0:1], v[1:2], a[1:4] ; encoding: [0x00,0x00,0xcd,0xd3,0x00,0x03,0x06,0x0c]
 257
 258 v_mfma_f32_16x16x16f16 a[0:3], a[0:1], v[1:2], a[1:4] cbsz:3 abid:2 blgp:7
 259 // GFX908: v_mfma_f32_16x16x16f16 a[0:3], a[0:1], v[1:2], a[1:4] cbsz:3 abid:2 blgp:7 ; encoding: [0x00,0x13,0xcd,0xd3,0x00,0x03,0x06,0xec]
 260
 261 v_mfma_f32_16x16x16f16 a[0:3], a[0:1], a[1:2], a[1:4]
 262 // GFX908: v_mfma_f32_16x16x16f16 a[0:3], a[0:1], a[1:2], a[1:4] ; encoding: [0x00,0x00,0xcd,0xd3,0x00,0x03,0x06,0x1c]
 263
 264 v_mfma_f32_16x16x16f16 a[0:3], a[0:1], a[1:2], a[1:4] cbsz:3 abid:2 blgp:7
 265 // GFX908: v_mfma_f32_16x16x16f16 a[0:3], a[0:1], a[1:2], a[1:4] cbsz:3 abid:2 blgp:7 ; encoding: [0x00,0x13,0xcd,0xd3,0x00,0x03,0x06,0xfc]
 266
 267 v_mfma_i32_32x32x4i8 a[0:31], v0, v1, a[1:32]
 268 // GFX908: v_mfma_i32_32x32x4i8 a[0:31], v0, v1, a[1:32] ; encoding: [0x00,0x00,0xd0,0xd3,0x00,0x03,0x06,0x04]
 269
 270 v_mfma_i32_32x32x4i8 a[0:31], v0, v1, a[1:32] cbsz:3 abid:2 blgp:7
 271 // GFX908: v_mfma_i32_32x32x4i8 a[0:31], v0, v1, a[1:32] cbsz:3 abid:2 blgp:7 ; encoding: [0x00,0x13,0xd0,0xd3,0x00,0x03,0x06,0xe4]
 272
 273 v_mfma_i32_32x32x4i8 a[0:31], v0, a1, a[1:32]
 274 // GFX908: v_mfma_i32_32x32x4i8 a[0:31], v0, a1, a[1:32] ; encoding: [0x00,0x00,0xd0,0xd3,0x00,0x03,0x06,0x14]
 275
 276 v_mfma_i32_32x32x4i8 a[0:31], v0, a1, a[1:32] cbsz:3 abid:2 blgp:7
 277 // GFX908: v_mfma_i32_32x32x4i8 a[0:31], v0, a1, a[1:32] cbsz:3 abid:2 blgp:7 ; encoding: [0x00,0x13,0xd0,0xd3,0x00,0x03,0x06,0xf4]
 278
 279 v_mfma_i32_32x32x4i8 a[0:31], a0, v1, a[1:32]
 280 // GFX908: v_mfma_i32_32x32x4i8 a[0:31], a0, v1, a[1:32] ; encoding: [0x00,0x00,0xd0,0xd3,0x00,0x03,0x06,0x0c]
 281
 282 v_mfma_i32_32x32x4i8 a[0:31], a0, v1, a[1:32] cbsz:3 abid:2 blgp:7
 283 // GFX908: v_mfma_i32_32x32x4i8 a[0:31], a0, v1, a[1:32] cbsz:3 abid:2 blgp:7 ; encoding: [0x00,0x13,0xd0,0xd3,0x00,0x03,0x06,0xec]
 284
 285 v_mfma_i32_32x32x4i8 a[0:31], a0, a1, a[1:32]
 286 // GFX908: v_mfma_i32_32x32x4i8 a[0:31], a0, a1, a[1:32] ; encoding: [0x00,0x00,0xd0,0xd3,0x00,0x03,0x06,0x1c]
 287
 288 v_mfma_i32_32x32x4i8 a[0:31], a0, a1, a[1:32] cbsz:3 abid:2 blgp:7
 289 // GFX908: v_mfma_i32_32x32x4i8 a[0:31], a0, a1, a[1:32] cbsz:3 abid:2 blgp:7 ; encoding: [0x00,0x13,0xd0,0xd3,0x00,0x03,0x06,0xfc]
 290
 291 v_mfma_i32_16x16x4i8 a[0:15], v0, v1, a[1:16]
 292 // GFX908: v_mfma_i32_16x16x4i8 a[0:15], v0, v1, a[1:16] ; encoding: [0x00,0x00,0xd1,0xd3,0x00,0x03,0x06,0x04]
 293
 294 v_mfma_i32_16x16x4i8 a[0:15], v0, v1, a[1:16] cbsz:3 abid:2 blgp:7
 295 // GFX908: v_mfma_i32_16x16x4i8 a[0:15], v0, v1, a[1:16] cbsz:3 abid:2 blgp:7 ; encoding: [0x00,0x13,0xd1,0xd3,0x00,0x03,0x06,0xe4]
 296
 297 v_mfma_i32_16x16x4i8 a[0:15], v0, a1, a[1:16]
 298 // GFX908: v_mfma_i32_16x16x4i8 a[0:15], v0, a1, a[1:16] ; encoding: [0x00,0x00,0xd1,0xd3,0x00,0x03,0x06,0x14]
 299
 300 v_mfma_i32_16x16x4i8 a[0:15], v0, a1, a[1:16] cbsz:3 abid:2 blgp:7
 301 // GFX908: v_mfma_i32_16x16x4i8 a[0:15], v0, a1, a[1:16] cbsz:3 abid:2 blgp:7 ; encoding: [0x00,0x13,0xd1,0xd3,0x00,0x03,0x06,0xf4]
 302
 303 v_mfma_i32_16x16x4i8 a[0:15], a0, v1, a[1:16]
 304 // GFX908: v_mfma_i32_16x16x4i8 a[0:15], a0, v1, a[1:16] ; encoding: [0x00,0x00,0xd1,0xd3,0x00,0x03,0x06,0x0c]
 305
 306 v_mfma_i32_16x16x4i8 a[0:15], a0, v1, a[1:16] cbsz:3 abid:2 blgp:7
 307 // GFX908: v_mfma_i32_16x16x4i8 a[0:15], a0, v1, a[1:16] cbsz:3 abid:2 blgp:7 ; encoding: [0x00,0x13,0xd1,0xd3,0x00,0x03,0x06,0xec]
 308
 309 v_mfma_i32_16x16x4i8 a[0:15], a0, a1, a[1:16]
 310 // GFX908: v_mfma_i32_16x16x4i8 a[0:15], a0, a1, a[1:16] ; encoding: [0x00,0x00,0xd1,0xd3,0x00,0x03,0x06,0x1c]
 311
 312 v_mfma_i32_16x16x4i8 a[0:15], a0, a1, a[1:16] cbsz:3 abid:2 blgp:7
 313 // GFX908: v_mfma_i32_16x16x4i8 a[0:15], a0, a1, a[1:16] cbsz:3 abid:2 blgp:7 ; encoding: [0x00,0x13,0xd1,0xd3,0x00,0x03,0x06,0xfc]
 314
 315 v_mfma_i32_4x4x4i8 a[0:3], v0, v1, a[1:4]
 316 // GFX908: v_mfma_i32_4x4x4i8 a[0:3], v0, v1, a[1:4] ; encoding: [0x00,0x00,0xd2,0xd3,0x00,0x03,0x06,0x04]
 317
 318 v_mfma_i32_4x4x4i8 a[0:3], v0, v1, a[1:4] cbsz:3 abid:2 blgp:7
 319 // GFX908: v_mfma_i32_4x4x4i8 a[0:3], v0, v1, a[1:4] cbsz:3 abid:2 blgp:7 ; encoding: [0x00,0x13,0xd2,0xd3,0x00,0x03,0x06,0xe4]
 320
 321 v_mfma_i32_4x4x4i8 a[0:3], v0, a1, a[1:4]
 322 // GFX908: v_mfma_i32_4x4x4i8 a[0:3], v0, a1, a[1:4] ; encoding: [0x00,0x00,0xd2,0xd3,0x00,0x03,0x06,0x14]
 323
 324 v_mfma_i32_4x4x4i8 a[0:3], v0, a1, a[1:4] cbsz:3 abid:2 blgp:7
 325 // GFX908: v_mfma_i32_4x4x4i8 a[0:3], v0, a1, a[1:4] cbsz:3 abid:2 blgp:7 ; encoding: [0x00,0x13,0xd2,0xd3,0x00,0x03,0x06,0xf4]
 326
 327 v_mfma_i32_4x4x4i8 a[0:3], a0, v1, a[1:4]
 328 // GFX908: v_mfma_i32_4x4x4i8 a[0:3], a0, v1, a[1:4] ; encoding: [0x00,0x00,0xd2,0xd3,0x00,0x03,0x06,0x0c]
 329
 330 v_mfma_i32_4x4x4i8 a[0:3], a0, v1, a[1:4] cbsz:3 abid:2 blgp:7
 331 // GFX908: v_mfma_i32_4x4x4i8 a[0:3], a0, v1, a[1:4] cbsz:3 abid:2 blgp:7 ; encoding: [0x00,0x13,0xd2,0xd3,0x00,0x03,0x06,0xec]
 332
 333 v_mfma_i32_4x4x4i8 a[0:3], a0, a1, a[1:4]
 334 // GFX908: v_mfma_i32_4x4x4i8 a[0:3], a0, a1, a[1:4] ; encoding: [0x00,0x00,0xd2,0xd3,0x00,0x03,0x06,0x1c]
 335
 336 v_mfma_i32_4x4x4i8 a[0:3], a0, a1, a[1:4] cbsz:3 abid:2 blgp:7
 337 // GFX908: v_mfma_i32_4x4x4i8 a[0:3], a0, a1, a[1:4] cbsz:3 abid:2 blgp:7 ; encoding: [0x00,0x13,0xd2,0xd3,0x00,0x03,0x06,0xfc]
 338
 339 v_mfma_i32_32x32x8i8 a[0:15], v0, v1, a[1:16]
 340 // GFX908: v_mfma_i32_32x32x8i8 a[0:15], v0, v1, a[1:16] ; encoding: [0x00,0x00,0xd4,0xd3,0x00,0x03,0x06,0x04]
 341
 342 v_mfma_i32_32x32x8i8 a[0:15], v0, v1, a[1:16] cbsz:3 abid:2 blgp:7
 343 // GFX908: v_mfma_i32_32x32x8i8 a[0:15], v0, v1, a[1:16] cbsz:3 abid:2 blgp:7 ; encoding: [0x00,0x13,0xd4,0xd3,0x00,0x03,0x06,0xe4]
 344
 345 v_mfma_i32_32x32x8i8 a[0:15], v0, a1, a[1:16]
 346 // GFX908: v_mfma_i32_32x32x8i8 a[0:15], v0, a1, a[1:16] ; encoding: [0x00,0x00,0xd4,0xd3,0x00,0x03,0x06,0x14]
 347
 348 v_mfma_i32_32x32x8i8 a[0:15], v0, a1, a[1:16] cbsz:3 abid:2 blgp:7
 349 // GFX908: v_mfma_i32_32x32x8i8 a[0:15], v0, a1, a[1:16] cbsz:3 abid:2 blgp:7 ; encoding: [0x00,0x13,0xd4,0xd3,0x00,0x03,0x06,0xf4]
 350
 351 v_mfma_i32_32x32x8i8 a[0:15], a0, v1, a[1:16]
 352 // GFX908: v_mfma_i32_32x32x8i8 a[0:15], a0, v1, a[1:16] ; encoding: [0x00,0x00,0xd4,0xd3,0x00,0x03,0x06,0x0c]
 353
 354 v_mfma_i32_32x32x8i8 a[0:15], a0, v1, a[1:16] cbsz:3 abid:2 blgp:7
 355 // GFX908: v_mfma_i32_32x32x8i8 a[0:15], a0, v1, a[1:16] cbsz:3 abid:2 blgp:7 ; encoding: [0x00,0x13,0xd4,0xd3,0x00,0x03,0x06,0xec]
 356
 357 v_mfma_i32_32x32x8i8 a[0:15], a0, a1, a[1:16]
 358 // GFX908: v_mfma_i32_32x32x8i8 a[0:15], a0, a1, a[1:16] ; encoding: [0x00,0x00,0xd4,0xd3,0x00,0x03,0x06,0x1c]
 359
 360 v_mfma_i32_32x32x8i8 a[0:15], a0, a1, a[1:16] cbsz:3 abid:2 blgp:7
 361 // GFX908: v_mfma_i32_32x32x8i8 a[0:15], a0, a1, a[1:16] cbsz:3 abid:2 blgp:7 ; encoding: [0x00,0x13,0xd4,0xd3,0x00,0x03,0x06,0xfc]
 362
 363 v_mfma_i32_16x16x16i8 a[0:3], v0, v1, a[1:4]
 364 // GFX908: v_mfma_i32_16x16x16i8 a[0:3], v0, v1, a[1:4] ; encoding: [0x00,0x00,0xd5,0xd3,0x00,0x03,0x06,0x04]
 365
 366 v_mfma_i32_16x16x16i8 a[0:3], v0, v1, a[1:4] cbsz:3 abid:2 blgp:7
 367 // GFX908: v_mfma_i32_16x16x16i8 a[0:3], v0, v1, a[1:4] cbsz:3 abid:2 blgp:7 ; encoding: [0x00,0x13,0xd5,0xd3,0x00,0x03,0x06,0xe4]
 368
 369 v_mfma_i32_16x16x16i8 a[0:3], v0, a1, a[1:4]
 370 // GFX908: v_mfma_i32_16x16x16i8 a[0:3], v0, a1, a[1:4] ; encoding: [0x00,0x00,0xd5,0xd3,0x00,0x03,0x06,0x14]
 371
 372 v_mfma_i32_16x16x16i8 a[0:3], v0, a1, a[1:4] cbsz:3 abid:2 blgp:7
 373 // GFX908: v_mfma_i32_16x16x16i8 a[0:3], v0, a1, a[1:4] cbsz:3 abid:2 blgp:7 ; encoding: [0x00,0x13,0xd5,0xd3,0x00,0x03,0x06,0xf4]
 374
 375 v_mfma_i32_16x16x16i8 a[0:3], a0, v1, a[1:4]
 376 // GFX908: v_mfma_i32_16x16x16i8 a[0:3], a0, v1, a[1:4] ; encoding: [0x00,0x00,0xd5,0xd3,0x00,0x03,0x06,0x0c]
 377
 378 v_mfma_i32_16x16x16i8 a[0:3], a0, v1, a[1:4] cbsz:3 abid:2 blgp:7
 379 // GFX908: v_mfma_i32_16x16x16i8 a[0:3], a0, v1, a[1:4] cbsz:3 abid:2 blgp:7 ; encoding: [0x00,0x13,0xd5,0xd3,0x00,0x03,0x06,0xec]
 380
 381 v_mfma_i32_16x16x16i8 a[0:3], a0, a1, a[1:4]
 382 // GFX908: v_mfma_i32_16x16x16i8 a[0:3], a0, a1, a[1:4] ; encoding: [0x00,0x00,0xd5,0xd3,0x00,0x03,0x06,0x1c]
 383
 384 v_mfma_i32_16x16x16i8 a[0:3], a0, a1, a[1:4] cbsz:3 abid:2 blgp:7
 385 // GFX908: v_mfma_i32_16x16x16i8 a[0:3], a0, a1, a[1:4] cbsz:3 abid:2 blgp:7 ; encoding: [0x00,0x13,0xd5,0xd3,0x00,0x03,0x06,0xfc]
 386
 387 v_mfma_f32_32x32x2bf16 a[0:31], v0, v1, a[1:32]
 388 // GFX908: v_mfma_f32_32x32x2bf16 a[0:31], v0, v1, a[1:32] ; encoding: [0x00,0x00,0xe8,0xd3,0x00,0x03,0x06,0x04]
 389
 390 v_mfma_f32_32x32x2bf16 a[0:31], v0, v1, a[1:32] cbsz:3 abid:2 blgp:7
 391 // GFX908: v_mfma_f32_32x32x2bf16 a[0:31], v0, v1, a[1:32] cbsz:3 abid:2 blgp:7 ; encoding: [0x00,0x13,0xe8,0xd3,0x00,0x03,0x06,0xe4]
 392
 393 v_mfma_f32_32x32x2bf16 a[0:31], v0, a1, a[1:32]
 394 // GFX908: v_mfma_f32_32x32x2bf16 a[0:31], v0, a1, a[1:32] ; encoding: [0x00,0x00,0xe8,0xd3,0x00,0x03,0x06,0x14]
 395
 396 v_mfma_f32_32x32x2bf16 a[0:31], v0, a1, a[1:32] cbsz:3 abid:2 blgp:7
 397 // GFX908: v_mfma_f32_32x32x2bf16 a[0:31], v0, a1, a[1:32] cbsz:3 abid:2 blgp:7 ; encoding: [0x00,0x13,0xe8,0xd3,0x00,0x03,0x06,0xf4]
 398
 399 v_mfma_f32_32x32x2bf16 a[0:31], a0, v1, a[1:32]
 400 // GFX908: v_mfma_f32_32x32x2bf16 a[0:31], a0, v1, a[1:32] ; encoding: [0x00,0x00,0xe8,0xd3,0x00,0x03,0x06,0x0c]
 401
 402 v_mfma_f32_32x32x2bf16 a[0:31], a0, v1, a[1:32] cbsz:3 abid:2 blgp:7
 403 // GFX908: v_mfma_f32_32x32x2bf16 a[0:31], a0, v1, a[1:32] cbsz:3 abid:2 blgp:7 ; encoding: [0x00,0x13,0xe8,0xd3,0x00,0x03,0x06,0xec]
 404
 405 v_mfma_f32_32x32x2bf16 a[0:31], a0, a1, a[1:32]
 406 // GFX908: v_mfma_f32_32x32x2bf16 a[0:31], a0, a1, a[1:32] ; encoding: [0x00,0x00,0xe8,0xd3,0x00,0x03,0x06,0x1c]
 407
 408 v_mfma_f32_32x32x2bf16 a[0:31], a0, a1, a[1:32] cbsz:3 abid:2 blgp:7
 409 // GFX908: v_mfma_f32_32x32x2bf16 a[0:31], a0, a1, a[1:32] cbsz:3 abid:2 blgp:7 ; encoding: [0x00,0x13,0xe8,0xd3,0x00,0x03,0x06,0xfc]
 410
 411 v_mfma_f32_16x16x2bf16 a[0:15], v0, v1, a[1:16]
 412 // GFX908: v_mfma_f32_16x16x2bf16 a[0:15], v0, v1, a[1:16] ; encoding: [0x00,0x00,0xe9,0xd3,0x00,0x03,0x06,0x04]
 413
 414 v_mfma_f32_16x16x2bf16 a[0:15], v0, v1, a[1:16] cbsz:3 abid:2 blgp:7
 415 // GFX908: v_mfma_f32_16x16x2bf16 a[0:15], v0, v1, a[1:16] cbsz:3 abid:2 blgp:7 ; encoding: [0x00,0x13,0xe9,0xd3,0x00,0x03,0x06,0xe4]
 416
 417 v_mfma_f32_16x16x2bf16 a[0:15], v0, a1, a[1:16]
 418 // GFX908: v_mfma_f32_16x16x2bf16 a[0:15], v0, a1, a[1:16] ; encoding: [0x00,0x00,0xe9,0xd3,0x00,0x03,0x06,0x14]
 419
 420 v_mfma_f32_16x16x2bf16 a[0:15], v0, a1, a[1:16] cbsz:3 abid:2 blgp:7
 421 // GFX908: v_mfma_f32_16x16x2bf16 a[0:15], v0, a1, a[1:16] cbsz:3 abid:2 blgp:7 ; encoding: [0x00,0x13,0xe9,0xd3,0x00,0x03,0x06,0xf4]
 422
 423 v_mfma_f32_16x16x2bf16 a[0:15], a0, v1, a[1:16]
 424 // GFX908: v_mfma_f32_16x16x2bf16 a[0:15], a0, v1, a[1:16] ; encoding: [0x00,0x00,0xe9,0xd3,0x00,0x03,0x06,0x0c]
 425
 426 v_mfma_f32_16x16x2bf16 a[0:15], a0, v1, a[1:16] cbsz:3 abid:2 blgp:7
 427 // GFX908: v_mfma_f32_16x16x2bf16 a[0:15], a0, v1, a[1:16] cbsz:3 abid:2 blgp:7 ; encoding: [0x00,0x13,0xe9,0xd3,0x00,0x03,0x06,0xec]
 428
 429 v_mfma_f32_16x16x2bf16 a[0:15], a0, a1, a[1:16]
 430 // GFX908: v_mfma_f32_16x16x2bf16 a[0:15], a0, a1, a[1:16] ; encoding: [0x00,0x00,0xe9,0xd3,0x00,0x03,0x06,0x1c]
 431
 432 v_mfma_f32_16x16x2bf16 a[0:15], a0, a1, a[1:16] cbsz:3 abid:2 blgp:7
 433 // GFX908: v_mfma_f32_16x16x2bf16 a[0:15], a0, a1, a[1:16] cbsz:3 abid:2 blgp:7 ; encoding: [0x00,0x13,0xe9,0xd3,0x00,0x03,0x06,0xfc]
 434
 435 v_mfma_f32_4x4x2bf16 a[0:3], v0, v1, a[1:4]
 436 // GFX908: v_mfma_f32_4x4x2bf16 a[0:3], v0, v1, a[1:4] ; encoding: [0x00,0x00,0xeb,0xd3,0x00,0x03,0x06,0x04]
 437
 438 v_mfma_f32_4x4x2bf16 a[0:3], v0, v1, a[1:4] cbsz:3 abid:2 blgp:7
 439 // GFX908: v_mfma_f32_4x4x2bf16 a[0:3], v0, v1, a[1:4] cbsz:3 abid:2 blgp:7 ; encoding: [0x00,0x13,0xeb,0xd3,0x00,0x03,0x06,0xe4]
 440
 441 v_mfma_f32_4x4x2bf16 a[0:3], v0, a1, a[1:4]
 442 // GFX908: v_mfma_f32_4x4x2bf16 a[0:3], v0, a1, a[1:4] ; encoding: [0x00,0x00,0xeb,0xd3,0x00,0x03,0x06,0x14]
 443
 444 v_mfma_f32_4x4x2bf16 a[0:3], v0, a1, a[1:4] cbsz:3 abid:2 blgp:7
 445 // GFX908: v_mfma_f32_4x4x2bf16 a[0:3], v0, a1, a[1:4] cbsz:3 abid:2 blgp:7 ; encoding: [0x00,0x13,0xeb,0xd3,0x00,0x03,0x06,0xf4]
 446
 447 v_mfma_f32_4x4x2bf16 a[0:3], a0, v1, a[1:4]
 448 // GFX908: v_mfma_f32_4x4x2bf16 a[0:3], a0, v1, a[1:4] ; encoding: [0x00,0x00,0xeb,0xd3,0x00,0x03,0x06,0x0c]
 449
 450 v_mfma_f32_4x4x2bf16 a[0:3], a0, v1, a[1:4] cbsz:3 abid:2 blgp:7
 451 // GFX908: v_mfma_f32_4x4x2bf16 a[0:3], a0, v1, a[1:4] cbsz:3 abid:2 blgp:7 ; encoding: [0x00,0x13,0xeb,0xd3,0x00,0x03,0x06,0xec]
 452
 453 v_mfma_f32_4x4x2bf16 a[0:3], a0, a1, a[1:4]
 454 // GFX908: v_mfma_f32_4x4x2bf16 a[0:3], a0, a1, a[1:4] ; encoding: [0x00,0x00,0xeb,0xd3,0x00,0x03,0x06,0x1c]
 455
 456 v_mfma_f32_4x4x2bf16 a[0:3], a0, a1, a[1:4] cbsz:3 abid:2 blgp:7
 457 // GFX908: v_mfma_f32_4x4x2bf16 a[0:3], a0, a1, a[1:4] cbsz:3 abid:2 blgp:7 ; encoding: [0x00,0x13,0xeb,0xd3,0x00,0x03,0x06,0xfc]
 458
 459 v_mfma_f32_32x32x4bf16 a[0:15], v0, v1, a[1:16]
 460 // GFX908: v_mfma_f32_32x32x4bf16 a[0:15], v0, v1, a[1:16] ; encoding: [0x00,0x00,0xec,0xd3,0x00,0x03,0x06,0x04]
 461
 462 v_mfma_f32_32x32x4bf16 a[0:15], v0, v1, a[1:16] cbsz:3 abid:2 blgp:7
 463 // GFX908: v_mfma_f32_32x32x4bf16 a[0:15], v0, v1, a[1:16] cbsz:3 abid:2 blgp:7 ; encoding: [0x00,0x13,0xec,0xd3,0x00,0x03,0x06,0xe4]
 464
 465 v_mfma_f32_32x32x4bf16 a[0:15], v0, a1, a[1:16]
 466 // GFX908: v_mfma_f32_32x32x4bf16 a[0:15], v0, a1, a[1:16] ; encoding: [0x00,0x00,0xec,0xd3,0x00,0x03,0x06,0x14]
 467
 468 v_mfma_f32_32x32x4bf16 a[0:15], v0, a1, a[1:16] cbsz:3 abid:2 blgp:7
 469 // GFX908: v_mfma_f32_32x32x4bf16 a[0:15], v0, a1, a[1:16] cbsz:3 abid:2 blgp:7 ; encoding: [0x00,0x13,0xec,0xd3,0x00,0x03,0x06,0xf4]
 470
 471 v_mfma_f32_32x32x4bf16 a[0:15], a0, v1, a[1:16]
 472 // GFX908: v_mfma_f32_32x32x4bf16 a[0:15], a0, v1, a[1:16] ; encoding: [0x00,0x00,0xec,0xd3,0x00,0x03,0x06,0x0c]
 473
 474 v_mfma_f32_32x32x4bf16 a[0:15], a0, v1, a[1:16] cbsz:3 abid:2 blgp:7
 475 // GFX908: v_mfma_f32_32x32x4bf16 a[0:15], a0, v1, a[1:16] cbsz:3 abid:2 blgp:7 ; encoding: [0x00,0x13,0xec,0xd3,0x00,0x03,0x06,0xec]
 476
 477 v_mfma_f32_32x32x4bf16 a[0:15], a0, a1, a[1:16]
 478 // GFX908: v_mfma_f32_32x32x4bf16 a[0:15], a0, a1, a[1:16] ; encoding: [0x00,0x00,0xec,0xd3,0x00,0x03,0x06,0x1c]
 479
 480 v_mfma_f32_32x32x4bf16 a[0:15], a0, a1, a[1:16] cbsz:3 abid:2 blgp:7
 481 // GFX908: v_mfma_f32_32x32x4bf16 a[0:15], a0, a1, a[1:16] cbsz:3 abid:2 blgp:7 ; encoding: [0x00,0x13,0xec,0xd3,0x00,0x03,0x06,0xfc]
 482
 483 v_mfma_f32_16x16x8bf16 a[0:3], v0, v1, a[1:4]
 484 // GFX908: v_mfma_f32_16x16x8bf16 a[0:3], v0, v1, a[1:4] ; encoding: [0x00,0x00,0xed,0xd3,0x00,0x03,0x06,0x04]
 485
 486 v_mfma_f32_16x16x8bf16 a[0:3], v0, v1, a[1:4] cbsz:3 abid:2 blgp:7
 487 // GFX908: v_mfma_f32_16x16x8bf16 a[0:3], v0, v1, a[1:4] cbsz:3 abid:2 blgp:7 ; encoding: [0x00,0x13,0xed,0xd3,0x00,0x03,0x06,0xe4]
 488
 489 v_mfma_f32_16x16x8bf16 a[0:3], v0, a1, a[1:4]
 490 // GFX908: v_mfma_f32_16x16x8bf16 a[0:3], v0, a1, a[1:4] ; encoding: [0x00,0x00,0xed,0xd3,0x00,0x03,0x06,0x14]
 491
 492 v_mfma_f32_16x16x8bf16 a[0:3], v0, a1, a[1:4] cbsz:3 abid:2 blgp:7
 493 // GFX908: v_mfma_f32_16x16x8bf16 a[0:3], v0, a1, a[1:4] cbsz:3 abid:2 blgp:7 ; encoding: [0x00,0x13,0xed,0xd3,0x00,0x03,0x06,0xf4]
 494
 495 v_mfma_f32_16x16x8bf16 a[0:3], a0, v1, a[1:4]
 496 // GFX908: v_mfma_f32_16x16x8bf16 a[0:3], a0, v1, a[1:4] ; encoding: [0x00,0x00,0xed,0xd3,0x00,0x03,0x06,0x0c]
 497
 498 v_mfma_f32_16x16x8bf16 a[0:3], a0, v1, a[1:4] cbsz:3 abid:2 blgp:7
 499 // GFX908: v_mfma_f32_16x16x8bf16 a[0:3], a0, v1, a[1:4] cbsz:3 abid:2 blgp:7 ; encoding: [0x00,0x13,0xed,0xd3,0x00,0x03,0x06,0xec]
 500
 501 v_mfma_f32_16x16x8bf16 a[0:3], a0, a1, a[1:4]
 502 // GFX908: v_mfma_f32_16x16x8bf16 a[0:3], a0, a1, a[1:4] ; encoding: [0x00,0x00,0xed,0xd3,0x00,0x03,0x06,0x1c]
 503
 504 v_mfma_f32_16x16x8bf16 a[0:3], a0, a1, a[1:4] cbsz:3 abid:2 blgp:7
 505 // GFX908: v_mfma_f32_16x16x8bf16 a[0:3], a0, a1, a[1:4] cbsz:3 abid:2 blgp:7 ; encoding: [0x00,0x13,0xed,0xd3,0x00,0x03,0x06,0xfc]