llvm/test/MC/AMDGPU/mai-gfx940.s

   1 // RUN: llvm-mc -triple=amdgcn -mcpu=gfx940 -show-encoding %s | FileCheck -check-prefix=GFX940 %s
   2 // RUN: not llvm-mc -triple=amdgcn -mcpu=gfx90a %s 2>&1 | FileCheck -check-prefix=GFX90A %s
   3
   4 //===----------------------------------------------------------------------===//
   5 // Misc opcodes.
   6 //===----------------------------------------------------------------------===//
   7
   8 v_accvgpr_write_b32 a10, s20
   9 // GFX940: v_accvgpr_write_b32 a10, s20    ; encoding: [0x0a,0x40,0xd9,0xd3,0x14,0x00,0x00,0x18]
  10
  11 //===----------------------------------------------------------------------===//
  12 // MFMA opcodes.
  13 //===----------------------------------------------------------------------===//
  14
  15 v_mfma_f64_4x4x4_4b_f64 a[0:1], v[0:1], a[2:3], a[2:3]
  16 // GFX940: v_mfma_f64_4x4x4_4b_f64 a[0:1], v[0:1], a[2:3], a[2:3] ; encoding: [0x00,0x80,0xef,0xd3,0x00,0x05,0x0a,0x14]
  17 // GFX90A: :[[@LINE-2]]:{{[0-9]+}}: error: instruction not supported on this GPU
  18
  19 v_mfma_f64_4x4x4_4b_f64 v[0:1], v[0:1], a[2:3], v[2:3]
  20 // GFX940: v_mfma_f64_4x4x4_4b_f64 v[0:1], v[0:1], a[2:3], v[2:3] ; encoding: [0x00,0x00,0xef,0xd3,0x00,0x05,0x0a,0x14]
  21 // GFX90A: :[[@LINE-2]]:{{[0-9]+}}: error: instruction not supported on this GPU
  22
  23 v_mfma_f64_4x4x4f64 a[0:1], v[0:1], a[2:3], a[2:3]
  24 // GFX940: v_mfma_f64_4x4x4_4b_f64 a[0:1], v[0:1], a[2:3], a[2:3] ; encoding: [0x00,0x80,0xef,0xd3,0x00,0x05,0x0a,0x14]
  25
  26 v_mfma_f64_4x4x4f64 v[0:1], v[0:1], a[2:3], v[2:3]
  27 // GFX940: v_mfma_f64_4x4x4_4b_f64 v[0:1], v[0:1], a[2:3], v[2:3] ; encoding: [0x00,0x00,0xef,0xd3,0x00,0x05,0x0a,0x14]
  28
  29 v_mfma_f64_4x4x4_4b_f64 a[0:1], v[0:1], a[2:3], a[2:3] neg:[1,0,0]
  30 // GFX940: v_mfma_f64_4x4x4_4b_f64 a[0:1], v[0:1], a[2:3], a[2:3] neg:[1,0,0] ; encoding: [0x00,0x80,0xef,0xd3,0x00,0x05,0x0a,0x34]
  31 // GFX90A: :[[@LINE-2]]:{{[0-9]+}}: error: instruction not supported on this GPU
  32
  33 v_mfma_f64_4x4x4_4b_f64 a[0:1], v[0:1], a[2:3], a[2:3] neg:[0,1,0]
  34 // GFX940: v_mfma_f64_4x4x4_4b_f64 a[0:1], v[0:1], a[2:3], a[2:3] neg:[0,1,0] ; encoding: [0x00,0x80,0xef,0xd3,0x00,0x05,0x0a,0x54]
  35 // GFX90A: :[[@LINE-2]]:{{[0-9]+}}: error: instruction not supported on this GPU
  36
  37 v_mfma_f64_4x4x4_4b_f64 a[0:1], v[0:1], a[2:3], a[2:3] neg:[0,0,1]
  38 // GFX940: v_mfma_f64_4x4x4_4b_f64 a[0:1], v[0:1], a[2:3], a[2:3] neg:[0,0,1] ; encoding: [0x00,0x80,0xef,0xd3,0x00,0x05,0x0a,0x94]
  39 // GFX90A: :[[@LINE-2]]:{{[0-9]+}}: error: instruction not supported on this GPU
  40
  41 v_mfma_f64_4x4x4_4b_f64 v[0:1], v[0:1], a[2:3], v[2:3] neg:[1,1,1]
  42 // GFX940: v_mfma_f64_4x4x4_4b_f64 v[0:1], v[0:1], a[2:3], v[2:3] neg:[1,1,1] ; encoding: [0x00,0x00,0xef,0xd3,0x00,0x05,0x0a,0xf4]
  43 // GFX90A: :[[@LINE-2]]:{{[0-9]+}}: error: instruction not supported on this GPU
  44
  45 v_mfma_f64_4x4x4f64 a[0:1], v[0:1], a[2:3], a[2:3] neg:[1,0,0]
  46 // GFX940: v_mfma_f64_4x4x4_4b_f64 a[0:1], v[0:1], a[2:3], a[2:3] neg:[1,0,0] ; encoding: [0x00,0x80,0xef,0xd3,0x00,0x05,0x0a,0x34]
  47 // GFX90A: :[[@LINE-2]]:{{[0-9]+}}: error: invalid modifier: neg is not supported
  48
  49 v_mfma_f64_4x4x4f64 v[0:1], v[0:1], a[2:3], v[2:3] neg:[1,0,0]
  50 // GFX940: v_mfma_f64_4x4x4_4b_f64 v[0:1], v[0:1], a[2:3], v[2:3] neg:[1,0,0] ; encoding: [0x00,0x00,0xef,0xd3,0x00,0x05,0x0a,0x34]
  51 // GFX90A: :[[@LINE-2]]:{{[0-9]+}}: error: invalid modifier: neg is not supported
  52
  53 v_mfma_f64_16x16x4_f64 a[0:7], v[0:1], v[2:3], a[0:7]
  54 // GFX940: v_mfma_f64_16x16x4_f64 a[0:7], v[0:1], v[2:3], a[0:7] ; encoding: [0x00,0x80,0xee,0xd3,0x00,0x05,0x02,0x04]
  55 // GFX90A: :[[@LINE-2]]:{{[0-9]+}}: error: instruction not supported on this GPU
  56
  57 v_mfma_f64_16x16x4_f64 v[0:7], v[0:1], v[2:3], v[0:7]
  58 // GFX940: v_mfma_f64_16x16x4_f64 v[0:7], v[0:1], v[2:3], v[0:7] ; encoding: [0x00,0x00,0xee,0xd3,0x00,0x05,0x02,0x04]
  59 // GFX90A: :[[@LINE-2]]:{{[0-9]+}}: error: instruction not supported on this GPU
  60
  61 v_mfma_f64_16x16x4f64 a[0:7], v[0:1], v[2:3], a[0:7]
  62 // GFX940: v_mfma_f64_16x16x4_f64 a[0:7], v[0:1], v[2:3], a[0:7] ; encoding: [0x00,0x80,0xee,0xd3,0x00,0x05,0x02,0x04]
  63
  64 v_mfma_f64_16x16x4f64 v[0:7], v[0:1], v[2:3], v[0:7]
  65 // GFX940: v_mfma_f64_16x16x4_f64 v[0:7], v[0:1], v[2:3], v[0:7] ; encoding: [0x00,0x00,0xee,0xd3,0x00,0x05,0x02,0x04]
  66
  67 v_mfma_f64_16x16x4_f64 a[0:7], v[0:1], v[2:3], a[0:7] neg:[1,1,1]
  68 // GFX940: v_mfma_f64_16x16x4_f64 a[0:7], v[0:1], v[2:3], a[0:7] neg:[1,1,1] ; encoding: [0x00,0x80,0xee,0xd3,0x00,0x05,0x02,0xe4]
  69 // GFX90A: :[[@LINE-2]]:{{[0-9]+}}: error: instruction not supported on this GPU
  70
  71 v_mfma_f64_16x16x4_f64 v[0:7], v[0:1], v[2:3], v[0:7] neg:[1,1,1]
  72 // GFX940: v_mfma_f64_16x16x4_f64 v[0:7], v[0:1], v[2:3], v[0:7] neg:[1,1,1] ; encoding: [0x00,0x00,0xee,0xd3,0x00,0x05,0x02,0xe4]
  73 // GFX90A: :[[@LINE-2]]:{{[0-9]+}}: error: instruction not supported on this GPU
  74
  75 v_mfma_f64_16x16x4f64 a[0:7], v[0:1], v[2:3], a[0:7] neg:[1,0,0]
  76 // GFX940: v_mfma_f64_16x16x4_f64 a[0:7], v[0:1], v[2:3], a[0:7] neg:[1,0,0] ; encoding: [0x00,0x80,0xee,0xd3,0x00,0x05,0x02,0x24]
  77 // GFX90A: :[[@LINE-2]]:{{[0-9]+}}: error: invalid modifier: neg is not supported
  78
  79 v_mfma_f64_16x16x4f64 v[0:7], v[0:1], v[2:3], v[0:7] neg:[1,0,0]
  80 // GFX940: v_mfma_f64_16x16x4_f64 v[0:7], v[0:1], v[2:3], v[0:7] neg:[1,0,0] ; encoding: [0x00,0x00,0xee,0xd3,0x00,0x05,0x02,0x24]
  81 // GFX90A: :[[@LINE-2]]:{{[0-9]+}}: error: invalid modifier: neg is not supported
  82
  83 v_mfma_f32_16x16x1_4b_f32 a[0:15], v0, v1, a[18:33]
  84 // GFX940: v_mfma_f32_16x16x1_4b_f32 a[0:15], v0, v1, a[18:33] ; encoding: [0x00,0x80,0xc1,0xd3,0x00,0x03,0x4a,0x04]
  85 // GFX90A: :[[@LINE-2]]:{{[0-9]+}}: error: instruction not supported on this GPU
  86
  87 v_mfma_f32_16x16x1_4b_f32 v[0:15], v0, v1, v[18:33]
  88 // GFX940: v_mfma_f32_16x16x1_4b_f32 v[0:15], v0, v1, v[18:33] ; encoding: [0x00,0x00,0xc1,0xd3,0x00,0x03,0x4a,0x04]
  89 // GFX90A: :[[@LINE-2]]:{{[0-9]+}}: error: instruction not supported on this GPU
  90
  91 v_mfma_f32_16x16x1f32 a[0:15], v0, v1, a[18:33]
  92 // GFX940: v_mfma_f32_16x16x1_4b_f32 a[0:15], v0, v1, a[18:33] ; encoding: [0x00,0x80,0xc1,0xd3,0x00,0x03,0x4a,0x04]
  93
  94 v_mfma_f32_16x16x1f32 v[0:15], v0, v1, v[18:33]
  95 // GFX940: v_mfma_f32_16x16x1_4b_f32 v[0:15], v0, v1, v[18:33] ; encoding: [0x00,0x00,0xc1,0xd3,0x00,0x03,0x4a,0x04]
  96
  97 v_mfma_f32_4x4x1_16b_f32 a[0:3], v0, v1, a[2:5]
  98 // GFX940: v_mfma_f32_4x4x1_16b_f32 a[0:3], v0, v1, a[2:5] ; encoding: [0x00,0x80,0xc2,0xd3,0x00,0x03,0x0a,0x04]
  99 // GFX90A: :[[@LINE-2]]:{{[0-9]+}}: error: instruction not supported on this GPU
 100
 101 v_mfma_f32_4x4x1_16b_f32 v[0:3], v0, v1, v[2:5]
 102 // GFX940: v_mfma_f32_4x4x1_16b_f32 v[0:3], v0, v1, v[2:5] ; encoding: [0x00,0x00,0xc2,0xd3,0x00,0x03,0x0a,0x04]
 103 // GFX90A: :[[@LINE-2]]:{{[0-9]+}}: error: instruction not supported on this GPU
 104
 105 v_mfma_f32_4x4x1f32 a[0:3], v0, v1, a[2:5]
 106 // GFX940: v_mfma_f32_4x4x1_16b_f32 a[0:3], v0, v1, a[2:5] ; encoding: [0x00,0x80,0xc2,0xd3,0x00,0x03,0x0a,0x04]
 107
 108 v_mfma_f32_4x4x1f32 v[0:3], v0, v1, v[2:5]
 109 // GFX940: v_mfma_f32_4x4x1_16b_f32 v[0:3], v0, v1, v[2:5] ; encoding: [0x00,0x00,0xc2,0xd3,0x00,0x03,0x0a,0x04]
 110
 111 v_mfma_f32_32x32x2_f32 a[0:15], v0, v1, a[18:33]
 112 // GFX940: v_mfma_f32_32x32x2_f32 a[0:15], v0, v1, a[18:33] ; encoding: [0x00,0x80,0xc4,0xd3,0x00,0x03,0x4a,0x04]
 113 // GFX90A: :[[@LINE-2]]:{{[0-9]+}}: error: instruction not supported on this GPU
 114
 115 v_mfma_f32_32x32x2_f32 v[0:15], v0, v1, v[18:33]
 116 // GFX940: v_mfma_f32_32x32x2_f32 v[0:15], v0, v1, v[18:33] ; encoding: [0x00,0x00,0xc4,0xd3,0x00,0x03,0x4a,0x04]
 117 // GFX90A: :[[@LINE-2]]:{{[0-9]+}}: error: instruction not supported on this GPU
 118
 119 v_mfma_f32_32x32x2f32 a[0:15], v0, v1, a[18:33]
 120 // GFX940: v_mfma_f32_32x32x2_f32 a[0:15], v0, v1, a[18:33] ; encoding: [0x00,0x80,0xc4,0xd3,0x00,0x03,0x4a,0x04]
 121
 122 v_mfma_f32_32x32x2f32 v[0:15], v0, v1, v[18:33]
 123 // GFX940: v_mfma_f32_32x32x2_f32 v[0:15], v0, v1, v[18:33] ; encoding: [0x00,0x00,0xc4,0xd3,0x00,0x03,0x4a,0x04]
 124
 125 v_mfma_f32_16x16x4_f32 a[0:3], v0, v1, a[2:5]
 126 // GFX940: v_mfma_f32_16x16x4_f32 a[0:3], v0, v1, a[2:5] ; encoding: [0x00,0x80,0xc5,0xd3,0x00,0x03,0x0a,0x04]
 127 // GFX90A: :[[@LINE-2]]:{{[0-9]+}}: error: instruction not supported on this GPU
 128
 129 v_mfma_f32_16x16x4_f32 v[0:3], v0, v1, v[2:5]
 130 // GFX940: v_mfma_f32_16x16x4_f32 v[0:3], v0, v1, v[2:5] ; encoding: [0x00,0x00,0xc5,0xd3,0x00,0x03,0x0a,0x04]
 131 // GFX90A: :[[@LINE-2]]:{{[0-9]+}}: error: instruction not supported on this GPU
 132
 133 v_mfma_f32_16x16x4f32 a[0:3], v0, v1, a[2:5]
 134 // GFX940: v_mfma_f32_16x16x4_f32 a[0:3], v0, v1, a[2:5] ; encoding: [0x00,0x80,0xc5,0xd3,0x00,0x03,0x0a,0x04]
 135
 136 v_mfma_f32_16x16x4f32 v[0:3], v0, v1, v[2:5]
 137 // GFX940: v_mfma_f32_16x16x4_f32 v[0:3], v0, v1, v[2:5] ; encoding: [0x00,0x00,0xc5,0xd3,0x00,0x03,0x0a,0x04]
 138
 139 v_mfma_f32_32x32x4_2b_f16 a[0:31], v[0:1], v[2:3], a[34:65]
 140 // GFX940: v_mfma_f32_32x32x4_2b_f16 a[0:31], v[0:1], v[2:3], a[34:65] ; encoding: [0x00,0x80,0xc8,0xd3,0x00,0x05,0x8a,0x04]
 141 // GFX90A: :[[@LINE-2]]:{{[0-9]+}}: error: instruction not supported on this GPU
 142
 143 v_mfma_f32_32x32x4_2b_f16 v[0:31], v[0:1], v[2:3], v[34:65]
 144 // GFX940: v_mfma_f32_32x32x4_2b_f16 v[0:31], v[0:1], v[2:3], v[34:65] ; encoding: [0x00,0x00,0xc8,0xd3,0x00,0x05,0x8a,0x04]
 145 // GFX90A: :[[@LINE-2]]:{{[0-9]+}}: error: instruction not supported on this GPU
 146
 147 v_mfma_f32_32x32x4f16 a[0:31], v[0:1], v[2:3], a[34:65]
 148 // GFX940: v_mfma_f32_32x32x4_2b_f16 a[0:31], v[0:1], v[2:3], a[34:65] ; encoding: [0x00,0x80,0xc8,0xd3,0x00,0x05,0x8a,0x04]
 149
 150 v_mfma_f32_32x32x4f16 v[0:31], v[0:1], v[2:3], v[34:65]
 151 // GFX940: v_mfma_f32_32x32x4_2b_f16 v[0:31], v[0:1], v[2:3], v[34:65] ; encoding: [0x00,0x00,0xc8,0xd3,0x00,0x05,0x8a,0x04]
 152
 153 v_mfma_f32_16x16x4_4b_f16 a[0:15], v[0:1], v[2:3], a[18:33]
 154 // GFX940: v_mfma_f32_16x16x4_4b_f16 a[0:15], v[0:1], v[2:3], a[18:33] ; encoding: [0x00,0x80,0xc9,0xd3,0x00,0x05,0x4a,0x04]
 155 // GFX90A: :[[@LINE-2]]:{{[0-9]+}}: error: instruction not supported on this GPU
 156
 157 v_mfma_f32_16x16x4_4b_f16 v[0:15], v[0:1], v[2:3], v[18:33]
 158 // GFX940: v_mfma_f32_16x16x4_4b_f16 v[0:15], v[0:1], v[2:3], v[18:33] ; encoding: [0x00,0x00,0xc9,0xd3,0x00,0x05,0x4a,0x04]
 159 // GFX90A: :[[@LINE-2]]:{{[0-9]+}}: error: instruction not supported on this GPU
 160
 161 v_mfma_f32_16x16x4f16 a[0:15], v[0:1], v[2:3], a[18:33]
 162 // GFX940: v_mfma_f32_16x16x4_4b_f16 a[0:15], v[0:1], v[2:3], a[18:33] ; encoding: [0x00,0x80,0xc9,0xd3,0x00,0x05,0x4a,0x04]
 163
 164 v_mfma_f32_16x16x4f16 v[0:15], v[0:1], v[2:3], v[18:33]
 165 // GFX940: v_mfma_f32_16x16x4_4b_f16 v[0:15], v[0:1], v[2:3], v[18:33] ; encoding: [0x00,0x00,0xc9,0xd3,0x00,0x05,0x4a,0x04]
 166
 167 v_mfma_f32_4x4x4_16b_f16 a[0:3], v[0:1], v[2:3], a[2:5]
 168 // GFX940: v_mfma_f32_4x4x4_16b_f16 a[0:3], v[0:1], v[2:3], a[2:5] ; encoding: [0x00,0x80,0xca,0xd3,0x00,0x05,0x0a,0x04]
 169 // GFX90A: :[[@LINE-2]]:{{[0-9]+}}: error: instruction not supported on this GPU
 170
 171 v_mfma_f32_4x4x4_16b_f16 v[0:3], v[0:1], v[2:3], v[2:5]
 172 // GFX940: v_mfma_f32_4x4x4_16b_f16 v[0:3], v[0:1], v[2:3], v[2:5] ; encoding: [0x00,0x00,0xca,0xd3,0x00,0x05,0x0a,0x04]
 173 // GFX90A: :[[@LINE-2]]:{{[0-9]+}}: error: instruction not supported on this GPU
 174
 175 v_mfma_f32_4x4x4f16 a[0:3], v[0:1], v[2:3], a[2:5]
 176 // GFX940: v_mfma_f32_4x4x4_16b_f16 a[0:3], v[0:1], v[2:3], a[2:5] ; encoding: [0x00,0x80,0xca,0xd3,0x00,0x05,0x0a,0x04]
 177
 178 v_mfma_f32_4x4x4f16 v[0:3], v[0:1], v[2:3], v[2:5]
 179 // GFX940: v_mfma_f32_4x4x4_16b_f16 v[0:3], v[0:1], v[2:3], v[2:5] ; encoding: [0x00,0x00,0xca,0xd3,0x00,0x05,0x0a,0x04]
 180
 181 v_mfma_f32_32x32x8_f16 a[0:15], v[0:1], v[2:3], a[18:33]
 182 // GFX940: v_mfma_f32_32x32x8_f16 a[0:15], v[0:1], v[2:3], a[18:33] ; encoding: [0x00,0x80,0xcc,0xd3,0x00,0x05,0x4a,0x04]
 183 // GFX90A: :[[@LINE-2]]:{{[0-9]+}}: error: instruction not supported on this GPU
 184
 185 v_mfma_f32_32x32x8_f16 v[0:15], v[0:1], v[2:3], v[18:33]
 186 // GFX940: v_mfma_f32_32x32x8_f16 v[0:15], v[0:1], v[2:3], v[18:33] ; encoding: [0x00,0x00,0xcc,0xd3,0x00,0x05,0x4a,0x04]
 187 // GFX90A: :[[@LINE-2]]:{{[0-9]+}}: error: instruction not supported on this GPU
 188
 189 v_mfma_f32_32x32x8f16 a[0:15], v[0:1], v[2:3], a[18:33]
 190 // GFX940: v_mfma_f32_32x32x8_f16 a[0:15], v[0:1], v[2:3], a[18:33] ; encoding: [0x00,0x80,0xcc,0xd3,0x00,0x05,0x4a,0x04]
 191
 192 v_mfma_f32_32x32x8f16 v[0:15], v[0:1], v[2:3], v[18:33]
 193 // GFX940: v_mfma_f32_32x32x8_f16 v[0:15], v[0:1], v[2:3], v[18:33] ; encoding: [0x00,0x00,0xcc,0xd3,0x00,0x05,0x4a,0x04]
 194
 195 v_mfma_f32_16x16x16_f16 a[0:3], v[0:1], v[2:3], a[2:5]
 196 // GFX940: v_mfma_f32_16x16x16_f16 a[0:3], v[0:1], v[2:3], a[2:5] ; encoding: [0x00,0x80,0xcd,0xd3,0x00,0x05,0x0a,0x04]
 197 // GFX90A: :[[@LINE-2]]:{{[0-9]+}}: error: instruction not supported on this GPU
 198
 199 v_mfma_f32_16x16x16_f16 v[0:3], v[0:1], v[2:3], v[2:5]
 200 // GFX940: v_mfma_f32_16x16x16_f16 v[0:3], v[0:1], v[2:3], v[2:5] ; encoding: [0x00,0x00,0xcd,0xd3,0x00,0x05,0x0a,0x04]
 201 // GFX90A: :[[@LINE-2]]:{{[0-9]+}}: error: instruction not supported on this GPU
 202
 203 v_mfma_f32_16x16x16f16 a[0:3], v[0:1], v[2:3], a[2:5]
 204 // GFX940: v_mfma_f32_16x16x16_f16 a[0:3], v[0:1], v[2:3], a[2:5] ; encoding: [0x00,0x80,0xcd,0xd3,0x00,0x05,0x0a,0x04]
 205
 206 v_mfma_f32_16x16x16f16 v[0:3], v[0:1], v[2:3], v[2:5]
 207 // GFX940: v_mfma_f32_16x16x16_f16 v[0:3], v[0:1], v[2:3], v[2:5] ; encoding: [0x00,0x00,0xcd,0xd3,0x00,0x05,0x0a,0x04]
 208
 209 v_mfma_i32_32x32x4_2b_i8 a[0:31], v0, v1, a[34:65]
 210 // GFX940: v_mfma_i32_32x32x4_2b_i8 a[0:31], v0, v1, a[34:65] ; encoding: [0x00,0x80,0xd0,0xd3,0x00,0x03,0x8a,0x04]
 211 // GFX90A: :[[@LINE-2]]:{{[0-9]+}}: error: instruction not supported on this GPU
 212
 213 v_mfma_i32_32x32x4_2b_i8 v[0:31], v0, a1, v[34:65]
 214 // GFX940: v_mfma_i32_32x32x4_2b_i8 v[0:31], v0, a1, v[34:65] ; encoding: [0x00,0x00,0xd0,0xd3,0x00,0x03,0x8a,0x14]
 215 // GFX90A: :[[@LINE-2]]:{{[0-9]+}}: error: instruction not supported on this GPU
 216
 217 v_mfma_i32_32x32x4i8 a[0:31], v0, v1, a[34:65]
 218 // GFX940: v_mfma_i32_32x32x4_2b_i8 a[0:31], v0, v1, a[34:65] ; encoding: [0x00,0x80,0xd0,0xd3,0x00,0x03,0x8a,0x04]
 219
 220 v_mfma_i32_32x32x4i8 v[0:31], v0, a1, v[34:65]
 221 // GFX940: v_mfma_i32_32x32x4_2b_i8 v[0:31], v0, a1, v[34:65] ; encoding: [0x00,0x00,0xd0,0xd3,0x00,0x03,0x8a,0x14]
 222
 223 v_mfma_i32_16x16x4_4b_i8 a[0:15], v0, v1, a[18:33]
 224 // GFX940: v_mfma_i32_16x16x4_4b_i8 a[0:15], v0, v1, a[18:33] ; encoding: [0x00,0x80,0xd1,0xd3,0x00,0x03,0x4a,0x04]
 225 // GFX90A: :[[@LINE-2]]:{{[0-9]+}}: error: instruction not supported on this GPU
 226
 227 v_mfma_i32_16x16x4_4b_i8 v[0:15], v0, v1, v[18:33]
 228 // GFX940: v_mfma_i32_16x16x4_4b_i8 v[0:15], v0, v1, v[18:33] ; encoding: [0x00,0x00,0xd1,0xd3,0x00,0x03,0x4a,0x04]
 229 // GFX90A: :[[@LINE-2]]:{{[0-9]+}}: error: instruction not supported on this GPU
 230
 231 v_mfma_i32_16x16x4i8 a[0:15], v0, v1, a[18:33]
 232 // GFX940: v_mfma_i32_16x16x4_4b_i8 a[0:15], v0, v1, a[18:33] ; encoding: [0x00,0x80,0xd1,0xd3,0x00,0x03,0x4a,0x04]
 233
 234 v_mfma_i32_16x16x4i8 v[0:15], v0, v1, v[18:33]
 235 // GFX940: v_mfma_i32_16x16x4_4b_i8 v[0:15], v0, v1, v[18:33] ; encoding: [0x00,0x00,0xd1,0xd3,0x00,0x03,0x4a,0x04]
 236
 237 v_mfma_i32_4x4x4_16b_i8 a[0:3], v0, v1, a[2:5]
 238 // GFX940: v_mfma_i32_4x4x4_16b_i8 a[0:3], v0, v1, a[2:5] ; encoding: [0x00,0x80,0xd2,0xd3,0x00,0x03,0x0a,0x04]
 239 // GFX90A: :[[@LINE-2]]:{{[0-9]+}}: error: instruction not supported on this GPU
 240
 241 v_mfma_i32_4x4x4_16b_i8 v[0:3], v0, v1, v[2:5]
 242 // GFX940: v_mfma_i32_4x4x4_16b_i8 v[0:3], v0, v1, v[2:5] ; encoding: [0x00,0x00,0xd2,0xd3,0x00,0x03,0x0a,0x04]
 243 // GFX90A: :[[@LINE-2]]:{{[0-9]+}}: error: instruction not supported on this GPU
 244
 245 v_mfma_i32_4x4x4i8 a[0:3], v0, v1, a[2:5]
 246 // GFX940: v_mfma_i32_4x4x4_16b_i8 a[0:3], v0, v1, a[2:5] ; encoding: [0x00,0x80,0xd2,0xd3,0x00,0x03,0x0a,0x04]
 247
 248 v_mfma_i32_4x4x4i8 v[0:3], v0, v1, v[2:5]
 249 // GFX940: v_mfma_i32_4x4x4_16b_i8 v[0:3], v0, v1, v[2:5] ; encoding: [0x00,0x00,0xd2,0xd3,0x00,0x03,0x0a,0x04]
 250
 251 v_mfma_f32_32x32x1_2b_f32 a[0:31], v0, v1, a[34:65] blgp:7
 252 // GFX940: v_mfma_f32_32x32x1_2b_f32 a[0:31], v0, v1, a[34:65] blgp:7 ; encoding: [0x00,0x80,0xc0,0xd3,0x00,0x03,0x8a,0xe4]
 253 // GFX90A: :[[@LINE-2]]:{{[0-9]+}}: error: instruction not supported on this GPU
 254
 255 v_mfma_f32_32x32x1_2b_f32 v[0:31], v0, v1, v[34:65] blgp:7
 256 // GFX940: v_mfma_f32_32x32x1_2b_f32 v[0:31], v0, v1, v[34:65] blgp:7 ; encoding: [0x00,0x00,0xc0,0xd3,0x00,0x03,0x8a,0xe4]
 257 // GFX90A: :[[@LINE-2]]:{{[0-9]+}}: error: instruction not supported on this GPU
 258
 259 v_mfma_f32_32x32x1_2b_f32 a[0:31], v0, v1, a[34:65] blgp:7
 260 // GFX940: v_mfma_f32_32x32x1_2b_f32 a[0:31], v0, v1, a[34:65] blgp:7 ; encoding: [0x00,0x80,0xc0,0xd3,0x00,0x03,0x8a,0xe4]
 261 // GFX90A: :[[@LINE-2]]:{{[0-9]+}}: error: instruction not supported on this GPU
 262
 263 v_mfma_f32_32x32x1_2b_f32 v[0:31], v0, v1, v[34:65] blgp:7
 264 // GFX940: v_mfma_f32_32x32x1_2b_f32 v[0:31], v0, v1, v[34:65] blgp:7 ; encoding: [0x00,0x00,0xc0,0xd3,0x00,0x03,0x8a,0xe4]
 265 // GFX90A: :[[@LINE-2]]:{{[0-9]+}}: error: instruction not supported on this GPU
 266
 267 v_mfma_f32_32x32x1f32 a[0:31], v0, v1, a[34:65] blgp:7
 268 // GFX940: v_mfma_f32_32x32x1_2b_f32 a[0:31], v0, v1, a[34:65] blgp:7 ; encoding: [0x00,0x80,0xc0,0xd3,0x00,0x03,0x8a,0xe4]
 269
 270 v_mfma_f32_32x32x1f32 v[0:31], v0, v1, v[34:65] blgp:7
 271 // GFX940: v_mfma_f32_32x32x1_2b_f32 v[0:31], v0, v1, v[34:65] blgp:7 ; encoding: [0x00,0x00,0xc0,0xd3,0x00,0x03,0x8a,0xe4]
 272
 273 v_mfma_i32_32x32x16_i8 v[0:15], v[2:3], v[4:5], v[0:15]
 274 // 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]
 275 // GFX90A: :[[@LINE-2]]:{{[0-9]+}}: error: instruction not supported on this GPU
 276
 277 v_mfma_i32_32x32x16_i8 a[0:15], v[2:3], v[4:5], a[0:15]
 278 // 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]
 279 // GFX90A: :[[@LINE-2]]:{{[0-9]+}}: error: instruction not supported on this GPU
 280
 281 v_mfma_i32_32x32x16_i8 v[0:15], v[2:3], v[4:5], v[0:15]
 282 // 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]
 283 // GFX90A: :[[@LINE-2]]:{{[0-9]+}}: error: instruction not supported on this GPU
 284
 285 v_mfma_i32_32x32x16_i8 a[0:15], v[2:3], v[4:5], a[0:15]
 286 // 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]
 287 // GFX90A: :[[@LINE-2]]:{{[0-9]+}}: error: instruction not supported on this GPU
 288
 289 v_mfma_i32_32x32x16_i8 a[0:15], v[2:3], v[4:5], a[0:15] blgp:5
 290 // 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]
 291 // GFX90A: :[[@LINE-2]]:{{[0-9]+}}: error: instruction not supported on this GPU
 292
 293 v_mfma_i32_32x32x16i8 a[0:15], v[2:3], v[4:5], a[0:15] blgp:5
 294 // 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]
 295 // GFX90A: :[[@LINE-2]]:{{[0-9]+}}: error: instruction not supported on this GPU
 296
 297 v_mfma_i32_32x32x16i8 v[0:15], v[2:3], v[4:5], v[0:15] blgp:5
 298 // GFX940: v_mfma_i32_32x32x16_i8 v[0:15], v[2:3], v[4:5], v[0:15] blgp:5 ; encoding: [0x00,0x00,0xd6,0xd3,0x02,0x09,0x02,0xa4]
 299 // GFX90A: :[[@LINE-2]]:{{[0-9]+}}: error: instruction not supported on this GPU
 300
 301 v_mfma_i32_16x16x32_i8 v[0:3], v[2:3], v[4:5], v[0:3]
 302 // 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]
 303 // GFX90A: :[[@LINE-2]]:{{[0-9]+}}: error: instruction not supported on this GPU
 304
 305 v_mfma_i32_16x16x32_i8 a[0:3], v[2:3], v[4:5], a[0:3]
 306 // 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]
 307 // GFX90A: :[[@LINE-2]]:{{[0-9]+}}: error: instruction not supported on this GPU
 308
 309 v_mfma_i32_16x16x32_i8 a[0:3], v[2:3], v[4:5], a[0:3] blgp:5
 310 // 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]
 311 // GFX90A: :[[@LINE-2]]:{{[0-9]+}}: error: instruction not supported on this GPU
 312
 313 v_mfma_i32_16x16x32i8 v[0:3], v[2:3], v[4:5], v[0:3] blgp:5
 314 // GFX940: v_mfma_i32_16x16x32_i8 v[0:3], v[2:3], v[4:5], v[0:3] blgp:5 ; encoding: [0x00,0x00,0xd7,0xd3,0x02,0x09,0x02,0xa4]
 315 // GFX90A: :[[@LINE-2]]:{{[0-9]+}}: error: instruction not supported on this GPU
 316
 317 v_mfma_i32_16x16x32i8 a[0:3], v[2:3], v[4:5], a[0:3] blgp:5
 318 // 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]
 319 // GFX90A: :[[@LINE-2]]:{{[0-9]+}}: error: instruction not supported on this GPU
 320
 321 v_mfma_f32_32x32x4_2b_bf16 v[0:31], v[2:3], v[4:5], v[34:65]
 322 // GFX940: v_mfma_f32_32x32x4_2b_bf16 v[0:31], v[2:3], v[4:5], v[34:65] ; encoding: [0x00,0x00,0xdd,0xd3,0x02,0x09,0x8a,0x04]
 323 // GFX90A: :[[@LINE-2]]:{{[0-9]+}}: error: instruction not supported on this GPU
 324
 325 v_mfma_f32_32x32x4_2b_bf16 a[0:31], v[2:3], v[4:5], a[34:65]
 326 // GFX940: v_mfma_f32_32x32x4_2b_bf16 a[0:31], v[2:3], v[4:5], a[34:65] ; encoding: [0x00,0x80,0xdd,0xd3,0x02,0x09,0x8a,0x04]
 327 // GFX90A: :[[@LINE-2]]:{{[0-9]+}}: error: instruction not supported on this GPU
 328
 329 v_mfma_f32_32x32x4_2b_bf16 v[0:31], v[2:3], v[4:5], v[34:65]
 330 // GFX940: v_mfma_f32_32x32x4_2b_bf16 v[0:31], v[2:3], v[4:5], v[34:65] ; encoding: [0x00,0x00,0xdd,0xd3,0x02,0x09,0x8a,0x04]
 331 // GFX90A: :[[@LINE-2]]:{{[0-9]+}}: error: instruction not supported on this GPU
 332
 333 v_mfma_f32_32x32x4_2b_bf16 a[0:31], v[2:3], v[4:5], a[34:65]
 334 // GFX940: v_mfma_f32_32x32x4_2b_bf16 a[0:31], v[2:3], v[4:5], a[34:65] ; encoding: [0x00,0x80,0xdd,0xd3,0x02,0x09,0x8a,0x04]
 335 // GFX90A: :[[@LINE-2]]:{{[0-9]+}}: error: instruction not supported on this GPU
 336
 337 v_mfma_f32_32x32x4bf16 v[0:31], v[2:3], v[4:5], v[34:65] blgp:5
 338 // GFX940: v_mfma_f32_32x32x4_2b_bf16 v[0:31], v[2:3], v[4:5], v[34:65] blgp:5 ; encoding: [0x00,0x00,0xdd,0xd3,0x02,0x09,0x8a,0xa4]
 339 // GFX90A: :[[@LINE-2]]:{{[0-9]+}}: error: operands are not valid for this GPU or mode
 340
 341 v_mfma_f32_32x32x4bf16 a[0:31], v[2:3], v[4:5], a[34:65] blgp:5
 342 // GFX940: v_mfma_f32_32x32x4_2b_bf16 a[0:31], v[2:3], v[4:5], a[34:65] blgp:5 ; encoding: [0x00,0x80,0xdd,0xd3,0x02,0x09,0x8a,0xa4]
 343 // GFX90A: :[[@LINE-2]]:{{[0-9]+}}: error: operands are not valid for this GPU or mode
 344
 345 v_mfma_f32_32x32x4bf16_1k v[0:31], v[2:3], v[4:5], v[34:65] blgp:5
 346 // GFX940: v_mfma_f32_32x32x4_2b_bf16 v[0:31], v[2:3], v[4:5], v[34:65] blgp:5 ; encoding: [0x00,0x00,0xdd,0xd3,0x02,0x09,0x8a,0xa4]
 347
 348 v_mfma_f32_32x32x4bf16_1k a[0:31], v[2:3], v[4:5], a[34:65] blgp:5
 349 // GFX940: v_mfma_f32_32x32x4_2b_bf16 a[0:31], v[2:3], v[4:5], a[34:65] blgp:5 ; encoding: [0x00,0x80,0xdd,0xd3,0x02,0x09,0x8a,0xa4]
 350
 351 v_mfma_f32_16x16x4_4b_bf16 v[0:15], v[2:3], v[4:5], v[18:33]
 352 // GFX940: v_mfma_f32_16x16x4_4b_bf16 v[0:15], v[2:3], v[4:5], v[18:33] ; encoding: [0x00,0x00,0xde,0xd3,0x02,0x09,0x4a,0x04]
 353 // GFX90A: :[[@LINE-2]]:{{[0-9]+}}: error: instruction not supported on this GPU
 354
 355 v_mfma_f32_16x16x4_4b_bf16 a[0:15], v[2:3], v[4:5], a[18:33]
 356 // GFX940: v_mfma_f32_16x16x4_4b_bf16 a[0:15], v[2:3], v[4:5], a[18:33] ; encoding: [0x00,0x80,0xde,0xd3,0x02,0x09,0x4a,0x04]
 357 // GFX90A: :[[@LINE-2]]:{{[0-9]+}}: error: instruction not supported on this GPU
 358
 359 v_mfma_f32_16x16x4bf16 v[0:15], v[2:3], v[4:5], v[18:33] blgp:5
 360 // GFX940: v_mfma_f32_16x16x4_4b_bf16 v[0:15], v[2:3], v[4:5], v[18:33] blgp:5 ; encoding: [0x00,0x00,0xde,0xd3,0x02,0x09,0x4a,0xa4]
 361 // GFX90A: :[[@LINE-2]]:{{[0-9]+}}: error: instruction not supported on this GPU
 362
 363 v_mfma_f32_16x16x4bf16 a[0:15], v[2:3], v[4:5], a[18:33] blgp:5
 364 // GFX940: v_mfma_f32_16x16x4_4b_bf16 a[0:15], v[2:3], v[4:5], a[18:33] blgp:5 ; encoding: [0x00,0x80,0xde,0xd3,0x02,0x09,0x4a,0xa4]
 365 // GFX90A: :[[@LINE-2]]:{{[0-9]+}}: error: instruction not supported on this GPU
 366
 367 v_mfma_f32_16x16x4bf16_1k v[0:15], v[2:3], v[4:5], v[18:33] blgp:5
 368 // GFX940: v_mfma_f32_16x16x4_4b_bf16 v[0:15], v[2:3], v[4:5], v[18:33] blgp:5 ; encoding: [0x00,0x00,0xde,0xd3,0x02,0x09,0x4a,0xa4]
 369
 370 v_mfma_f32_16x16x4bf16_1k a[0:15], v[2:3], v[4:5], a[18:33] blgp:5
 371 // GFX940: v_mfma_f32_16x16x4_4b_bf16 a[0:15], v[2:3], v[4:5], a[18:33] blgp:5 ; encoding: [0x00,0x80,0xde,0xd3,0x02,0x09,0x4a,0xa4]
 372
 373 v_mfma_f32_4x4x4_16b_bf16 v[0:3], v[2:3], v[4:5], v[2:5]
 374 // 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]
 375 // GFX90A: :[[@LINE-2]]:{{[0-9]+}}: error: instruction not supported on this GPU
 376
 377 v_mfma_f32_4x4x4_16b_bf16 a[0:3], v[2:3], v[4:5], a[2:5]
 378 // 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]
 379 // GFX90A: :[[@LINE-2]]:{{[0-9]+}}: error: instruction not supported on this GPU
 380
 381 v_mfma_f32_4x4x4bf16 v[0:3], v[2:3], v[4:5], v[2:5]
 382 // 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]
 383 // GFX90A: :[[@LINE-2]]:{{[0-9]+}}: error: instruction not supported on this GPU
 384
 385 v_mfma_f32_4x4x4bf16 a[0:3], v[2:3], v[4:5], a[2:5]
 386 // 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]
 387 // GFX90A: :[[@LINE-2]]:{{[0-9]+}}: error: instruction not supported on this GPU
 388
 389 v_mfma_f32_4x4x4bf16_1k v[0:3], v[2:3], v[4:5], v[2:5]
 390 // 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]
 391
 392 v_mfma_f32_4x4x4bf16_1k a[0:3], v[2:3], v[4:5], a[2:5]
 393 // 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]
 394
 395 v_mfma_f32_32x32x8_bf16 v[0:15], v[2:3], v[4:5], v[18:33]
 396 // GFX940: v_mfma_f32_32x32x8_bf16 v[0:15], v[2:3], v[4:5], v[18:33] ; encoding: [0x00,0x00,0xe0,0xd3,0x02,0x09,0x4a,0x04]
 397 // GFX90A: :[[@LINE-2]]:{{[0-9]+}}: error: instruction not supported on this GPU
 398
 399 v_mfma_f32_32x32x8_bf16 a[0:15], v[2:3], v[4:5], a[18:33]
 400 // GFX940: v_mfma_f32_32x32x8_bf16 a[0:15], v[2:3], v[4:5], a[18:33] ; encoding: [0x00,0x80,0xe0,0xd3,0x02,0x09,0x4a,0x04]
 401 // GFX90A: :[[@LINE-2]]:{{[0-9]+}}: error: instruction not supported on this GPU
 402
 403 v_mfma_f32_32x32x8bf16 v[0:15], v[2:3], v[4:5], v[18:33]
 404 // GFX940: v_mfma_f32_32x32x8_bf16 v[0:15], v[2:3], v[4:5], v[18:33] ; encoding: [0x00,0x00,0xe0,0xd3,0x02,0x09,0x4a,0x04]
 405 // GFX90A: :[[@LINE-2]]:{{[0-9]+}}: error: instruction not supported on this GPU
 406
 407 v_mfma_f32_32x32x8bf16 a[0:15], v[2:3], v[4:5], a[18:33]
 408 // GFX940: v_mfma_f32_32x32x8_bf16 a[0:15], v[2:3], v[4:5], a[18:33] ; encoding: [0x00,0x80,0xe0,0xd3,0x02,0x09,0x4a,0x04]
 409 // GFX90A: :[[@LINE-2]]:{{[0-9]+}}: error: instruction not supported on this GPU
 410
 411 v_mfma_f32_32x32x8bf16_1k v[0:15], v[2:3], v[4:5], v[18:33]
 412 // GFX940: v_mfma_f32_32x32x8_bf16 v[0:15], v[2:3], v[4:5], v[18:33] ; encoding: [0x00,0x00,0xe0,0xd3,0x02,0x09,0x4a,0x04]
 413
 414 v_mfma_f32_32x32x8bf16_1k a[0:15], v[2:3], v[4:5], a[18:33]
 415 // GFX940: v_mfma_f32_32x32x8_bf16 a[0:15], v[2:3], v[4:5], a[18:33] ; encoding: [0x00,0x80,0xe0,0xd3,0x02,0x09,0x4a,0x04]
 416
 417 v_mfma_f32_16x16x16_bf16 v[0:3], v[2:3], v[4:5], v[2:5]
 418 // 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]
 419 // GFX90A: :[[@LINE-2]]:{{[0-9]+}}: error: instruction not supported on this GPU
 420
 421 v_mfma_f32_16x16x16_bf16 a[0:3], v[2:3], v[4:5], a[2:5]
 422 // 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]
 423 // GFX90A: :[[@LINE-2]]:{{[0-9]+}}: error: instruction not supported on this GPU
 424
 425 v_mfma_f32_16x16x16bf16 v[0:3], v[2:3], v[4:5], v[2:5]
 426 // 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]
 427 // GFX90A: :[[@LINE-2]]:{{[0-9]+}}: error: instruction not supported on this GPU
 428
 429 v_mfma_f32_16x16x16bf16 a[0:3], v[2:3], v[4:5], a[2:5]
 430 // 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]
 431 // GFX90A: :[[@LINE-2]]:{{[0-9]+}}: error: instruction not supported on this GPU
 432
 433 v_mfma_f32_16x16x16bf16_1k v[0:3], v[2:3], v[4:5], v[2:5]
 434 // 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]
 435
 436 v_mfma_f32_16x16x16bf16_1k a[0:3], v[2:3], v[4:5], a[2:5]
 437 // 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]
 438
 439 v_mfma_f32_16x16x8_xf32 a[0:3], v[2:3], v[4:5], a[2:5]
 440 // 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]
 441 // GFX90A: :[[@LINE-2]]:{{[0-9]+}}: error: instruction not supported on this GPU
 442
 443 v_mfma_f32_16x16x8_xf32 v[0:3], v[2:3], v[4:5], v[2:5]
 444 // 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]
 445 // GFX90A: :[[@LINE-2]]:{{[0-9]+}}: error: instruction not supported on this GPU
 446
 447 v_mfma_f32_16x16x8xf32 a[0:3], v[2:3], v[4:5], a[2:5]
 448 // 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]
 449 // GFX90A: :[[@LINE-2]]:{{[0-9]+}}: error: instruction not supported on this GPU
 450
 451 v_mfma_f32_16x16x8xf32 v[0:3], v[2:3], v[4:5], v[2:5]
 452 // 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]
 453 // GFX90A: :[[@LINE-2]]:{{[0-9]+}}: error: instruction not supported on this GPU
 454
 455 v_mfma_f32_32x32x4_xf32 v[0:15], v[2:3], v[4:5], v[18:33]
 456 // GFX940: v_mfma_f32_32x32x4_xf32 v[0:15], v[2:3], v[4:5], v[18:33] ; encoding: [0x00,0x00,0xbf,0xd3,0x02,0x09,0x4a,0x04]
 457 // GFX90A: :[[@LINE-2]]:{{[0-9]+}}: error: instruction not supported on this GPU
 458
 459 v_mfma_f32_32x32x4_xf32 a[0:15], v[2:3], v[4:5], a[18:33]
 460 // GFX940: v_mfma_f32_32x32x4_xf32 a[0:15], v[2:3], v[4:5], a[18:33] ; encoding: [0x00,0x80,0xbf,0xd3,0x02,0x09,0x4a,0x04]
 461 // GFX90A: :[[@LINE-2]]:{{[0-9]+}}: error: instruction not supported on this GPU
 462
 463 v_mfma_f32_32x32x4xf32 v[0:15], v[2:3], v[4:5], v[18:33]
 464 // GFX940: v_mfma_f32_32x32x4_xf32 v[0:15], v[2:3], v[4:5], v[18:33] ; encoding: [0x00,0x00,0xbf,0xd3,0x02,0x09,0x4a,0x04]
 465 // GFX90A: :[[@LINE-2]]:{{[0-9]+}}: error: instruction not supported on this GPU
 466
 467 v_mfma_f32_32x32x4xf32 a[0:15], v[2:3], v[4:5], a[18:33]
 468 // GFX940: v_mfma_f32_32x32x4_xf32 a[0:15], v[2:3], v[4:5], a[18:33] ; encoding: [0x00,0x80,0xbf,0xd3,0x02,0x09,0x4a,0x04]
 469 // GFX90A: :[[@LINE-2]]:{{[0-9]+}}: error: instruction not supported on this GPU
 470
 471 v_mfma_f32_16x16x32_bf8_bf8 v[0:3], v[2:3], v[4:5], v[0:3]
 472 // 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]
 473 // GFX90A: :[[@LINE-2]]:{{[0-9]+}}: error: instruction not supported on this GPU
 474
 475 v_mfma_f32_16x16x32_bf8_bf8 a[0:3], v[2:3], v[4:5], a[0:3]
 476 // 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]
 477 // GFX90A: :[[@LINE-2]]:{{[0-9]+}}: error: instruction not supported on this GPU
 478
 479 v_mfma_f32_16x16x32_bf8_bf8 a[0:3], v[2:3], v[4:5], a[0:3] blgp:5
 480 // 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]
 481 // GFX90A: :[[@LINE-2]]:{{[0-9]+}}: error: instruction not supported on this GPU
 482
 483 v_mfma_f32_16x16x32_bf8_fp8 v[0:3], v[2:3], v[4:5], v[0:3]
 484 // 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]
 485 // GFX90A: :[[@LINE-2]]:{{[0-9]+}}: error: instruction not supported on this GPU
 486
 487 v_mfma_f32_16x16x32_bf8_fp8 a[0:3], v[2:3], v[4:5], a[0:3]
 488 // 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]
 489 // GFX90A: :[[@LINE-2]]:{{[0-9]+}}: error: instruction not supported on this GPU
 490
 491 v_mfma_f32_16x16x32_bf8_fp8 a[0:3], v[2:3], v[4:5], a[0:3] blgp:5
 492 // 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]
 493 // GFX90A: :[[@LINE-2]]:{{[0-9]+}}: error: instruction not supported on this GPU
 494
 495 v_mfma_f32_16x16x32_fp8_bf8 v[0:3], v[2:3], v[4:5], v[0:3]
 496 // 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]
 497 // GFX90A: :[[@LINE-2]]:{{[0-9]+}}: error: instruction not supported on this GPU
 498
 499 v_mfma_f32_16x16x32_fp8_bf8 a[0:3], v[2:3], v[4:5], a[0:3]
 500 // 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]
 501 // GFX90A: :[[@LINE-2]]:{{[0-9]+}}: error: instruction not supported on this GPU
 502
 503 v_mfma_f32_16x16x32_fp8_bf8 a[0:3], v[2:3], v[4:5], a[0:3] blgp:5
 504 // 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]
 505 // GFX90A: :[[@LINE-2]]:{{[0-9]+}}: error: instruction not supported on this GPU
 506
 507 v_mfma_f32_16x16x32_fp8_fp8 v[0:3], v[2:3], v[4:5], v[0:3]
 508 // 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]
 509 // GFX90A: :[[@LINE-2]]:{{[0-9]+}}: error: instruction not supported on this GPU
 510
 511 v_mfma_f32_16x16x32_fp8_fp8 a[0:3], v[2:3], v[4:5], a[0:3]
 512 // 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]
 513 // GFX90A: :[[@LINE-2]]:{{[0-9]+}}: error: instruction not supported on this GPU
 514
 515 v_mfma_f32_16x16x32_fp8_fp8 a[0:3], v[2:3], v[4:5], a[0:3] blgp:5
 516 // 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]
 517 // GFX90A: :[[@LINE-2]]:{{[0-9]+}}: error: instruction not supported on this GPU
 518
 519 v_mfma_f32_32x32x16_bf8_bf8 v[0:15], v[2:3], v[4:5], v[0:15]
 520 // 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]
 521 // GFX90A: :[[@LINE-2]]:{{[0-9]+}}: error: instruction not supported on this GPU
 522
 523 v_mfma_f32_32x32x16_bf8_bf8 a[0:15], v[2:3], v[4:5], a[0:15]
 524 // 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]
 525 // GFX90A: :[[@LINE-2]]:{{[0-9]+}}: error: instruction not supported on this GPU
 526
 527 v_mfma_f32_32x32x16_bf8_bf8 a[0:15], v[2:3], v[4:5], a[0:15] blgp:5
 528 // 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]
 529 // GFX90A: :[[@LINE-2]]:{{[0-9]+}}: error: instruction not supported on this GPU
 530
 531 v_mfma_f32_32x32x16_bf8_fp8 v[0:15], v[2:3], v[4:5], v[0:15]
 532 // 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]
 533 // GFX90A: :[[@LINE-2]]:{{[0-9]+}}: error: instruction not supported on this GPU
 534
 535 v_mfma_f32_32x32x16_bf8_fp8 a[0:15], v[2:3], v[4:5], a[0:15]
 536 // 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]
 537 // GFX90A: :[[@LINE-2]]:{{[0-9]+}}: error: instruction not supported on this GPU
 538
 539 v_mfma_f32_32x32x16_bf8_fp8 a[0:15], v[2:3], v[4:5], a[0:15] blgp:5
 540 // 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]
 541 // GFX90A: :[[@LINE-2]]:{{[0-9]+}}: error: instruction not supported on this GPU
 542
 543 v_mfma_f32_32x32x16_fp8_bf8 v[0:15], v[2:3], v[4:5], v[0:15]
 544 // 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]
 545 // GFX90A: :[[@LINE-2]]:{{[0-9]+}}: error: instruction not supported on this GPU
 546
 547 v_mfma_f32_32x32x16_fp8_bf8 a[0:15], v[2:3], v[4:5], a[0:15]
 548 // 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]
 549 // GFX90A: :[[@LINE-2]]:{{[0-9]+}}: error: instruction not supported on this GPU
 550
 551 v_mfma_f32_32x32x16_fp8_bf8 a[0:15], v[2:3], v[4:5], a[0:15] blgp:5
 552 // 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]
 553 // GFX90A: :[[@LINE-2]]:{{[0-9]+}}: error: instruction not supported on this GPU
 554
 555 v_mfma_f32_32x32x16_fp8_fp8 v[0:15], v[2:3], v[4:5], v[0:15]
 556 // 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]
 557 // GFX90A: :[[@LINE-2]]:{{[0-9]+}}: error: instruction not supported on this GPU
 558
 559 v_mfma_f32_32x32x16_fp8_fp8 a[0:15], v[2:3], v[4:5], a[0:15]
 560 // 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]
 561 // GFX90A: :[[@LINE-2]]:{{[0-9]+}}: error: instruction not supported on this GPU
 562
 563 v_mfma_f32_32x32x16_fp8_fp8 a[0:15], v[2:3], v[4:5], a[0:15] blgp:5
 564 // 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]
 565 // GFX90A: :[[@LINE-2]]:{{[0-9]+}}: error: instruction not supported on this GPU
 566
 567 //===----------------------------------------------------------------------===//
 568 // SMFMAC opcodes.
 569 //===----------------------------------------------------------------------===//
 570
 571 v_smfmac_f32_16x16x32_f16 v[10:13], a[2:3], v[4:7], v0 cbsz:3 abid:1
 572 // 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]
 573 // GFX90A: :[[@LINE-2]]:{{[0-9]+}}: error: instruction not supported on this GPU
 574
 575 v_smfmac_f32_16x16x32_f16 a[10:13], v[2:3], a[4:7], v1
 576 // 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]
 577 // GFX90A: :[[@LINE-2]]:{{[0-9]+}}: error: instruction not supported on this GPU
 578
 579 v_smfmac_f32_32x32x16_f16 v[10:25], a[2:3], v[4:7], v2 cbsz:3 abid:1
 580 // 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]
 581 // GFX90A: :[[@LINE-2]]:{{[0-9]+}}: error: instruction not supported on this GPU
 582
 583 v_smfmac_f32_32x32x16_f16 a[10:25], v[2:3], a[4:7], v3
 584 // 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]
 585 // GFX90A: :[[@LINE-2]]:{{[0-9]+}}: error: instruction not supported on this GPU
 586
 587 v_smfmac_f32_16x16x32_bf16 v[10:13], a[2:3], v[4:7], v4 cbsz:3 abid:1
 588 // 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]
 589 // GFX90A: :[[@LINE-2]]:{{[0-9]+}}: error: instruction not supported on this GPU
 590
 591 v_smfmac_f32_16x16x32_bf16 a[10:13], v[2:3], a[4:7], v5
 592 // 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]
 593 // GFX90A: :[[@LINE-2]]:{{[0-9]+}}: error: instruction not supported on this GPU
 594
 595 v_smfmac_f32_32x32x16_bf16 v[10:25], a[2:3], v[4:7], v6 cbsz:3 abid:1
 596 // 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]
 597 // GFX90A: :[[@LINE-2]]:{{[0-9]+}}: error: instruction not supported on this GPU
 598
 599 v_smfmac_f32_32x32x16_bf16 a[10:25], v[2:3], a[4:7], v7
 600 // 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]
 601 // GFX90A: :[[@LINE-2]]:{{[0-9]+}}: error: instruction not supported on this GPU
 602
 603 v_smfmac_i32_16x16x64_i8 v[10:13], a[2:3], v[4:7], v8 cbsz:3 abid:1
 604 // GFX940: v_smfmac_i32_16x16x64_i8 v[10:13], a[2:3], v[4:7], v8 cbsz:3 abid:1 ; encoding: [0x0a,0x0b,0xea,0xd3,0x02,0x09,0x22,0x0c]
 605 // GFX90A: :[[@LINE-2]]:{{[0-9]+}}: error: instruction not supported on this GPU
 606
 607 v_smfmac_i32_16x16x64_i8 a[10:13], v[2:3], a[4:7], v9
 608 // GFX940: v_smfmac_i32_16x16x64_i8 a[10:13], v[2:3], a[4:7], v9 ; encoding: [0x0a,0x80,0xea,0xd3,0x02,0x09,0x26,0x14]
 609 // GFX90A: :[[@LINE-2]]:{{[0-9]+}}: error: instruction not supported on this GPU
 610
 611 v_smfmac_i32_32x32x32_i8 v[10:25], a[2:3], v[4:7], v10 cbsz:3 abid:1
 612 // GFX940: v_smfmac_i32_32x32x32_i8 v[10:25], a[2:3], v[4:7], v10 cbsz:3 abid:1 ; encoding: [0x0a,0x0b,0xec,0xd3,0x02,0x09,0x2a,0x0c]
 613 // GFX90A: :[[@LINE-2]]:{{[0-9]+}}: error: instruction not supported on this GPU
 614
 615 v_smfmac_i32_32x32x32_i8 a[10:25], v[2:3], a[4:7], v11
 616 // GFX940: v_smfmac_i32_32x32x32_i8 a[10:25], v[2:3], a[4:7], v11 ; encoding: [0x0a,0x80,0xec,0xd3,0x02,0x09,0x2e,0x14]
 617 // GFX90A: :[[@LINE-2]]:{{[0-9]+}}: error: instruction not supported on this GPU
 618
 619 v_smfmac_f32_16x16x64_bf8_bf8 v[0:3], a[2:3], v[4:7], v1 cbsz:3 abid:1
 620 // 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]
 621 // GFX90A: :[[@LINE-2]]:{{[0-9]+}}: error: instruction not supported on this GPU
 622
 623 v_smfmac_f32_16x16x64_bf8_bf8 a[0:3], v[2:3], a[4:7], v1
 624 // 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]
 625 // GFX90A: :[[@LINE-2]]:{{[0-9]+}}: error: instruction not supported on this GPU
 626
 627 v_smfmac_f32_16x16x64_bf8_fp8 v[0:3], a[2:3], v[4:7], v1 cbsz:3 abid:1
 628 // 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]
 629 // GFX90A: :[[@LINE-2]]:{{[0-9]+}}: error: instruction not supported on this GPU
 630
 631 v_smfmac_f32_16x16x64_bf8_fp8 a[0:3], v[2:3], a[4:7], v1
 632 // 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]
 633 // GFX90A: :[[@LINE-2]]:{{[0-9]+}}: error: instruction not supported on this GPU
 634
 635 v_smfmac_f32_16x16x64_fp8_bf8 v[0:3], a[2:3], v[4:7], v1 cbsz:3 abid:1
 636 // 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]
 637 // GFX90A: :[[@LINE-2]]:{{[0-9]+}}: error: instruction not supported on this GPU
 638
 639 v_smfmac_f32_16x16x64_fp8_bf8 a[0:3], v[2:3], a[4:7], v1
 640 // 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]
 641 // GFX90A: :[[@LINE-2]]:{{[0-9]+}}: error: instruction not supported on this GPU
 642
 643 v_smfmac_f32_16x16x64_fp8_fp8 v[0:3], a[2:3], v[4:7], v1 cbsz:3 abid:1
 644 // 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]
 645 // GFX90A: :[[@LINE-2]]:{{[0-9]+}}: error: instruction not supported on this GPU
 646
 647 v_smfmac_f32_16x16x64_fp8_fp8 a[0:3], v[2:3], a[4:7], v1
 648 // 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]
 649 // GFX90A: :[[@LINE-2]]:{{[0-9]+}}: error: instruction not supported on this GPU
 650
 651 v_smfmac_f32_32x32x32_bf8_bf8 v[0:15], a[2:3], v[4:7], v1 cbsz:3 abid:1
 652 // 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]
 653 // GFX90A: :[[@LINE-2]]:{{[0-9]+}}: error: instruction not supported on this GPU
 654
 655 v_smfmac_f32_32x32x32_bf8_bf8 a[0:15], v[2:3], a[4:7], v1
 656 // 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]
 657 // GFX90A: :[[@LINE-2]]:{{[0-9]+}}: error: instruction not supported on this GPU
 658
 659 v_smfmac_f32_32x32x32_bf8_fp8 v[0:15], a[2:3], v[4:7], v1 cbsz:3 abid:1
 660 // 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]
 661 // GFX90A: :[[@LINE-2]]:{{[0-9]+}}: error: instruction not supported on this GPU
 662
 663 v_smfmac_f32_32x32x32_bf8_fp8 a[0:15], v[2:3], a[4:7], v1
 664 // 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]
 665 // GFX90A: :[[@LINE-2]]:{{[0-9]+}}: error: instruction not supported on this GPU
 666
 667 v_smfmac_f32_32x32x32_fp8_bf8 v[0:15], a[2:3], v[4:7], v1 cbsz:3 abid:1
 668 // 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]
 669 // GFX90A: :[[@LINE-2]]:{{[0-9]+}}: error: instruction not supported on this GPU
 670
 671 v_smfmac_f32_32x32x32_fp8_bf8 a[0:15], v[2:3], a[4:7], v1
 672 // 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]
 673 // GFX90A: :[[@LINE-2]]:{{[0-9]+}}: error: instruction not supported on this GPU
 674
 675 v_smfmac_f32_32x32x32_fp8_fp8 v[0:15], a[2:3], v[4:7], v1 cbsz:3 abid:1
 676 // 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]
 677 // GFX90A: :[[@LINE-2]]:{{[0-9]+}}: error: instruction not supported on this GPU
 678
 679 v_smfmac_f32_32x32x32_fp8_fp8 a[0:15], v[2:3], a[4:7], v1
 680 // 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]
 681 // GFX90A: :[[@LINE-2]]:{{[0-9]+}}: error: instruction not supported on this GPU
 682
 683 //===----------------------------------------------------------------------===//
 684 // SMFMAC aliases.
 685 //===----------------------------------------------------------------------===//
 686
 687 v_smfmac_f32_16x16x32f16 v[10:13], a[2:3], v[4:7], v0 cbsz:3 abid:1
 688 // 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]
 689 // GFX90A: :[[@LINE-2]]:{{[0-9]+}}: error: instruction not supported on this GPU
 690
 691 v_smfmac_f32_32x32x16f16 v[10:25], a[2:3], v[4:7], v2 cbsz:3 abid:1
 692 // 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]
 693 // GFX90A: :[[@LINE-2]]:{{[0-9]+}}: error: instruction not supported on this GPU
 694
 695 v_smfmac_f32_16x16x32bf16 v[10:13], a[2:3], v[4:7], v4 cbsz:3 abid:1
 696 // 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]
 697 // GFX90A: :[[@LINE-2]]:{{[0-9]+}}: error: instruction not supported on this GPU
 698
 699 v_smfmac_f32_32x32x16bf16 v[10:25], a[2:3], v[4:7], v6 cbsz:3 abid:1
 700 // 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]
 701 // GFX90A: :[[@LINE-2]]:{{[0-9]+}}: error: instruction not supported on this GPU
 702
 703 v_smfmac_i32_16x16x64i8 v[10:13], a[2:3], v[4:7], v8 cbsz:3 abid:1
 704 // GFX940: v_smfmac_i32_16x16x64_i8 v[10:13], a[2:3], v[4:7], v8 cbsz:3 abid:1 ; encoding: [0x0a,0x0b,0xea,0xd3,0x02,0x09,0x22,0x0c]
 705 // GFX90A: :[[@LINE-2]]:{{[0-9]+}}: error: instruction not supported on this GPU
 706
 707 v_smfmac_i32_32x32x32i8 a[10:25], v[2:3], a[4:7], v11
 708 // GFX940: v_smfmac_i32_32x32x32_i8 a[10:25], v[2:3], a[4:7], v11 ; encoding: [0x0a,0x80,0xec,0xd3,0x02,0x09,0x2e,0x14]
 709 // GFX90A: :[[@LINE-2]]:{{[0-9]+}}: error: instruction not supported on this GPU
 710
 711 v_smfmac_f32_16x16x64bf8bf8 v[0:3], a[2:3], v[4:7], v1 cbsz:3 abid:1
 712 // 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]
 713 // GFX90A: :[[@LINE-2]]:{{[0-9]+}}: error: instruction not supported on this GPU
 714
 715 v_smfmac_f32_16x16x64bf8fp8 v[0:3], a[2:3], v[4:7], v1 cbsz:3 abid:1
 716 // 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]
 717 // GFX90A: :[[@LINE-2]]:{{[0-9]+}}: error: instruction not supported on this GPU
 718
 719 v_smfmac_f32_16x16x64fp8bf8 v[0:3], a[2:3], v[4:7], v1 cbsz:3 abid:1
 720 // 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]
 721 // GFX90A: :[[@LINE-2]]:{{[0-9]+}}: error: instruction not supported on this GPU
 722
 723 v_smfmac_f32_16x16x64fp8fp8 v[0:3], a[2:3], v[4:7], v1 cbsz:3 abid:1
 724 // 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]
 725 // GFX90A: :[[@LINE-2]]:{{[0-9]+}}: error: instruction not supported on this GPU
 726
 727 v_smfmac_f32_32x32x32bf8bf8 v[0:15], a[2:3], v[4:7], v1 cbsz:3 abid:1
 728 // 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]
 729 // GFX90A: :[[@LINE-2]]:{{[0-9]+}}: error: instruction not supported on this GPU
 730
 731 v_smfmac_f32_32x32x32bf8fp8 v[0:15], a[2:3], v[4:7], v1 cbsz:3 abid:1
 732 // 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]
 733 // GFX90A: :[[@LINE-2]]:{{[0-9]+}}: error: instruction not supported on this GPU
 734
 735 v_smfmac_f32_32x32x32fp8bf8 v[0:15], a[2:3], v[4:7], v1 cbsz:3 abid:1
 736 // 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]
 737 // GFX90A: :[[@LINE-2]]:{{[0-9]+}}: error: instruction not supported on this GPU
 738
 739 v_smfmac_f32_32x32x32fp8fp8 v[0:15], a[2:3], v[4:7], v1 cbsz:3 abid:1
 740 // 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]
 741 // GFX90A: :[[@LINE-2]]:{{[0-9]+}}: error: instruction not supported on this GPU