llvm/test/MC/AMDGPU/mai.s

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