1629 lines · plain
1// RUN: llvm-mc -triple=amdgcn -mcpu=gfx950 -show-encoding %s | FileCheck -check-prefix=GFX950 %s2// RUN: not llvm-mc -triple=amdgcn -mcpu=gfx942 %s 2>&1 | FileCheck -check-prefix=ERR %s3// RUN: not llvm-mc -triple=amdgcn -mcpu=gfx90a %s 2>&1 | FileCheck -check-prefix=ERR %s4// RUN: not llvm-mc -triple=amdgcn -mcpu=gfx1010 %s 2>&1 | FileCheck -check-prefix=ERR %s5 6// cbsz = SrcA7// blgp = SrcB8 9// 0 = fp8. 1 = bf810// 2 = fp6, 3 = bf611// 4 = fp412 13//===----------------------------------------------------------------------===//14// MFMA opcodes.15//===----------------------------------------------------------------------===//16 17 18//===----------------------------------------------------------------------===//19// v_mfma_f32_16x16x32_f1620//===----------------------------------------------------------------------===//21 22// GFX950: v_mfma_f32_16x16x32_f16 v[0:3], v[0:3], v[0:3], v[0:3] ; encoding: [0x00,0x00,0xd4,0xd3,0x00,0x01,0x02,0x04]23// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU24v_mfma_f32_16x16x32_f16 v[0:3], v[0:3], v[0:3], v[0:3]25 26// GFX950: v_mfma_f32_16x16x32_f16 v[0:3], v[0:3], v[0:3], v[0:3] ; encoding: [0x00,0x00,0xd4,0xd3,0x00,0x01,0x02,0x04]27// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU28v_mfma_f32_16x16x32f16 v[0:3], v[0:3], v[0:3], v[0:3]29 30// GFX950: v_mfma_f32_16x16x32_f16 a[0:3], a[0:3], a[0:3], a[0:3] ; encoding: [0x00,0x80,0xd4,0xd3,0x00,0x01,0x02,0x1c]31// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU32v_mfma_f32_16x16x32_f16 a[0:3], a[0:3], a[0:3], a[0:3]33 34// GFX950: v_mfma_f32_16x16x32_f16 a[0:3], a[0:3], a[0:3], a[0:3] ; encoding: [0x00,0x80,0xd4,0xd3,0x00,0x01,0x02,0x1c]35// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU36v_mfma_f32_16x16x32f16 a[0:3], a[0:3], a[0:3], a[0:3]37 38// GFX950: v_mfma_f32_16x16x32_f16 v[0:3], a[0:3], v[0:3], 1.0 ; encoding: [0x00,0x00,0xd4,0xd3,0x00,0x01,0xca,0x0b]39// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU40v_mfma_f32_16x16x32_f16 v[0:3], a[0:3], v[0:3], 1.041 42// GFX950: v_mfma_f32_16x16x32_f16 a[0:3], v[0:3], a[0:3], 1.0 ; encoding: [0x00,0x80,0xd4,0xd3,0x00,0x01,0xca,0x13]43// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU44v_mfma_f32_16x16x32_f16 a[0:3], v[0:3], a[0:3], 1.045 46// GFX950: v_mfma_f32_16x16x32_f16 v[0:3], v[0:3], v[0:3], v[0:3] blgp:5 ; encoding: [0x00,0x00,0xd4,0xd3,0x00,0x01,0x02,0xa4]47// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU48v_mfma_f32_16x16x32_f16 v[0:3], v[0:3], v[0:3], v[0:3] blgp:549 50// GFX950: v_mfma_f32_16x16x32_f16 a[0:3], a[0:3], a[0:3], a[0:3] blgp:1 ; encoding: [0x00,0x80,0xd4,0xd3,0x00,0x01,0x02,0x3c]51// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU52v_mfma_f32_16x16x32_f16 a[0:3], a[0:3], a[0:3], a[0:3] blgp:153 54// GFX950: v_mfma_f32_16x16x32_f16 a[0:3], a[0:3], a[0:3], a[0:3] cbsz:3 ; encoding: [0x00,0x83,0xd4,0xd3,0x00,0x01,0x02,0x1c]55// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU56v_mfma_f32_16x16x32_f16 a[0:3], a[0:3], a[0:3], a[0:3] cbsz:357 58// GFX950: v_mfma_f32_16x16x32_f16 a[0:3], a[0:3], a[0:3], a[0:3] abid:1 ; encoding: [0x00,0x88,0xd4,0xd3,0x00,0x01,0x02,0x1c]59// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU60v_mfma_f32_16x16x32_f16 a[0:3], a[0:3], a[0:3], a[0:3] abid:161 62// GFX950: v_mfma_f32_16x16x32_f16 a[0:3], a[0:3], a[0:3], a[0:3] cbsz:3 abid:1 ; encoding: [0x00,0x8b,0xd4,0xd3,0x00,0x01,0x02,0x1c]63// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU64v_mfma_f32_16x16x32_f16 a[0:3], a[0:3], a[0:3], a[0:3] cbsz:3 abid:165 66// GFX950: v_mfma_f32_16x16x32_f16 a[0:3], v[0:3], v[0:3], a[4:7] ; encoding: [0x00,0x80,0xd4,0xd3,0x00,0x01,0x12,0x04]67// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU68v_mfma_f32_16x16x32_f16 a[0:3], v[0:3], v[0:3], a[4:7]69 70// GFX950: v_mfma_f32_16x16x32_f16 v[0:3], a[0:3], a[0:3], v[4:7] ; encoding: [0x00,0x00,0xd4,0xd3,0x00,0x01,0x12,0x1c]71// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU72v_mfma_f32_16x16x32_f16 v[0:3], a[0:3], a[0:3], v[4:7]73 74//===----------------------------------------------------------------------===//75// v_mfma_f32_32x32x16_f1676//===----------------------------------------------------------------------===//77 78// GFX950: v_mfma_f32_32x32x16_f16 v[0:15], v[0:3], v[0:3], v[0:15] ; encoding: [0x00,0x00,0xd5,0xd3,0x00,0x01,0x02,0x04]79// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU80v_mfma_f32_32x32x16_f16 v[0:15], v[0:3], v[0:3], v[0:15]81 82// GFX950: v_mfma_f32_32x32x16_f16 a[0:15], a[0:3], a[0:3], a[0:15] ; encoding: [0x00,0x80,0xd5,0xd3,0x00,0x01,0x02,0x1c]83// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU84v_mfma_f32_32x32x16_f16 a[0:15], a[0:3], a[0:3], a[0:15]85 86// GFX950: v_mfma_f32_32x32x16_f16 v[0:15], v[0:3], v[0:3], v[0:15] ; encoding: [0x00,0x00,0xd5,0xd3,0x00,0x01,0x02,0x04]87// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU88v_mfma_f32_32x32x16f16 v[0:15], v[0:3], v[0:3], v[0:15]89 90// GFX950: v_mfma_f32_32x32x16_f16 a[0:15], a[0:3], a[0:3], a[0:15] ; encoding: [0x00,0x80,0xd5,0xd3,0x00,0x01,0x02,0x1c]91// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU92v_mfma_f32_32x32x16f16 a[0:15], a[0:3], a[0:3], a[0:15]93 94// GFX950: v_mfma_f32_32x32x16_f16 v[0:15], v[0:3], v[0:3], 1.0 ; encoding: [0x00,0x00,0xd5,0xd3,0x00,0x01,0xca,0x03]95// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU96v_mfma_f32_32x32x16_f16 v[0:15], v[0:3], v[0:3], 1.097 98// GFX950: v_mfma_f32_32x32x16_f16 a[0:15], a[0:3], a[0:3], 1.0 ; encoding: [0x00,0x80,0xd5,0xd3,0x00,0x01,0xca,0x1b]99// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU100v_mfma_f32_32x32x16_f16 a[0:15], a[0:3], a[0:3], 1.0101 102// GFX950: v_mfma_f32_32x32x16_f16 v[0:15], v[0:3], v[0:3], v[0:15] blgp:5 ; encoding: [0x00,0x00,0xd5,0xd3,0x00,0x01,0x02,0xa4]103// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU104v_mfma_f32_32x32x16_f16 v[0:15], v[0:3], v[0:3], v[0:15] blgp:5105 106// GFX950: v_mfma_f32_32x32x16_f16 a[0:15], a[0:3], a[0:3], a[0:15] blgp:2 ; encoding: [0x00,0x80,0xd5,0xd3,0x00,0x01,0x02,0x5c]107// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU108v_mfma_f32_32x32x16_f16 a[0:15], a[0:3], a[0:3], a[0:15] blgp:2109 110// GFX950: v_mfma_f32_32x32x16_f16 v[0:15], v[0:3], v[0:3], v[0:15] cbsz:3 ; encoding: [0x00,0x03,0xd5,0xd3,0x00,0x01,0x02,0x04]111// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU112v_mfma_f32_32x32x16_f16 v[0:15], v[0:3], v[0:3], v[0:15] cbsz:3113 114// GFX950: v_mfma_f32_32x32x16_f16 v[0:15], v[0:3], v[0:3], v[0:15] abid:1 ; encoding: [0x00,0x08,0xd5,0xd3,0x00,0x01,0x02,0x04]115// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU116v_mfma_f32_32x32x16_f16 v[0:15], v[0:3], v[0:3], v[0:15] abid:1117 118// GFX950: v_mfma_f32_32x32x16_f16 a[0:15], a[0:3], a[0:3], a[0:15] cbsz:3 abid:1 ; encoding: [0x00,0x8b,0xd5,0xd3,0x00,0x01,0x02,0x1c]119// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU120v_mfma_f32_32x32x16_f16 a[0:15], a[0:3], a[0:3], a[0:15] cbsz:3 abid:1121 122//===----------------------------------------------------------------------===//123// v_mfma_f32_32x32x16_bf16124//===----------------------------------------------------------------------===//125 126// GFX950: v_mfma_f32_32x32x16_bf16 v[0:15], v[0:3], v[0:3], v[0:15] ; encoding: [0x00,0x00,0xb7,0xd3,0x00,0x01,0x02,0x04]127// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU128v_mfma_f32_32x32x16_bf16 v[0:15], v[0:3], v[0:3], v[0:15]129 130// GFX950: v_mfma_f32_32x32x16_bf16 a[0:15], a[0:3], a[0:3], a[0:15] ; encoding: [0x00,0x80,0xb7,0xd3,0x00,0x01,0x02,0x1c]131// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU132v_mfma_f32_32x32x16_bf16 a[0:15], a[0:3], a[0:3], a[0:15]133 134// GFX950: v_mfma_f32_32x32x16_bf16 v[0:15], v[0:3], v[0:3], v[0:15] ; encoding: [0x00,0x00,0xb7,0xd3,0x00,0x01,0x02,0x04]135// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU136v_mfma_f32_32x32x16bf16 v[0:15], v[0:3], v[0:3], v[0:15]137 138// GFX950: v_mfma_f32_32x32x16_bf16 a[0:15], a[0:3], a[0:3], a[0:15] ; encoding: [0x00,0x80,0xb7,0xd3,0x00,0x01,0x02,0x1c]139// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU140v_mfma_f32_32x32x16bf16 a[0:15], a[0:3], a[0:3], a[0:15]141 142// GFX950: v_mfma_f32_32x32x16_bf16 v[0:15], v[0:3], v[0:3], 1.0 ; encoding: [0x00,0x00,0xb7,0xd3,0x00,0x01,0xca,0x03]143// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU144v_mfma_f32_32x32x16_bf16 v[0:15], v[0:3], v[0:3], 1.0145 146// GFX950: v_mfma_f32_32x32x16_bf16 a[0:15], a[0:3], a[0:3], 1.0 ; encoding: [0x00,0x80,0xb7,0xd3,0x00,0x01,0xca,0x1b]147// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU148v_mfma_f32_32x32x16_bf16 a[0:15], a[0:3], a[0:3], 1.0149 150// GFX950: v_mfma_f32_32x32x16_bf16 v[0:15], v[0:3], v[0:3], v[0:15] blgp:5 ; encoding: [0x00,0x00,0xb7,0xd3,0x00,0x01,0x02,0xa4]151// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU152v_mfma_f32_32x32x16_bf16 v[0:15], v[0:3], v[0:3], v[0:15] blgp:5153 154// GFX950: v_mfma_f32_32x32x16_bf16 a[0:15], a[0:3], a[0:3], a[0:15] blgp:2 ; encoding: [0x00,0x80,0xb7,0xd3,0x00,0x01,0x02,0x5c]155// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU156v_mfma_f32_32x32x16_bf16 a[0:15], a[0:3], a[0:3], a[0:15] blgp:2157 158// GFX950: v_mfma_f32_32x32x16_bf16 v[0:15], v[0:3], v[0:3], v[0:15] cbsz:3 ; encoding: [0x00,0x03,0xb7,0xd3,0x00,0x01,0x02,0x04]159// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU160v_mfma_f32_32x32x16_bf16 v[0:15], v[0:3], v[0:3], v[0:15] cbsz:3161 162// GFX950: v_mfma_f32_32x32x16_bf16 v[0:15], v[0:3], v[0:3], v[0:15] abid:1 ; encoding: [0x00,0x08,0xb7,0xd3,0x00,0x01,0x02,0x04]163// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU164v_mfma_f32_32x32x16_bf16 v[0:15], v[0:3], v[0:3], v[0:15] abid:1165 166// GFX950: v_mfma_f32_32x32x16_bf16 a[0:15], a[0:3], a[0:3], a[0:15] cbsz:3 abid:1 ; encoding: [0x00,0x8b,0xb7,0xd3,0x00,0x01,0x02,0x1c]167// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU168v_mfma_f32_32x32x16_bf16 a[0:15], a[0:3], a[0:3], a[0:15] cbsz:3 abid:1169 170//===----------------------------------------------------------------------===//171// v_mfma_ld_scale_b32172//===----------------------------------------------------------------------===//173 174// GFX950: v_mfma_ld_scale_b32 v0, 64 ; encoding: [0x00,0x40,0xac,0xd3,0x00,0x81,0x01,0x18]175// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU176v_mfma_ld_scale_b32 v0, 64177 178// GFX950: v_mfma_ld_scale_b32 64, v0 ; encoding: [0x00,0x40,0xac,0xd3,0xc0,0x00,0x02,0x18]179// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU180v_mfma_ld_scale_b32 64, v0181 182// GFX950: v_mfma_ld_scale_b32 64, 64 ; encoding: [0x00,0x40,0xac,0xd3,0xc0,0x80,0x01,0x18]183// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU184v_mfma_ld_scale_b32 64, 64185 186// GFX950: v_mfma_ld_scale_b32 s0, s0 ; encoding: [0x00,0x40,0xac,0xd3,0x00,0x00,0x00,0x18]187// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU188v_mfma_ld_scale_b32 s0, s0189 190// GFX950: v_mfma_ld_scale_b32 s0, v0 ; encoding: [0x00,0x40,0xac,0xd3,0x00,0x00,0x02,0x18]191// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU192v_mfma_ld_scale_b32 s0, v0193 194// GFX950: v_mfma_ld_scale_b32 v0, s0 ; encoding: [0x00,0x40,0xac,0xd3,0x00,0x01,0x00,0x18]195// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU196v_mfma_ld_scale_b32 v0, s0197 198// GFX950: v_mfma_ld_scale_b32 vcc_lo, vcc_lo ; encoding: [0x00,0x40,0xac,0xd3,0x6a,0xd4,0x00,0x18]199// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU200v_mfma_ld_scale_b32 vcc_lo, vcc_lo201 202// GFX950: v_mfma_ld_scale_b32 m0, m0 ; encoding: [0x00,0x40,0xac,0xd3,0x7c,0xf8,0x00,0x18]203// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU204v_mfma_ld_scale_b32 m0, m0205 206// GFX950: v_mfma_ld_scale_b32 src_vccz, src_vccz ; encoding: [0x00,0x40,0xac,0xd3,0xfb,0xf6,0x01,0x18]207// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU208v_mfma_ld_scale_b32 vccz, vccz209 210// GFX950: v_mfma_ld_scale_b32 src_execz, src_execz ; encoding: [0x00,0x40,0xac,0xd3,0xfc,0xf8,0x01,0x18]211// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU212v_mfma_ld_scale_b32 execz, execz213 214// GFX950: v_mfma_ld_scale_b32 v0, v0 ; encoding: [0x00,0x40,0xac,0xd3,0x00,0x01,0x02,0x18]215// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU216v_mfma_ld_scale_b32 v0, v0217 218// GFX950: v_mfma_ld_scale_b32 v1, v1 ; encoding: [0x00,0x40,0xac,0xd3,0x01,0x03,0x02,0x18]219// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU220v_mfma_ld_scale_b32 v1, v1221 222// GFX950: v_mfma_ld_scale_b32 0, 0 ; encoding: [0x00,0x40,0xac,0xd3,0x80,0x00,0x01,0x18]223// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU224v_mfma_ld_scale_b32 0, 0225 226// GFX950: v_mfma_ld_scale_b32 1, 0 ; encoding: [0x00,0x40,0xac,0xd3,0x81,0x00,0x01,0x18]227// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU228v_mfma_ld_scale_b32 1, 0229 230// GFX950: v_mfma_ld_scale_b32 v1, v1 op_sel:[1,0] ; encoding: [0x00,0x48,0xac,0xd3,0x01,0x03,0x02,0x18]231// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU232v_mfma_ld_scale_b32 v1, v1 op_sel:[1, 0]233 234// GFX950: v_mfma_ld_scale_b32 v1, v1 op_sel:[0,1] ; encoding: [0x00,0x50,0xac,0xd3,0x01,0x03,0x02,0x18]235// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU236v_mfma_ld_scale_b32 v1, v1 op_sel:[0, 1]237 238// GFX950: v_mfma_ld_scale_b32 v1, v1 op_sel:[1,1] ; encoding: [0x00,0x58,0xac,0xd3,0x01,0x03,0x02,0x18]239// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU240v_mfma_ld_scale_b32 v1, v1 op_sel:[1, 1]241 242// GFX950: v_mfma_ld_scale_b32 v1, v1 op_sel_hi:[1,0] ; encoding: [0x00,0x40,0xac,0xd3,0x01,0x03,0x02,0x08]243// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU244v_mfma_ld_scale_b32 v1, v1 op_sel_hi:[1, 0]245 246// GFX950: v_mfma_ld_scale_b32 v1, v1 op_sel_hi:[0,1] ; encoding: [0x00,0x40,0xac,0xd3,0x01,0x03,0x02,0x10]247// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU248v_mfma_ld_scale_b32 v1, v1 op_sel_hi:[0, 1]249 250// GFX950: v_mfma_ld_scale_b32 v1, v1 ; encoding: [0x00,0x40,0xac,0xd3,0x01,0x03,0x02,0x18]251// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU252v_mfma_ld_scale_b32 v1, v1 op_sel_hi:[1, 1]253 254// GFX950: v_mfma_ld_scale_b32 v1, v1 op_sel_hi:[0,0] ; encoding: [0x00,0x40,0xac,0xd3,0x01,0x03,0x02,0x00]255// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU256v_mfma_ld_scale_b32 v1, v1 op_sel:[0,0] op_sel_hi:[0,0]257 258// GFX950: v_mfma_ld_scale_b32 v1, v1 op_sel:[1,0] op_sel_hi:[1,0] ; encoding: [0x00,0x48,0xac,0xd3,0x01,0x03,0x02,0x08]259// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU260v_mfma_ld_scale_b32 v1, v1 op_sel:[1,0] op_sel_hi:[1,0]261 262// GFX950: v_mfma_ld_scale_b32 v1, v1 op_sel:[0,1] op_sel_hi:[0,1] ; encoding: [0x00,0x50,0xac,0xd3,0x01,0x03,0x02,0x10]263// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU264v_mfma_ld_scale_b32 v1, v1 op_sel:[0,1] op_sel_hi:[0,1]265 266// GFX950: v_mfma_ld_scale_b32 v1, v1 op_sel:[0,1] ; encoding: [0x00,0x50,0xac,0xd3,0x01,0x03,0x02,0x18]267// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU268v_mfma_ld_scale_b32 v1, v1 op_sel:[0,1] op_sel_hi:[1,1]269 270// GFX950: v_mfma_ld_scale_b32 v1, v1 op_sel:[0,1] ; encoding: [0x00,0x50,0xac,0xd3,0x01,0x03,0x02,0x18]271// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU272v_mfma_ld_scale_b32 v1, v1 op_sel:[0,1] op_sel_hi:[1,1]273 274// GFX950: v_mfma_ld_scale_b32 v1, v1 op_sel:[1,1] ; encoding: [0x00,0x58,0xac,0xd3,0x01,0x03,0x02,0x18]275// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU276v_mfma_ld_scale_b32 v1, v1 op_sel:[1,1] op_sel_hi:[1,1]277 278// GFX950: v_mfma_ld_scale_b32 v1, v1 op_sel:[1,0] op_sel_hi:[0,1] ; encoding: [0x00,0x48,0xac,0xd3,0x01,0x03,0x02,0x10]279// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU280v_mfma_ld_scale_b32 v1, v1 op_sel:[1,0] op_sel_hi:[0,1]281 282// GFX950: v_mfma_ld_scale_b32 v1, v1 op_sel:[0,1] op_sel_hi:[1,0] ; encoding: [0x00,0x50,0xac,0xd3,0x01,0x03,0x02,0x08]283// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU284v_mfma_ld_scale_b32 v1, v1 op_sel:[0,1] op_sel_hi:[1,0]285 286 287//===----------------------------------------------------------------------===//288// v_mfma_f32_16x16x128_f8f6f4289//===----------------------------------------------------------------------===//290 291// GFX950: v_mfma_f32_16x16x128_f8f6f4 v[0:3], v[4:11], v[4:11], v[0:3] ; encoding: [0x00,0x00,0xad,0xd3,0x04,0x09,0x02,0x04]292// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU293v_mfma_f32_16x16x128_f8f6f4 v[0:3], v[4:11], v[4:11], v[0:3]294 295// GFX950: v_mfma_f32_16x16x128_f8f6f4 v[0:3], v[4:11], v[4:11], v[0:3] blgp:1 ; encoding: [0x00,0x00,0xad,0xd3,0x04,0x09,0x02,0x24]296// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU297v_mfma_f32_16x16x128_f8f6f4 v[0:3], v[4:11], v[4:11], v[0:3] blgp:1298 299// GFX950: v_mfma_f32_16x16x128_f8f6f4 v[0:3], v[4:9], v[4:11], v[0:3] cbsz:3 ; encoding: [0x00,0x03,0xad,0xd3,0x04,0x09,0x02,0x04]300// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU301v_mfma_f32_16x16x128_f8f6f4 v[0:3], v[4:9], v[4:11], v[0:3] cbsz:3302 303// GFX950: v_mfma_f32_16x16x128_f8f6f4 v[0:3], v[4:11], v[4:11], v[0:3] ; encoding: [0x00,0x00,0xad,0xd3,0x04,0x09,0x02,0x04]304// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU305v_mfma_f32_16x16x128_f8f6f4 v[0:3], v[4:11], v[4:11], v[0:3]306 307// GFX950: v_mfma_f32_16x16x128_f8f6f4 v[0:3], v[4:9], v[4:11], v[0:3] cbsz:3 blgp:1 ; encoding: [0x00,0x03,0xad,0xd3,0x04,0x09,0x02,0x24]308// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU309v_mfma_f32_16x16x128_f8f6f4 v[0:3], v[4:9], v[4:11], v[0:3] cbsz:3 blgp:1310 311// GFX950: v_mfma_f32_16x16x128_f8f6f4 a[0:3], a[4:11], a[4:11], a[0:3] ; encoding: [0x00,0x80,0xad,0xd3,0x04,0x09,0x02,0x1c]312// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU313v_mfma_f32_16x16x128_f8f6f4 a[0:3], a[4:11], a[4:11], a[0:3]314 315// GFX950: v_mfma_f32_16x16x128_f8f6f4 a[0:3], a[4:11], a[4:11], a[0:3] blgp:1 ; encoding: [0x00,0x80,0xad,0xd3,0x04,0x09,0x02,0x3c]316// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU317v_mfma_f32_16x16x128_f8f6f4 a[0:3], a[4:11], a[4:11], a[0:3] blgp:1318 319// GFX950: v_mfma_f32_16x16x128_f8f6f4 a[0:3], a[4:9], a[4:11], a[0:3] cbsz:3 ; encoding: [0x00,0x83,0xad,0xd3,0x04,0x09,0x02,0x1c]320// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU321v_mfma_f32_16x16x128_f8f6f4 a[0:3], a[4:9], a[4:11], a[0:3] cbsz:3322 323// GFX950: v_mfma_f32_16x16x128_f8f6f4 a[0:3], a[4:11], a[4:11], a[0:3] ; encoding: [0x00,0x80,0xad,0xd3,0x04,0x09,0x02,0x1c]324// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU325v_mfma_f32_16x16x128_f8f6f4 a[0:3], a[4:11], a[4:11], a[0:3]326 327// GFX950: v_mfma_f32_16x16x128_f8f6f4 a[0:3], a[4:11], a[4:9], a[0:3] cbsz:1 blgp:3 ; encoding: [0x00,0x81,0xad,0xd3,0x04,0x09,0x02,0x7c]328// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU329v_mfma_f32_16x16x128_f8f6f4 a[0:3], a[4:11], a[4:9], a[0:3] cbsz:1 blgp:3330 331//===----------------------------------------------------------------------===//332// v_mfma_f32_32x32x64_f8f6f4333//===----------------------------------------------------------------------===//334 335// GFX950: v_mfma_f32_32x32x64_f8f6f4 v[0:15], v[4:11], v[4:11], v[0:15] ; encoding: [0x00,0x00,0xae,0xd3,0x04,0x09,0x02,0x04]336// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU337v_mfma_f32_32x32x64_f8f6f4 v[0:15], v[4:11], v[4:11], v[0:15]338 339// GFX950: v_mfma_f32_32x32x64_f8f6f4 v[0:15], v[4:11], v[4:11], v[0:15] blgp:1 ; encoding: [0x00,0x00,0xae,0xd3,0x04,0x09,0x02,0x24]340// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU341v_mfma_f32_32x32x64_f8f6f4 v[0:15], v[4:11], v[4:11], v[0:15] blgp:1342 343// GFX950: v_mfma_f32_32x32x64_f8f6f4 v[0:15], v[4:9], v[4:11], v[0:15] cbsz:3 ; encoding: [0x00,0x03,0xae,0xd3,0x04,0x09,0x02,0x04]344// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU345v_mfma_f32_32x32x64_f8f6f4 v[0:15], v[4:9], v[4:11], v[0:15] cbsz:3346 347// GFX950: v_mfma_f32_32x32x64_f8f6f4 v[0:15], v[4:11], v[4:11], v[0:15] ; encoding: [0x00,0x00,0xae,0xd3,0x04,0x09,0x02,0x04]348// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU349v_mfma_f32_32x32x64_f8f6f4 v[0:15], v[4:11], v[4:11], v[0:15]350 351// GFX950: v_mfma_f32_32x32x64_f8f6f4 v[0:15], v[4:9], v[4:11], v[0:15] cbsz:3 blgp:1 ; encoding: [0x00,0x03,0xae,0xd3,0x04,0x09,0x02,0x24]352// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU353v_mfma_f32_32x32x64_f8f6f4 v[0:15], v[4:9], v[4:11], v[0:15] cbsz:3 blgp:1354 355// GFX950: v_mfma_f32_32x32x64_f8f6f4 a[0:15], a[4:11], a[4:11], a[0:15] ; encoding: [0x00,0x80,0xae,0xd3,0x04,0x09,0x02,0x1c]356// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU357v_mfma_f32_32x32x64_f8f6f4 a[0:15], a[4:11], a[4:11], a[0:15]358 359// GFX950: v_mfma_f32_32x32x64_f8f6f4 a[0:15], a[4:11], a[4:11], a[0:15] blgp:1 ; encoding: [0x00,0x80,0xae,0xd3,0x04,0x09,0x02,0x3c]360// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU361v_mfma_f32_32x32x64_f8f6f4 a[0:15], a[4:11], a[4:11], a[0:15] blgp:1362 363// GFX950: v_mfma_f32_32x32x64_f8f6f4 a[0:15], a[4:9], a[4:11], a[0:15] cbsz:3 ; encoding: [0x00,0x83,0xae,0xd3,0x04,0x09,0x02,0x1c]364// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU365v_mfma_f32_32x32x64_f8f6f4 a[0:15], a[4:9], a[4:11], a[0:15] cbsz:3366 367// GFX950: v_mfma_f32_32x32x64_f8f6f4 a[0:15], a[4:11], a[4:11], a[0:15] ; encoding: [0x00,0x80,0xae,0xd3,0x04,0x09,0x02,0x1c]368// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU369v_mfma_f32_32x32x64_f8f6f4 a[0:15], a[4:11], a[4:11], a[0:15]370 371// GFX950: v_mfma_f32_32x32x64_f8f6f4 a[0:15], a[4:11], a[4:9], a[0:15] cbsz:1 blgp:3 ; encoding: [0x00,0x81,0xae,0xd3,0x04,0x09,0x02,0x7c]372// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU373v_mfma_f32_32x32x64_f8f6f4 a[0:15], a[4:11], a[4:9], a[0:15] cbsz:1 blgp:3374 375//===----------------------------------------------------------------------===//376// v_mfma_scale_f32_16x16x128_f8f6f4377//===----------------------------------------------------------------------===//378// FIXME: Test neg, clamp379 380// GFX950: v_mfma_scale_f32_16x16x128_f8f6f4 v[0:3], v[4:11], v[12:19], v[20:23], v24, v25 op_sel_hi:[0,0,0] ; encoding: [0x00,0x00,0xac,0xd3,0x18,0x33,0x02,0x00,0x00,0x08,0xad,0xd3,0x04,0x19,0x52,0x04]381// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU382v_mfma_scale_f32_16x16x128_f8f6f4 v[0:3], v[4:11], v[12:19], v[20:23], v24, v25383 384// GFX950: v_mfma_scale_f32_16x16x128_f8f6f4 a[0:3], a[4:11], a[12:19], a[20:23], v24, v25 op_sel_hi:[0,0,0] ; encoding: [0x00,0x00,0xac,0xd3,0x18,0x33,0x02,0x00,0x00,0x88,0xad,0xd3,0x04,0x19,0x52,0x1c]385// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU386v_mfma_scale_f32_16x16x128_f8f6f4 a[0:3], a[4:11], a[12:19], a[20:23], v24, v25387 388// GFX950: v_mfma_scale_f32_16x16x128_f8f6f4 a[0:3], v[4:11], v[12:19], a[20:23], v24, v25 op_sel_hi:[0,0,0] ; encoding: [0x00,0x00,0xac,0xd3,0x18,0x33,0x02,0x00,0x00,0x88,0xad,0xd3,0x04,0x19,0x52,0x04]389// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU390v_mfma_scale_f32_16x16x128_f8f6f4 a[0:3], v[4:11], v[12:19], a[20:23], v24, v25391 392// GFX950: v_mfma_scale_f32_16x16x128_f8f6f4 v[0:3], a[4:11], a[12:19], v[20:23], v24, v25 op_sel_hi:[0,0,0] ; encoding: [0x00,0x00,0xac,0xd3,0x18,0x33,0x02,0x00,0x00,0x08,0xad,0xd3,0x04,0x19,0x52,0x1c]393// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU394v_mfma_scale_f32_16x16x128_f8f6f4 v[0:3], a[4:11], a[12:19], v[20:23], v24, v25395 396// GFX950: v_mfma_scale_f32_16x16x128_f8f6f4 v[0:3], v[4:11], a[12:19], v[20:23], v24, v25 op_sel_hi:[0,0,0] ; encoding: [0x00,0x00,0xac,0xd3,0x18,0x33,0x02,0x00,0x00,0x08,0xad,0xd3,0x04,0x19,0x52,0x14]397// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU398v_mfma_scale_f32_16x16x128_f8f6f4 v[0:3], v[4:11], a[12:19], v[20:23], v24, v25399 400// GFX950: v_mfma_scale_f32_16x16x128_f8f6f4 v[0:3], a[4:11], v[12:19], v[20:23], v24, v25 op_sel_hi:[0,0,0] ; encoding: [0x00,0x00,0xac,0xd3,0x18,0x33,0x02,0x00,0x00,0x08,0xad,0xd3,0x04,0x19,0x52,0x0c]401// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU402v_mfma_scale_f32_16x16x128_f8f6f4 v[0:3], a[4:11], v[12:19], v[20:23], v24, v25403 404// GFX950: v_mfma_scale_f32_16x16x128_f8f6f4 v[50:53], v[4:11], v[12:19], v[20:23], v24, v25 op_sel_hi:[0,0,0] ; encoding: [0x00,0x00,0xac,0xd3,0x18,0x33,0x02,0x00,0x32,0x08,0xad,0xd3,0x04,0x19,0x52,0x04]405// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU406v_mfma_scale_f32_16x16x128_f8f6f4 v[50:53], v[4:11], v[12:19], v[20:23], v24, v25407 408// GFX950: v_mfma_scale_f32_16x16x128_f8f6f4 v[0:3], v[4:9], v[12:19], v[20:23], v24, v25 op_sel_hi:[0,0,0] cbsz:3 blgp:1 ; encoding: [0x00,0x00,0xac,0xd3,0x18,0x33,0x02,0x00,0x00,0x0b,0xad,0xd3,0x04,0x19,0x52,0x24]409// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU410v_mfma_scale_f32_16x16x128_f8f6f4 v[0:3], v[4:9], v[12:19], v[20:23], v24, v25 cbsz:3 blgp:1411 412// GFX950: v_mfma_scale_f32_16x16x128_f8f6f4 v[0:3], v[4:9], v[12:19], v[20:23], v24, v25 op_sel:[0,1,0] op_sel_hi:[0,0,0] cbsz:3 blgp:1 ; encoding: [0x00,0x10,0xac,0xd3,0x18,0x33,0x02,0x00,0x00,0x0b,0xad,0xd3,0x04,0x19,0x52,0x24]413// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU414v_mfma_scale_f32_16x16x128_f8f6f4 v[0:3], v[4:9], v[12:19], v[20:23], v24, v25 op_sel:[0,1,0] cbsz:3 blgp:1415 416// GFX950: v_mfma_scale_f32_16x16x128_f8f6f4 v[0:3], v[4:9], v[12:19], v[20:23], v24, v25 op_sel:[0,1,0] op_sel_hi:[0,1,0] cbsz:3 blgp:1 ; encoding: [0x00,0x10,0xac,0xd3,0x18,0x33,0x02,0x10,0x00,0x0b,0xad,0xd3,0x04,0x19,0x52,0x24]417// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU418v_mfma_scale_f32_16x16x128_f8f6f4 v[0:3], v[4:9], v[12:19], v[20:23], v24, v25 op_sel:[0,1,0] op_sel_hi:[0,1,0] cbsz:3 blgp:1419 420// GFX950: v_mfma_scale_f32_16x16x128_f8f6f4 v[0:3], v[4:11], v[12:17], v[20:23], v24, v25 op_sel:[0,1,0] op_sel_hi:[0,1,0] blgp:2 ; encoding: [0x00,0x10,0xac,0xd3,0x18,0x33,0x02,0x10,0x00,0x08,0xad,0xd3,0x04,0x19,0x52,0x44]421// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU422v_mfma_scale_f32_16x16x128_f8f6f4 v[0:3], v[4:11], v[12:17], v[20:23], v24, v25 op_sel:[0,1,0] op_sel_hi:[0,1,0] blgp:2423 424// GFX950: v_mfma_scale_f32_16x16x128_f8f6f4 v[0:3], v[4:11], v[12:19], v[20:23], v24, v25 op_sel:[0,1,0] op_sel_hi:[0,1,0] ; encoding: [0x00,0x10,0xac,0xd3,0x18,0x33,0x02,0x10,0x00,0x08,0xad,0xd3,0x04,0x19,0x52,0x04]425// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU426v_mfma_scale_f32_16x16x128_f8f6f4 v[0:3], v[4:11], v[12:19], v[20:23], v24, v25 op_sel:[0,1,0] op_sel_hi:[0,1,0]427 428// GFX950: v_mfma_scale_f32_16x16x128_f8f6f4 v[0:3], v[4:9], v[12:19], v[20:23], v24, v25 op_sel:[0,1,0] op_sel_hi:[0,1,0] cbsz:3 ; encoding: [0x00,0x10,0xac,0xd3,0x18,0x33,0x02,0x10,0x00,0x0b,0xad,0xd3,0x04,0x19,0x52,0x04]429// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU430v_mfma_scale_f32_16x16x128_f8f6f4 v[0:3], v[4:9], v[12:19], v[20:23], v24, v25 op_sel:[0,1,0] op_sel_hi:[0,1,0] cbsz:3431 432// op_sel combinations433 434// GFX950: v_mfma_scale_f32_16x16x128_f8f6f4 v[0:3], v[4:9], v[12:19], v[20:23], v24, v25 op_sel:[0,1,0] op_sel_hi:[0,0,0] cbsz:3 blgp:1 ; encoding: [0x00,0x10,0xac,0xd3,0x18,0x33,0x02,0x00,0x00,0x0b,0xad,0xd3,0x04,0x19,0x52,0x24]435// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU436v_mfma_scale_f32_16x16x128_f8f6f4 v[0:3], v[4:9], v[12:19], v[20:23], v24, v25 op_sel:[0,1,0] cbsz:3 blgp:1437 438// GFX950: v_mfma_scale_f32_16x16x128_f8f6f4 v[0:3], v[4:9], v[12:19], v[20:23], v24, v25 op_sel:[1,0,0] op_sel_hi:[0,0,0] cbsz:3 blgp:1 ; encoding: [0x00,0x08,0xac,0xd3,0x18,0x33,0x02,0x00,0x00,0x0b,0xad,0xd3,0x04,0x19,0x52,0x24]439// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU440v_mfma_scale_f32_16x16x128_f8f6f4 v[0:3], v[4:9], v[12:19], v[20:23], v24, v25 op_sel:[1,0,0] cbsz:3 blgp:1441 442// GFX950: v_mfma_scale_f32_16x16x128_f8f6f4 v[0:3], v[4:9], v[12:19], v[20:23], v24, v25 op_sel:[1,1,0] op_sel_hi:[0,0,0] cbsz:3 blgp:1 ; encoding: [0x00,0x18,0xac,0xd3,0x18,0x33,0x02,0x00,0x00,0x0b,0xad,0xd3,0x04,0x19,0x52,0x24]443// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU444v_mfma_scale_f32_16x16x128_f8f6f4 v[0:3], v[4:9], v[12:19], v[20:23], v24, v25 op_sel:[1,1,0] cbsz:3 blgp:1445 446// GFX950: v_mfma_scale_f32_16x16x128_f8f6f4 v[0:3], v[4:9], v[12:19], v[20:23], v24, v25 op_sel_hi:[0,1,0] cbsz:3 blgp:1 ; encoding: [0x00,0x00,0xac,0xd3,0x18,0x33,0x02,0x10,0x00,0x0b,0xad,0xd3,0x04,0x19,0x52,0x24]447// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU448v_mfma_scale_f32_16x16x128_f8f6f4 v[0:3], v[4:9], v[12:19], v[20:23], v24, v25 op_sel_hi:[0,1,0] cbsz:3 blgp:1449 450// GFX950: v_mfma_scale_f32_16x16x128_f8f6f4 v[0:3], v[4:9], v[12:19], v[20:23], v24, v25 op_sel_hi:[1,0,0] cbsz:3 blgp:1 ; encoding: [0x00,0x00,0xac,0xd3,0x18,0x33,0x02,0x08,0x00,0x0b,0xad,0xd3,0x04,0x19,0x52,0x24]451// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU452v_mfma_scale_f32_16x16x128_f8f6f4 v[0:3], v[4:9], v[12:19], v[20:23], v24, v25 op_sel_hi:[1,0,0] cbsz:3 blgp:1453 454// GFX950: v_mfma_scale_f32_16x16x128_f8f6f4 v[0:3], v[4:9], v[12:19], v[20:23], v24, v25 op_sel_hi:[1,1,0] cbsz:3 blgp:1 ; encoding: [0x00,0x00,0xac,0xd3,0x18,0x33,0x02,0x18,0x00,0x0b,0xad,0xd3,0x04,0x19,0x52,0x24]455// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU456v_mfma_scale_f32_16x16x128_f8f6f4 v[0:3], v[4:9], v[12:19], v[20:23], v24, v25 op_sel_hi:[1,1,0] cbsz:3 blgp:1457 458// GFX950: v_mfma_scale_f32_16x16x128_f8f6f4 v[0:3], v[4:9], v[12:19], v[20:23], v24, v25 op_sel:[0,1,0] op_sel_hi:[1,1,0] cbsz:3 blgp:1 ; encoding: [0x00,0x10,0xac,0xd3,0x18,0x33,0x02,0x18,0x00,0x0b,0xad,0xd3,0x04,0x19,0x52,0x24]459// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU460v_mfma_scale_f32_16x16x128_f8f6f4 v[0:3], v[4:9], v[12:19], v[20:23], v24, v25 op_sel:[0,1,0] op_sel_hi:[1,1,0] cbsz:3 blgp:1461 462// GFX950: v_mfma_scale_f32_16x16x128_f8f6f4 v[0:3], v[4:9], v[12:19], v[20:23], v24, v25 op_sel:[0,1,0] op_sel_hi:[1,0,0] cbsz:3 blgp:1 ; encoding: [0x00,0x10,0xac,0xd3,0x18,0x33,0x02,0x08,0x00,0x0b,0xad,0xd3,0x04,0x19,0x52,0x24]463// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU464v_mfma_scale_f32_16x16x128_f8f6f4 v[0:3], v[4:9], v[12:19], v[20:23], v24, v25 op_sel:[0,1,0] op_sel_hi:[1,0,0] cbsz:3 blgp:1465 466// GFX950: v_mfma_scale_f32_16x16x128_f8f6f4 v[0:3], v[4:9], v[12:19], v[20:23], v24, v25 op_sel:[0,1,0] op_sel_hi:[0,1,0] cbsz:3 blgp:1 ; encoding: [0x00,0x10,0xac,0xd3,0x18,0x33,0x02,0x10,0x00,0x0b,0xad,0xd3,0x04,0x19,0x52,0x24]467// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU468v_mfma_scale_f32_16x16x128_f8f6f4 v[0:3], v[4:9], v[12:19], v[20:23], v24, v25 op_sel:[0,1,0] op_sel_hi:[0,1,0] cbsz:3 blgp:1469 470// GFX950: v_mfma_scale_f32_16x16x128_f8f6f4 v[0:3], v[4:9], v[12:19], v[20:23], v24, v25 op_sel:[1,0,0] op_sel_hi:[0,1,0] cbsz:3 blgp:1 ; encoding: [0x00,0x08,0xac,0xd3,0x18,0x33,0x02,0x10,0x00,0x0b,0xad,0xd3,0x04,0x19,0x52,0x24]471// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU472v_mfma_scale_f32_16x16x128_f8f6f4 v[0:3], v[4:9], v[12:19], v[20:23], v24, v25 op_sel:[1,0,0] op_sel_hi:[0,1,0] cbsz:3 blgp:1473 474// GFX950: v_mfma_scale_f32_16x16x128_f8f6f4 v[0:3], v[4:9], v[12:19], v[20:23], v24, v25 op_sel:[1,0,0] op_sel_hi:[1,0,0] cbsz:3 blgp:1 ; encoding: [0x00,0x08,0xac,0xd3,0x18,0x33,0x02,0x08,0x00,0x0b,0xad,0xd3,0x04,0x19,0x52,0x24]475// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU476v_mfma_scale_f32_16x16x128_f8f6f4 v[0:3], v[4:9], v[12:19], v[20:23], v24, v25 op_sel:[1,0,0] op_sel_hi:[1,0,0] cbsz:3 blgp:1477 478// GFX950: v_mfma_scale_f32_16x16x128_f8f6f4 v[0:3], v[4:9], v[12:19], v[20:23], v24, v25 op_sel:[1,0,0] op_sel_hi:[1,1,0] cbsz:3 blgp:1 ; encoding: [0x00,0x08,0xac,0xd3,0x18,0x33,0x02,0x18,0x00,0x0b,0xad,0xd3,0x04,0x19,0x52,0x24]479// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU480v_mfma_scale_f32_16x16x128_f8f6f4 v[0:3], v[4:9], v[12:19], v[20:23], v24, v25 op_sel:[1,0,0] op_sel_hi:[1,1,0] cbsz:3 blgp:1481 482 483//===----------------------------------------------------------------------===//484// v_mfma_scale_f32_32x32x64_f8f6f4485//===----------------------------------------------------------------------===//486// FIXME: Test neg, clamp487 488// GFX950: v_mfma_scale_f32_32x32x64_f8f6f4 v[0:15], v[16:23], v[24:31], v[32:47], v48, v49 op_sel_hi:[0,0,0] ; encoding: [0x00,0x00,0xac,0xd3,0x30,0x63,0x02,0x00,0x00,0x08,0xae,0xd3,0x10,0x31,0x82,0x04]489// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU490v_mfma_scale_f32_32x32x64_f8f6f4 v[0:15], v[16:23], v[24:31], v[32:47], v48, v49491 492// GFX950: v_mfma_scale_f32_32x32x64_f8f6f4 a[0:15], a[16:23], a[24:31], a[32:47], v48, v49 op_sel_hi:[0,0,0] ; encoding: [0x00,0x00,0xac,0xd3,0x30,0x63,0x02,0x00,0x00,0x88,0xae,0xd3,0x10,0x31,0x82,0x1c]493// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU494v_mfma_scale_f32_32x32x64_f8f6f4 a[0:15], a[16:23], a[24:31], a[32:47], v48, v49495 496// GFX950: v_mfma_scale_f32_32x32x64_f8f6f4 a[0:15], v[16:23], v[24:31], a[32:47], v48, v49 op_sel_hi:[0,0,0] ; encoding: [0x00,0x00,0xac,0xd3,0x30,0x63,0x02,0x00,0x00,0x88,0xae,0xd3,0x10,0x31,0x82,0x04]497// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU498v_mfma_scale_f32_32x32x64_f8f6f4 a[0:15], v[16:23], v[24:31], a[32:47], v48, v49499 500// GFX950: v_mfma_scale_f32_32x32x64_f8f6f4 v[0:15], a[16:23], a[24:31], v[32:47], v48, v49 op_sel_hi:[0,0,0] ; encoding: [0x00,0x00,0xac,0xd3,0x30,0x63,0x02,0x00,0x00,0x08,0xae,0xd3,0x10,0x31,0x82,0x1c]501// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU502v_mfma_scale_f32_32x32x64_f8f6f4 v[0:15], a[16:23], a[24:31], v[32:47], v48, v49503 504// GFX950: v_mfma_scale_f32_32x32x64_f8f6f4 v[0:15], v[16:23], a[24:31], v[32:47], v48, v49 op_sel_hi:[0,0,0] ; encoding: [0x00,0x00,0xac,0xd3,0x30,0x63,0x02,0x00,0x00,0x08,0xae,0xd3,0x10,0x31,0x82,0x14]505// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU506v_mfma_scale_f32_32x32x64_f8f6f4 v[0:15], v[16:23], a[24:31], v[32:47], v48, v49507 508// GFX950: v_mfma_scale_f32_32x32x64_f8f6f4 a[0:15], a[16:23], v[24:31], a[32:47], v48, v49 op_sel_hi:[0,0,0] ; encoding: [0x00,0x00,0xac,0xd3,0x30,0x63,0x02,0x00,0x00,0x88,0xae,0xd3,0x10,0x31,0x82,0x0c]509// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU510v_mfma_scale_f32_32x32x64_f8f6f4 a[0:15], a[16:23], v[24:31], a[32:47], v48, v49511 512// GFX950: v_mfma_scale_f32_32x32x64_f8f6f4 v[50:65], v[16:23], v[24:31], v[32:47], v48, v49 op_sel_hi:[0,0,0] ; encoding: [0x00,0x00,0xac,0xd3,0x30,0x63,0x02,0x00,0x32,0x08,0xae,0xd3,0x10,0x31,0x82,0x04]513// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU514v_mfma_scale_f32_32x32x64_f8f6f4 v[50:65], v[16:23], v[24:31], v[32:47], v48, v49515 516// GFX950: v_mfma_scale_f32_32x32x64_f8f6f4 v[0:15], v[16:21], v[24:31], v[32:47], v48, v49 op_sel_hi:[0,0,0] cbsz:3 blgp:1 ; encoding: [0x00,0x00,0xac,0xd3,0x30,0x63,0x02,0x00,0x00,0x0b,0xae,0xd3,0x10,0x31,0x82,0x24]517// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU518v_mfma_scale_f32_32x32x64_f8f6f4 v[0:15], v[16:21], v[24:31], v[32:47], v48, v49 cbsz:3 blgp:1519 520// GFX950: v_mfma_scale_f32_32x32x64_f8f6f4 v[0:15], v[16:21], v[24:31], v[32:47], v48, v49 op_sel_hi:[0,0,0] cbsz:2 ; encoding: [0x00,0x00,0xac,0xd3,0x30,0x63,0x02,0x00,0x00,0x0a,0xae,0xd3,0x10,0x31,0x82,0x04]521// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU522v_mfma_scale_f32_32x32x64_f8f6f4 v[0:15], v[16:21], v[24:31], v[32:47], v48, v49 cbsz:2523 524// GFX950: v_mfma_scale_f32_32x32x64_f8f6f4 v[0:15], v[16:23], v[24:31], v[32:47], v48, v49 op_sel_hi:[0,0,0] ; encoding: [0x00,0x00,0xac,0xd3,0x30,0x63,0x02,0x00,0x00,0x08,0xae,0xd3,0x10,0x31,0x82,0x04]525// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU526v_mfma_scale_f32_32x32x64_f8f6f4 v[0:15], v[16:23], v[24:31], v[32:47], v48, v49527 528// GFX950: v_mfma_scale_f32_32x32x64_f8f6f4 v[0:15], v[16:23], v[24:29], v[32:47], v48, v49 op_sel_hi:[0,0,0] blgp:2 ; encoding: [0x00,0x00,0xac,0xd3,0x30,0x63,0x02,0x00,0x00,0x08,0xae,0xd3,0x10,0x31,0x82,0x44]529// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU530v_mfma_scale_f32_32x32x64_f8f6f4 v[0:15], v[16:23], v[24:29], v[32:47], v48, v49 blgp:2531 532// GFX950: v_mfma_scale_f32_32x32x64_f8f6f4 v[0:15], v[16:21], v[24:29], v[32:47], v48, v49 op_sel:[0,1,0] op_sel_hi:[0,1,0] cbsz:2 blgp:3 ; encoding: [0x00,0x10,0xac,0xd3,0x30,0x63,0x02,0x10,0x00,0x0a,0xae,0xd3,0x10,0x31,0x82,0x64]533// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU534v_mfma_scale_f32_32x32x64_f8f6f4 v[0:15], v[16:21], v[24:29], v[32:47], v48, v49 op_sel:[0,1,0] op_sel_hi:[0,1,0] cbsz:2 blgp:3535 536// op_sel combinations537 538// GFX950: v_mfma_scale_f32_32x32x64_f8f6f4 v[0:15], v[16:23], v[24:31], v[32:47], v48, v49 op_sel:[0,1,0] op_sel_hi:[0,0,0] ; encoding: [0x00,0x10,0xac,0xd3,0x30,0x63,0x02,0x00,0x00,0x08,0xae,0xd3,0x10,0x31,0x82,0x04]539// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU540v_mfma_scale_f32_32x32x64_f8f6f4 v[0:15], v[16:23], v[24:31], v[32:47], v48, v49 op_sel:[0,1,0]541 542// GFX950: v_mfma_scale_f32_32x32x64_f8f6f4 v[0:15], v[16:23], v[24:31], v[32:47], v48, v49 op_sel:[1,0,0] op_sel_hi:[0,0,0] ; encoding: [0x00,0x08,0xac,0xd3,0x30,0x63,0x02,0x00,0x00,0x08,0xae,0xd3,0x10,0x31,0x82,0x04]543// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU544v_mfma_scale_f32_32x32x64_f8f6f4 v[0:15], v[16:23], v[24:31], v[32:47], v48, v49 op_sel:[1,0,0]545 546// GFX950: v_mfma_scale_f32_32x32x64_f8f6f4 v[0:15], v[16:23], v[24:31], v[32:47], v48, v49 op_sel:[1,1,0] op_sel_hi:[0,0,0] ; encoding: [0x00,0x18,0xac,0xd3,0x30,0x63,0x02,0x00,0x00,0x08,0xae,0xd3,0x10,0x31,0x82,0x04]547// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU548v_mfma_scale_f32_32x32x64_f8f6f4 v[0:15], v[16:23], v[24:31], v[32:47], v48, v49 op_sel:[1,1,0]549 550// GFX950: v_mfma_scale_f32_32x32x64_f8f6f4 v[0:15], v[16:23], v[24:31], v[32:47], v48, v49 op_sel_hi:[0,1,0] ; encoding: [0x00,0x00,0xac,0xd3,0x30,0x63,0x02,0x10,0x00,0x08,0xae,0xd3,0x10,0x31,0x82,0x04]551// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU552v_mfma_scale_f32_32x32x64_f8f6f4 v[0:15], v[16:23], v[24:31], v[32:47], v48, v49 op_sel_hi:[0,1,0]553 554// GFX950: v_mfma_scale_f32_32x32x64_f8f6f4 v[0:15], v[16:23], v[24:31], v[32:47], v48, v49 op_sel_hi:[1,0,0] ; encoding: [0x00,0x00,0xac,0xd3,0x30,0x63,0x02,0x08,0x00,0x08,0xae,0xd3,0x10,0x31,0x82,0x04]555// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU556v_mfma_scale_f32_32x32x64_f8f6f4 v[0:15], v[16:23], v[24:31], v[32:47], v48, v49 op_sel_hi:[1,0,0]557 558// GFX950: v_mfma_scale_f32_32x32x64_f8f6f4 v[0:15], v[16:23], v[24:31], v[32:47], v48, v49 op_sel_hi:[1,1,0] ; encoding: [0x00,0x00,0xac,0xd3,0x30,0x63,0x02,0x18,0x00,0x08,0xae,0xd3,0x10,0x31,0x82,0x04]559// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU560v_mfma_scale_f32_32x32x64_f8f6f4 v[0:15], v[16:23], v[24:31], v[32:47], v48, v49 op_sel_hi:[1,1,0]561 562// GFX950: v_mfma_scale_f32_32x32x64_f8f6f4 v[0:15], v[16:23], v[24:31], v[32:47], v48, v49 op_sel:[0,1,0] op_sel_hi:[0,1,0] ; encoding: [0x00,0x10,0xac,0xd3,0x30,0x63,0x02,0x10,0x00,0x08,0xae,0xd3,0x10,0x31,0x82,0x04]563// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU564v_mfma_scale_f32_32x32x64_f8f6f4 v[0:15], v[16:23], v[24:31], v[32:47], v48, v49 op_sel:[0,1,0] op_sel_hi:[0,1,0]565 566// GFX950: v_mfma_scale_f32_32x32x64_f8f6f4 v[0:15], v[16:23], v[24:31], v[32:47], v48, v49 op_sel:[0,1,0] op_sel_hi:[1,1,0] ; encoding: [0x00,0x10,0xac,0xd3,0x30,0x63,0x02,0x18,0x00,0x08,0xae,0xd3,0x10,0x31,0x82,0x04]567// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU568v_mfma_scale_f32_32x32x64_f8f6f4 v[0:15], v[16:23], v[24:31], v[32:47], v48, v49 op_sel:[0,1,0] op_sel_hi:[1,1,0]569 570// GFX950: v_mfma_scale_f32_32x32x64_f8f6f4 v[0:15], v[16:23], v[24:31], v[32:47], v48, v49 op_sel:[0,1,0] op_sel_hi:[1,0,0] ; encoding: [0x00,0x10,0xac,0xd3,0x30,0x63,0x02,0x08,0x00,0x08,0xae,0xd3,0x10,0x31,0x82,0x04]571// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU572v_mfma_scale_f32_32x32x64_f8f6f4 v[0:15], v[16:23], v[24:31], v[32:47], v48, v49 op_sel:[0,1,0] op_sel_hi:[1,0,0]573 574// GFX950: v_mfma_scale_f32_32x32x64_f8f6f4 v[0:15], v[16:23], v[24:31], v[32:47], v48, v49 op_sel:[1,0,0] op_sel_hi:[0,1,0] ; encoding: [0x00,0x08,0xac,0xd3,0x30,0x63,0x02,0x10,0x00,0x08,0xae,0xd3,0x10,0x31,0x82,0x04]575// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU576v_mfma_scale_f32_32x32x64_f8f6f4 v[0:15], v[16:23], v[24:31], v[32:47], v48, v49 op_sel:[1,0,0] op_sel_hi:[0,1,0]577 578// GFX950: v_mfma_scale_f32_32x32x64_f8f6f4 v[0:15], v[16:23], v[24:31], v[32:47], v48, v49 op_sel:[1,0,0] op_sel_hi:[1,0,0] ; encoding: [0x00,0x08,0xac,0xd3,0x30,0x63,0x02,0x08,0x00,0x08,0xae,0xd3,0x10,0x31,0x82,0x04]579// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU580v_mfma_scale_f32_32x32x64_f8f6f4 v[0:15], v[16:23], v[24:31], v[32:47], v48, v49 op_sel:[1,0,0] op_sel_hi:[1,0,0]581 582// GFX950: v_mfma_scale_f32_32x32x64_f8f6f4 v[0:15], v[16:23], v[24:31], v[32:47], v48, v49 op_sel:[1,0,0] op_sel_hi:[1,1,0] ; encoding: [0x00,0x08,0xac,0xd3,0x30,0x63,0x02,0x18,0x00,0x08,0xae,0xd3,0x10,0x31,0x82,0x04]583// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU584v_mfma_scale_f32_32x32x64_f8f6f4 v[0:15], v[16:23], v[24:31], v[32:47], v48, v49 op_sel:[1,0,0] op_sel_hi:[1,1,0]585 586 587//===----------------------------------------------------------------------===//588// v_mfma_f32_16x16x128_f8f6f4 with appropriate register widths589//===----------------------------------------------------------------------===//590 591// bf8 x fp4592// GFX950: v_mfma_f32_16x16x128_f8f6f4 v[0:3], v[4:11], v[12:15], v[20:23] cbsz:1 blgp:4 ; encoding: [0x00,0x01,0xad,0xd3,0x04,0x19,0x52,0x84]593v_mfma_f32_16x16x128_f8f6f4 v[0:3], v[4:11], v[12:15], v[20:23] cbsz:1 blgp:4594 595// GFX950: v_mfma_f32_16x16x128_f8f6f4 v[0:3], v[4:11], v[12:15], v[20:23] cbsz:1 blgp:4 ; encoding: [0x00,0x01,0xad,0xd3,0x04,0x19,0x52,0x84]596v_mfma_f32_16x16x128_f8f6f4 v[0:3], v[4:11], v[12:15], v[20:23] cbsz:1 blgp:4597 598// GFX950: v_mfma_f32_16x16x128_f8f6f4 a[0:3], a[4:11], a[12:15], a[20:23] cbsz:1 blgp:4 ; encoding: [0x00,0x81,0xad,0xd3,0x04,0x19,0x52,0x9c]599v_mfma_f32_16x16x128_f8f6f4 a[0:3], a[4:11], a[12:15], a[20:23] cbsz:1 blgp:4600 601 602// fp4 x bf8603// GFX950: v_mfma_f32_16x16x128_f8f6f4 v[0:3], v[4:7], v[12:19], v[20:23] cbsz:4 blgp:1 ; encoding: [0x00,0x04,0xad,0xd3,0x04,0x19,0x52,0x24]604v_mfma_f32_16x16x128_f8f6f4 v[0:3], v[4:7], v[12:19], v[20:23] cbsz:4 blgp:1605 606// GFX950: v_mfma_f32_16x16x128_f8f6f4 a[0:3], a[4:7], a[12:19], a[20:23] cbsz:4 blgp:1 ; encoding: [0x00,0x84,0xad,0xd3,0x04,0x19,0x52,0x3c]607v_mfma_f32_16x16x128_f8f6f4 a[0:3], a[4:7], a[12:19], a[20:23] cbsz:4 blgp:1608 609 610// bf6 x bf8611// GFX950: v_mfma_f32_16x16x128_f8f6f4 v[0:3], v[4:9], v[12:19], v[20:23] cbsz:3 blgp:1 ; encoding: [0x00,0x03,0xad,0xd3,0x04,0x19,0x52,0x24]612v_mfma_f32_16x16x128_f8f6f4 v[0:3], v[4:9], v[12:19], v[20:23] cbsz:3 blgp:1613 614// GFX950: v_mfma_f32_16x16x128_f8f6f4 a[0:3], a[4:9], a[12:19], a[20:23] cbsz:3 blgp:1 ; encoding: [0x00,0x83,0xad,0xd3,0x04,0x19,0x52,0x3c]615v_mfma_f32_16x16x128_f8f6f4 a[0:3], a[4:9], a[12:19], a[20:23] cbsz:3 blgp:1616 617 618// bf8 x bf6619// GFX950: v_mfma_f32_16x16x128_f8f6f4 v[0:3], v[4:11], v[12:17], v[20:23] cbsz:1 blgp:3 ; encoding: [0x00,0x01,0xad,0xd3,0x04,0x19,0x52,0x64]620v_mfma_f32_16x16x128_f8f6f4 v[0:3], v[4:11], v[12:17], v[20:23] cbsz:1 blgp:3621 622// GFX950: v_mfma_f32_16x16x128_f8f6f4 a[0:3], a[4:11], a[12:17], a[20:23] cbsz:1 blgp:3 ; encoding: [0x00,0x81,0xad,0xd3,0x04,0x19,0x52,0x7c]623v_mfma_f32_16x16x128_f8f6f4 a[0:3], a[4:11], a[12:17], a[20:23] cbsz:1 blgp:3624 625 626// bf6 x bf6627// GFX950: v_mfma_f32_16x16x128_f8f6f4 v[0:3], v[4:9], v[12:17], v[20:23] cbsz:3 blgp:3 ; encoding: [0x00,0x03,0xad,0xd3,0x04,0x19,0x52,0x64]628v_mfma_f32_16x16x128_f8f6f4 v[0:3], v[4:9], v[12:17], v[20:23] cbsz:3 blgp:3629 630// GFX950: v_mfma_f32_16x16x128_f8f6f4 a[0:3], a[4:9], a[12:17], a[20:23] cbsz:3 blgp:3 ; encoding: [0x00,0x83,0xad,0xd3,0x04,0x19,0x52,0x7c]631v_mfma_f32_16x16x128_f8f6f4 a[0:3], a[4:9], a[12:17], a[20:23] cbsz:3 blgp:3632 633// bf6 x fp6634// GFX950: v_mfma_f32_16x16x128_f8f6f4 v[0:3], v[4:9], v[12:17], v[20:23] cbsz:3 blgp:2 ; encoding: [0x00,0x03,0xad,0xd3,0x04,0x19,0x52,0x44]635v_mfma_f32_16x16x128_f8f6f4 v[0:3], v[4:9], v[12:17], v[20:23] cbsz:3 blgp:2636 637// GFX950: v_mfma_f32_16x16x128_f8f6f4 a[0:3], a[4:9], a[12:17], a[20:23] cbsz:3 blgp:2 ; encoding: [0x00,0x83,0xad,0xd3,0x04,0x19,0x52,0x5c]638v_mfma_f32_16x16x128_f8f6f4 a[0:3], a[4:9], a[12:17], a[20:23] cbsz:3 blgp:2639 640 641// fp6 x bf6642// GFX950: v_mfma_f32_16x16x128_f8f6f4 v[0:3], v[4:9], v[12:17], v[20:23] cbsz:2 blgp:3 ; encoding: [0x00,0x02,0xad,0xd3,0x04,0x19,0x52,0x64]643v_mfma_f32_16x16x128_f8f6f4 v[0:3], v[4:9], v[12:17], v[20:23] cbsz:2 blgp:3644 645// GFX950: v_mfma_f32_16x16x128_f8f6f4 a[0:3], a[4:9], a[12:17], a[20:23] cbsz:2 blgp:3 ; encoding: [0x00,0x82,0xad,0xd3,0x04,0x19,0x52,0x7c]646v_mfma_f32_16x16x128_f8f6f4 a[0:3], a[4:9], a[12:17], a[20:23] cbsz:2 blgp:3647 648 649// fp6 x fp4650// GFX950: v_mfma_f32_16x16x128_f8f6f4 v[0:3], v[4:9], v[12:15], v[20:23] cbsz:2 blgp:4 ; encoding: [0x00,0x02,0xad,0xd3,0x04,0x19,0x52,0x84]651v_mfma_f32_16x16x128_f8f6f4 v[0:3], v[4:9], v[12:15], v[20:23] cbsz:2 blgp:4652 653// GFX950: v_mfma_f32_16x16x128_f8f6f4 a[0:3], a[4:9], a[12:15], a[20:23] cbsz:2 blgp:4 ; encoding: [0x00,0x82,0xad,0xd3,0x04,0x19,0x52,0x9c]654v_mfma_f32_16x16x128_f8f6f4 a[0:3], a[4:9], a[12:15], a[20:23] cbsz:2 blgp:4655 656 657// fp4 x fp6658// GFX950: v_mfma_f32_16x16x128_f8f6f4 v[0:3], v[4:7], v[12:17], v[20:23] cbsz:4 blgp:2 ; encoding: [0x00,0x04,0xad,0xd3,0x04,0x19,0x52,0x44]659v_mfma_f32_16x16x128_f8f6f4 v[0:3], v[4:7], v[12:17], v[20:23] cbsz:4 blgp:2660 661// GFX950: v_mfma_f32_16x16x128_f8f6f4 a[0:3], a[4:7], a[12:17], a[20:23] cbsz:4 blgp:2 ; encoding: [0x00,0x84,0xad,0xd3,0x04,0x19,0x52,0x5c]662v_mfma_f32_16x16x128_f8f6f4 a[0:3], a[4:7], a[12:17], a[20:23] cbsz:4 blgp:2663 664 665// fp4 x fp4666// GFX950: v_mfma_f32_16x16x128_f8f6f4 v[0:3], v[4:7], v[12:15], v[20:23] cbsz:4 blgp:4 ; encoding: [0x00,0x04,0xad,0xd3,0x04,0x19,0x52,0x84]667v_mfma_f32_16x16x128_f8f6f4 v[0:3], v[4:7], v[12:15], v[20:23] cbsz:4 blgp:4668 669// GFX950: v_mfma_f32_16x16x128_f8f6f4 a[0:3], a[4:7], a[12:15], a[20:23] cbsz:4 blgp:4 ; encoding: [0x00,0x84,0xad,0xd3,0x04,0x19,0x52,0x9c]670v_mfma_f32_16x16x128_f8f6f4 a[0:3], a[4:7], a[12:15], a[20:23] cbsz:4 blgp:4671 672//===----------------------------------------------------------------------===//673// v_mfma_f32_32x32x64_f8f6f4 with appropriate register widths674//===----------------------------------------------------------------------===//675 676// bf8 x fp4677// GFX950: v_mfma_f32_32x32x64_f8f6f4 v[0:15], v[4:11], v[12:15], v[16:31] cbsz:1 blgp:4 ; encoding: [0x00,0x01,0xae,0xd3,0x04,0x19,0x42,0x84]678v_mfma_f32_32x32x64_f8f6f4 v[0:15], v[4:11], v[12:15], v[16:31] cbsz:1 blgp:4679 680// GFX950: v_mfma_f32_32x32x64_f8f6f4 a[0:15], a[4:11], a[12:15], a[16:31] cbsz:1 blgp:4 ; encoding: [0x00,0x81,0xae,0xd3,0x04,0x19,0x42,0x9c]681v_mfma_f32_32x32x64_f8f6f4 a[0:15], a[4:11], a[12:15], a[16:31] cbsz:1 blgp:4682 683 684// fp4 x bf8685// GFX950: v_mfma_f32_32x32x64_f8f6f4 v[0:15], v[4:7], v[8:15], v[16:31] cbsz:4 blgp:1 ; encoding: [0x00,0x04,0xae,0xd3,0x04,0x11,0x42,0x24]686v_mfma_f32_32x32x64_f8f6f4 v[0:15], v[4:7], v[8:15], v[16:31] cbsz:4 blgp:1687 688// GFX950: v_mfma_f32_32x32x64_f8f6f4 a[0:15], a[4:7], a[8:15], a[16:31] cbsz:4 blgp:1 ; encoding: [0x00,0x84,0xae,0xd3,0x04,0x11,0x42,0x3c]689v_mfma_f32_32x32x64_f8f6f4 a[0:15], a[4:7], a[8:15], a[16:31] cbsz:4 blgp:1690 691 692// bf6 x bf8693// GFX950: v_mfma_f32_32x32x64_f8f6f4 v[0:15], v[4:9], v[12:19], v[16:31] cbsz:3 blgp:1 ; encoding: [0x00,0x03,0xae,0xd3,0x04,0x19,0x42,0x24]694v_mfma_f32_32x32x64_f8f6f4 v[0:15], v[4:9], v[12:19], v[16:31] cbsz:3 blgp:1695 696// GFX950: v_mfma_f32_32x32x64_f8f6f4 a[0:15], a[4:9], a[12:19], a[16:31] cbsz:3 blgp:1 ; encoding: [0x00,0x83,0xae,0xd3,0x04,0x19,0x42,0x3c]697v_mfma_f32_32x32x64_f8f6f4 a[0:15], a[4:9], a[12:19], a[16:31] cbsz:3 blgp:1698 699// bf8 x bf6700// GFX950: v_mfma_f32_32x32x64_f8f6f4 v[0:15], v[4:11], v[12:17], v[16:31] cbsz:1 blgp:3 ; encoding: [0x00,0x01,0xae,0xd3,0x04,0x19,0x42,0x64]701v_mfma_f32_32x32x64_f8f6f4 v[0:15], v[4:11], v[12:17], v[16:31] cbsz:1 blgp:3702 703// GFX950: v_mfma_f32_32x32x64_f8f6f4 a[0:15], a[4:11], a[12:17], a[16:31] cbsz:1 blgp:3 ; encoding: [0x00,0x81,0xae,0xd3,0x04,0x19,0x42,0x7c]704v_mfma_f32_32x32x64_f8f6f4 a[0:15], a[4:11], a[12:17], a[16:31] cbsz:1 blgp:3705 706 707// bf6 x bf6708// GFX950: v_mfma_f32_32x32x64_f8f6f4 v[0:15], v[4:9], v[12:17], v[16:31] cbsz:3 blgp:3 ; encoding: [0x00,0x03,0xae,0xd3,0x04,0x19,0x42,0x64]709v_mfma_f32_32x32x64_f8f6f4 v[0:15], v[4:9], v[12:17], v[16:31] cbsz:3 blgp:3710 711// GFX950: v_mfma_f32_32x32x64_f8f6f4 a[0:15], a[4:9], a[12:17], a[16:31] cbsz:3 blgp:3 ; encoding: [0x00,0x83,0xae,0xd3,0x04,0x19,0x42,0x7c]712v_mfma_f32_32x32x64_f8f6f4 a[0:15], a[4:9], a[12:17], a[16:31] cbsz:3 blgp:3713 714 715// bf6 x fp6716// GFX950: v_mfma_f32_32x32x64_f8f6f4 v[0:15], v[4:9], v[12:17], v[16:31] cbsz:3 blgp:2 ; encoding: [0x00,0x03,0xae,0xd3,0x04,0x19,0x42,0x44]717v_mfma_f32_32x32x64_f8f6f4 v[0:15], v[4:9], v[12:17], v[16:31] cbsz:3 blgp:2718 719// GFX950: v_mfma_f32_32x32x64_f8f6f4 a[0:15], a[4:9], a[12:17], a[16:31] cbsz:3 blgp:2 ; encoding: [0x00,0x83,0xae,0xd3,0x04,0x19,0x42,0x5c]720v_mfma_f32_32x32x64_f8f6f4 a[0:15], a[4:9], a[12:17], a[16:31] cbsz:3 blgp:2721 722 723// fp6 x bf6724// GFX950: v_mfma_f32_32x32x64_f8f6f4 v[0:15], v[4:9], v[12:17], v[16:31] cbsz:2 blgp:3 ; encoding: [0x00,0x02,0xae,0xd3,0x04,0x19,0x42,0x64]725v_mfma_f32_32x32x64_f8f6f4 v[0:15], v[4:9], v[12:17], v[16:31] cbsz:2 blgp:3726 727// GFX950: v_mfma_f32_32x32x64_f8f6f4 a[0:15], a[4:9], a[12:17], a[16:31] cbsz:2 blgp:3 ; encoding: [0x00,0x82,0xae,0xd3,0x04,0x19,0x42,0x7c]728v_mfma_f32_32x32x64_f8f6f4 a[0:15], a[4:9], a[12:17], a[16:31] cbsz:2 blgp:3729 730 731// fp6 x fp4732// GFX950: v_mfma_f32_32x32x64_f8f6f4 v[0:15], v[4:9], v[12:15], v[16:31] cbsz:2 blgp:4 ; encoding: [0x00,0x02,0xae,0xd3,0x04,0x19,0x42,0x84]733v_mfma_f32_32x32x64_f8f6f4 v[0:15], v[4:9], v[12:15], v[16:31] cbsz:2 blgp:4734 735// GFX950: v_mfma_f32_32x32x64_f8f6f4 a[0:15], a[4:9], a[12:15], a[16:31] cbsz:2 blgp:4 ; encoding: [0x00,0x82,0xae,0xd3,0x04,0x19,0x42,0x9c]736v_mfma_f32_32x32x64_f8f6f4 a[0:15], a[4:9], a[12:15], a[16:31] cbsz:2 blgp:4737 738 739// fp4 x fp6740// GFX950: v_mfma_f32_32x32x64_f8f6f4 v[0:15], v[4:7], v[12:17], v[16:31] cbsz:4 blgp:2 ; encoding: [0x00,0x04,0xae,0xd3,0x04,0x19,0x42,0x44]741v_mfma_f32_32x32x64_f8f6f4 v[0:15], v[4:7], v[12:17], v[16:31] cbsz:4 blgp:2742 743// GFX950: v_mfma_f32_32x32x64_f8f6f4 a[0:15], a[4:7], a[12:17], a[16:31] cbsz:4 blgp:2 ; encoding: [0x00,0x84,0xae,0xd3,0x04,0x19,0x42,0x5c]744v_mfma_f32_32x32x64_f8f6f4 a[0:15], a[4:7], a[12:17], a[16:31] cbsz:4 blgp:2745 746 747// fp4 x fp4748// GFX950: v_mfma_f32_32x32x64_f8f6f4 v[0:15], v[4:7], v[12:15], v[16:31] cbsz:4 blgp:4 ; encoding: [0x00,0x04,0xae,0xd3,0x04,0x19,0x42,0x84]749v_mfma_f32_32x32x64_f8f6f4 v[0:15], v[4:7], v[12:15], v[16:31] cbsz:4 blgp:4750 751// GFX950: v_mfma_f32_32x32x64_f8f6f4 a[0:15], a[4:7], a[12:15], a[16:31] cbsz:4 blgp:4 ; encoding: [0x00,0x84,0xae,0xd3,0x04,0x19,0x42,0x9c]752v_mfma_f32_32x32x64_f8f6f4 a[0:15], a[4:7], a[12:15], a[16:31] cbsz:4 blgp:4753 754//===----------------------------------------------------------------------===//755// v_mfma_scale_f32_16x16x128_f8f6f4 corrected widths756//===----------------------------------------------------------------------===//757 758// fp8 x fp8759// GFX950: v_mfma_scale_f32_16x16x128_f8f6f4 v[0:3], v[4:11], v[12:19], v[20:23], v24, v25 op_sel_hi:[0,0,0] ; encoding: [0x00,0x00,0xac,0xd3,0x18,0x33,0x02,0x00,0x00,0x08,0xad,0xd3,0x04,0x19,0x52,0x04]760v_mfma_scale_f32_16x16x128_f8f6f4 v[0:3], v[4:11], v[12:19], v[20:23], v24, v25 cbsz:0 blgp:0761 762// GFX950: v_mfma_scale_f32_16x16x128_f8f6f4 a[0:3], a[4:11], a[12:19], a[20:23], v24, v25 op_sel_hi:[0,0,0] ; encoding: [0x00,0x00,0xac,0xd3,0x18,0x33,0x02,0x00,0x00,0x88,0xad,0xd3,0x04,0x19,0x52,0x1c]763v_mfma_scale_f32_16x16x128_f8f6f4 a[0:3], a[4:11], a[12:19], a[20:23], v24, v25 cbsz:0 blgp:0764 765 766// bf8 x fp8767// GFX950: v_mfma_scale_f32_16x16x128_f8f6f4 v[0:3], v[4:11], v[12:19], v[20:23], v24, v25 op_sel_hi:[0,0,0] cbsz:1 ; encoding: [0x00,0x00,0xac,0xd3,0x18,0x33,0x02,0x00,0x00,0x09,0xad,0xd3,0x04,0x19,0x52,0x04]768v_mfma_scale_f32_16x16x128_f8f6f4 v[0:3], v[4:11], v[12:19], v[20:23], v24, v25 cbsz:1 blgp:0769 770// GFX950: v_mfma_scale_f32_16x16x128_f8f6f4 a[0:3], a[4:11], a[12:19], a[20:23], v24, v25 op_sel_hi:[0,0,0] cbsz:1 ; encoding: [0x00,0x00,0xac,0xd3,0x18,0x33,0x02,0x00,0x00,0x89,0xad,0xd3,0x04,0x19,0x52,0x1c]771v_mfma_scale_f32_16x16x128_f8f6f4 a[0:3], a[4:11], a[12:19], a[20:23], v24, v25 cbsz:1 blgp:0772 773// fp8 x bf8774// GFX950: v_mfma_scale_f32_16x16x128_f8f6f4 v[0:3], v[4:11], v[12:19], v[20:23], v24, v25 op_sel_hi:[0,0,0] blgp:1 ; encoding: [0x00,0x00,0xac,0xd3,0x18,0x33,0x02,0x00,0x00,0x08,0xad,0xd3,0x04,0x19,0x52,0x24]775v_mfma_scale_f32_16x16x128_f8f6f4 v[0:3], v[4:11], v[12:19], v[20:23], v24, v25 cbsz:0 blgp:1776 777// GFX950: v_mfma_scale_f32_16x16x128_f8f6f4 a[0:3], a[4:11], a[12:19], a[20:23], v24, v25 op_sel_hi:[0,0,0] blgp:1 ; encoding: [0x00,0x00,0xac,0xd3,0x18,0x33,0x02,0x00,0x00,0x88,0xad,0xd3,0x04,0x19,0x52,0x3c]778v_mfma_scale_f32_16x16x128_f8f6f4 a[0:3], a[4:11], a[12:19], a[20:23], v24, v25 cbsz:0 blgp:1779 780 781// bf8 x fp4782// GFX950: v_mfma_scale_f32_16x16x128_f8f6f4 v[0:3], v[4:11], v[12:15], v[20:23], v24, v25 op_sel_hi:[0,0,0] cbsz:1 blgp:4 ; encoding: [0x00,0x00,0xac,0xd3,0x18,0x33,0x02,0x00,0x00,0x09,0xad,0xd3,0x04,0x19,0x52,0x84]783v_mfma_scale_f32_16x16x128_f8f6f4 v[0:3], v[4:11], v[12:15], v[20:23], v24, v25 cbsz:1 blgp:4784 785// GFX950: v_mfma_scale_f32_16x16x128_f8f6f4 a[0:3], a[4:11], a[12:15], a[20:23], v24, v25 op_sel_hi:[0,0,0] cbsz:1 blgp:4 ; encoding: [0x00,0x00,0xac,0xd3,0x18,0x33,0x02,0x00,0x00,0x89,0xad,0xd3,0x04,0x19,0x52,0x9c]786v_mfma_scale_f32_16x16x128_f8f6f4 a[0:3], a[4:11], a[12:15], a[20:23], v24, v25 cbsz:1 blgp:4787 788// fp4 x bf8789// GFX950: v_mfma_scale_f32_16x16x128_f8f6f4 v[0:3], v[4:7], v[12:19], v[20:23], v24, v25 op_sel_hi:[0,0,0] cbsz:4 blgp:1 ; encoding: [0x00,0x00,0xac,0xd3,0x18,0x33,0x02,0x00,0x00,0x0c,0xad,0xd3,0x04,0x19,0x52,0x24]790v_mfma_scale_f32_16x16x128_f8f6f4 v[0:3], v[4:7], v[12:19], v[20:23], v24, v25 cbsz:4 blgp:1791 792// GFX950: v_mfma_scale_f32_16x16x128_f8f6f4 a[0:3], a[4:7], a[12:19], a[20:23], v24, v25 op_sel_hi:[0,0,0] cbsz:4 blgp:1 ; encoding: [0x00,0x00,0xac,0xd3,0x18,0x33,0x02,0x00,0x00,0x8c,0xad,0xd3,0x04,0x19,0x52,0x3c]793v_mfma_scale_f32_16x16x128_f8f6f4 a[0:3], a[4:7], a[12:19], a[20:23], v24, v25 cbsz:4 blgp:1794 795 796// bf6 x bf8797// GFX950: v_mfma_scale_f32_16x16x128_f8f6f4 v[0:3], v[4:9], v[12:19], v[20:23], v24, v25 op_sel_hi:[0,0,0] cbsz:3 blgp:1 ; encoding: [0x00,0x00,0xac,0xd3,0x18,0x33,0x02,0x00,0x00,0x0b,0xad,0xd3,0x04,0x19,0x52,0x24]798v_mfma_scale_f32_16x16x128_f8f6f4 v[0:3], v[4:9], v[12:19], v[20:23], v24, v25 cbsz:3 blgp:1799 800// GFX950: v_mfma_scale_f32_16x16x128_f8f6f4 a[0:3], a[4:9], a[12:19], a[20:23], v24, v25 op_sel_hi:[0,0,0] cbsz:3 blgp:1 ; encoding: [0x00,0x00,0xac,0xd3,0x18,0x33,0x02,0x00,0x00,0x8b,0xad,0xd3,0x04,0x19,0x52,0x3c]801v_mfma_scale_f32_16x16x128_f8f6f4 a[0:3], a[4:9], a[12:19], a[20:23], v24, v25 cbsz:3 blgp:1802 803 804// bf8 x bf6805// GFX950: v_mfma_scale_f32_16x16x128_f8f6f4 v[0:3], v[4:11], v[12:17], v[20:23], v24, v25 op_sel_hi:[0,0,0] cbsz:1 blgp:3 ; encoding: [0x00,0x00,0xac,0xd3,0x18,0x33,0x02,0x00,0x00,0x09,0xad,0xd3,0x04,0x19,0x52,0x64]806v_mfma_scale_f32_16x16x128_f8f6f4 v[0:3], v[4:11], v[12:17], v[20:23], v24, v25 cbsz:1 blgp:3807 808// GFX950: v_mfma_scale_f32_16x16x128_f8f6f4 a[0:3], a[4:11], a[12:17], a[20:23], v24, v25 op_sel_hi:[0,0,0] cbsz:1 blgp:3 ; encoding: [0x00,0x00,0xac,0xd3,0x18,0x33,0x02,0x00,0x00,0x89,0xad,0xd3,0x04,0x19,0x52,0x7c]809v_mfma_scale_f32_16x16x128_f8f6f4 a[0:3], a[4:11], a[12:17], a[20:23], v24, v25 cbsz:1 blgp:3810 811 812// bf6 x bf6813// GFX950: v_mfma_scale_f32_16x16x128_f8f6f4 v[0:3], v[4:9], v[12:17], v[20:23], v24, v25 op_sel_hi:[0,0,0] cbsz:3 blgp:3 ; encoding: [0x00,0x00,0xac,0xd3,0x18,0x33,0x02,0x00,0x00,0x0b,0xad,0xd3,0x04,0x19,0x52,0x64]814v_mfma_scale_f32_16x16x128_f8f6f4 v[0:3], v[4:9], v[12:17], v[20:23], v24, v25 cbsz:3 blgp:3815 816// GFX950: v_mfma_scale_f32_16x16x128_f8f6f4 a[0:3], a[4:9], a[12:17], a[20:23], v24, v25 op_sel_hi:[0,0,0] cbsz:3 blgp:3 ; encoding: [0x00,0x00,0xac,0xd3,0x18,0x33,0x02,0x00,0x00,0x8b,0xad,0xd3,0x04,0x19,0x52,0x7c]817v_mfma_scale_f32_16x16x128_f8f6f4 a[0:3], a[4:9], a[12:17], a[20:23], v24, v25 cbsz:3 blgp:3818 819 820// bf6 x fp6821// GFX950: v_mfma_scale_f32_16x16x128_f8f6f4 v[0:3], v[4:9], v[12:17], v[20:23], v24, v25 op_sel_hi:[0,0,0] cbsz:3 blgp:2 ; encoding: [0x00,0x00,0xac,0xd3,0x18,0x33,0x02,0x00,0x00,0x0b,0xad,0xd3,0x04,0x19,0x52,0x44]822v_mfma_scale_f32_16x16x128_f8f6f4 v[0:3], v[4:9], v[12:17], v[20:23], v24, v25 cbsz:3 blgp:2823 824// GFX950: v_mfma_scale_f32_16x16x128_f8f6f4 a[0:3], a[4:9], a[12:17], a[20:23], v24, v25 op_sel_hi:[0,0,0] cbsz:3 blgp:2 ; encoding: [0x00,0x00,0xac,0xd3,0x18,0x33,0x02,0x00,0x00,0x8b,0xad,0xd3,0x04,0x19,0x52,0x5c]825v_mfma_scale_f32_16x16x128_f8f6f4 a[0:3], a[4:9], a[12:17], a[20:23], v24, v25 cbsz:3 blgp:2826 827 828// fp6 x bf6829// GFX950: v_mfma_scale_f32_16x16x128_f8f6f4 v[0:3], v[4:9], v[12:17], v[20:23], v24, v25 op_sel_hi:[0,0,0] cbsz:2 blgp:3 ; encoding: [0x00,0x00,0xac,0xd3,0x18,0x33,0x02,0x00,0x00,0x0a,0xad,0xd3,0x04,0x19,0x52,0x64830v_mfma_scale_f32_16x16x128_f8f6f4 v[0:3], v[4:9], v[12:17], v[20:23], v24, v25 cbsz:2 blgp:3831 832// GFX950: v_mfma_scale_f32_16x16x128_f8f6f4 a[0:3], a[4:9], a[12:17], a[20:23], v24, v25 op_sel_hi:[0,0,0] cbsz:2 blgp:3 ; encoding: [0x00,0x00,0xac,0xd3,0x18,0x33,0x02,0x00,0x00,0x8a,0xad,0xd3,0x04,0x19,0x52,0x7c]833v_mfma_scale_f32_16x16x128_f8f6f4 a[0:3], a[4:9], a[12:17], a[20:23], v24, v25 cbsz:2 blgp:3834 835 836// fp6 x fp4837// GFX950: v_mfma_scale_f32_16x16x128_f8f6f4 v[0:3], v[4:9], v[12:15], v[20:23], v24, v25 op_sel_hi:[0,0,0] cbsz:2 blgp:4 ; encoding: [0x00,0x00,0xac,0xd3,0x18,0x33,0x02,0x00,0x00,0x0a,0xad,0xd3,0x04,0x19,0x52,0x84]838v_mfma_scale_f32_16x16x128_f8f6f4 v[0:3], v[4:9], v[12:15], v[20:23], v24, v25 cbsz:2 blgp:4839 840// GFX950: v_mfma_scale_f32_16x16x128_f8f6f4 a[0:3], a[4:9], a[12:15], a[20:23], v24, v25 op_sel_hi:[0,0,0] cbsz:2 blgp:4 ; encoding: [0x00,0x00,0xac,0xd3,0x18,0x33,0x02,0x00,0x00,0x8a,0xad,0xd3,0x04,0x19,0x52,0x9c]841v_mfma_scale_f32_16x16x128_f8f6f4 a[0:3], a[4:9], a[12:15], a[20:23], v24, v25 cbsz:2 blgp:4842 843// fp4 x fp6844// GFX950: v_mfma_scale_f32_16x16x128_f8f6f4 v[0:3], v[4:7], v[12:17], v[20:23], v24, v25 op_sel_hi:[0,0,0] cbsz:4 blgp:2 ; encoding: [0x00,0x00,0xac,0xd3,0x18,0x33,0x02,0x00,0x00,0x0c,0xad,0xd3,0x04,0x19,0x52,0x44]845v_mfma_scale_f32_16x16x128_f8f6f4 v[0:3], v[4:7], v[12:17], v[20:23], v24, v25 cbsz:4 blgp:2846 847// GFX950: v_mfma_scale_f32_16x16x128_f8f6f4 a[0:3], a[4:7], a[12:17], a[20:23], v24, v25 op_sel_hi:[0,0,0] cbsz:4 blgp:2 ; encoding: [0x00,0x00,0xac,0xd3,0x18,0x33,0x02,0x00,0x00,0x8c,0xad,0xd3,0x04,0x19,0x52,0x5c]848v_mfma_scale_f32_16x16x128_f8f6f4 a[0:3], a[4:7], a[12:17], a[20:23], v24, v25 cbsz:4 blgp:2849 850// fp4 x fp4851// GFX950: v_mfma_scale_f32_16x16x128_f8f6f4 v[0:3], v[4:7], v[12:15], v[20:23], v24, v25 op_sel_hi:[0,0,0] cbsz:4 blgp:4 ; encoding: [0x00,0x00,0xac,0xd3,0x18,0x33,0x02,0x00,0x00,0x0c,0xad,0xd3,0x04,0x19,0x52,0x84]852v_mfma_scale_f32_16x16x128_f8f6f4 v[0:3], v[4:7], v[12:15], v[20:23], v24, v25 cbsz:4 blgp:4853 854// GFX950: v_mfma_scale_f32_16x16x128_f8f6f4 a[0:3], a[4:7], a[12:15], a[20:23], v24, v25 op_sel_hi:[0,0,0] cbsz:4 blgp:4 ; encoding: [0x00,0x00,0xac,0xd3,0x18,0x33,0x02,0x00,0x00,0x8c,0xad,0xd3,0x04,0x19,0x52,0x9c]855v_mfma_scale_f32_16x16x128_f8f6f4 a[0:3], a[4:7], a[12:15], a[20:23], v24, v25 cbsz:4 blgp:4856 857 858//===----------------------------------------------------------------------===//859// v_mfma_scale_f32_32x32x64_f8f6f4 corrected widths860//===----------------------------------------------------------------------===//861 862// fp8 x fp8863// GFX950: v_mfma_scale_f32_32x32x64_f8f6f4 v[0:15], v[16:23], v[24:31], v[32:47], v48, v49 op_sel_hi:[0,0,0] ; encoding: [0x00,0x00,0xac,0xd3,0x30,0x63,0x02,0x00,0x00,0x08,0xae,0xd3,0x10,0x31,0x82,0x04]864v_mfma_scale_f32_32x32x64_f8f6f4 v[0:15], v[16:23], v[24:31], v[32:47], v48, v49 cbsz:0 blgp:0865 866// GFX950: v_mfma_scale_f32_32x32x64_f8f6f4 a[0:15], a[16:23], a[24:31], a[32:47], v48, v49 op_sel_hi:[0,0,0] ; encoding: [0x00,0x00,0xac,0xd3,0x30,0x63,0x02,0x00,0x00,0x88,0xae,0xd3,0x10,0x31,0x82,0x1c]867v_mfma_scale_f32_32x32x64_f8f6f4 a[0:15], a[16:23], a[24:31], a[32:47], v48, v49 cbsz:0 blgp:0868 869// bf8 x fp8870// GFX950: v_mfma_scale_f32_32x32x64_f8f6f4 v[0:15], v[16:23], v[24:31], v[32:47], v48, v49 op_sel_hi:[0,0,0] cbsz:1 ; encoding: [0x00,0x00,0xac,0xd3,0x30,0x63,0x02,0x00,0x00,0x09,0xae,0xd3,0x10,0x31,0x82,0x04]871v_mfma_scale_f32_32x32x64_f8f6f4 v[0:15], v[16:23], v[24:31], v[32:47], v48, v49 cbsz:1 blgp:0872 873// GFX950: v_mfma_scale_f32_32x32x64_f8f6f4 a[0:15], a[16:23], a[24:31], a[32:47], v48, v49 op_sel_hi:[0,0,0] cbsz:1 ; encoding: [0x00,0x00,0xac,0xd3,0x30,0x63,0x02,0x00,0x00,0x89,0xae,0xd3,0x10,0x31,0x82,0x1c]874v_mfma_scale_f32_32x32x64_f8f6f4 a[0:15], a[16:23], a[24:31], a[32:47], v48, v49 cbsz:1 blgp:0875 876 877// fp8 x bf8878// GFX950: v_mfma_scale_f32_32x32x64_f8f6f4 v[0:15], v[16:23], v[24:31], v[32:47], v48, v49 op_sel_hi:[0,0,0] blgp:1 ; encoding: [0x00,0x00,0xac,0xd3,0x30,0x63,0x02,0x00,0x00,0x08,0xae,0xd3,0x10,0x31,0x82,0x24]879v_mfma_scale_f32_32x32x64_f8f6f4 v[0:15], v[16:23], v[24:31], v[32:47], v48, v49 cbsz:0 blgp:1880 881// GFX950: v_mfma_scale_f32_32x32x64_f8f6f4 a[0:15], a[16:23], a[24:31], a[32:47], v48, v49 op_sel_hi:[0,0,0] blgp:1 ; encoding: [0x00,0x00,0xac,0xd3,0x30,0x63,0x02,0x00,0x00,0x88,0xae,0xd3,0x10,0x31,0x82,0x3c]882v_mfma_scale_f32_32x32x64_f8f6f4 a[0:15], a[16:23], a[24:31], a[32:47], v48, v49 cbsz:0 blgp:1883 884// bf8 x fp4885// GFX950: v_mfma_scale_f32_32x32x64_f8f6f4 v[0:15], v[16:23], v[24:27], v[32:47], v48, v49 op_sel_hi:[0,0,0] cbsz:1 blgp:4 ; encoding: [0x00,0x00,0xac,0xd3,0x30,0x63,0x02,0x00,0x00,0x09,0xae,0xd3,0x10,0x31,0x82,0x84]886v_mfma_scale_f32_32x32x64_f8f6f4 v[0:15], v[16:23], v[24:27], v[32:47], v48, v49 cbsz:1 blgp:4887 888// GFX950: v_mfma_scale_f32_32x32x64_f8f6f4 a[0:15], a[16:23], a[24:27], a[32:47], v48, v49 op_sel_hi:[0,0,0] cbsz:1 blgp:4 ; encoding: [0x00,0x00,0xac,0xd3,0x30,0x63,0x02,0x00,0x00,0x89,0xae,0xd3,0x10,0x31,0x82,0x9c]889v_mfma_scale_f32_32x32x64_f8f6f4 a[0:15], a[16:23], a[24:27], a[32:47], v48, v49 cbsz:1 blgp:4890 891 892// fp4 x bf8893// GFX950: v_mfma_scale_f32_32x32x64_f8f6f4 v[0:15], v[16:19], v[24:31], v[32:47], v48, v49 op_sel_hi:[0,0,0] cbsz:4 blgp:1 ; encoding: [0x00,0x00,0xac,0xd3,0x30,0x63,0x02,0x00,0x00,0x0c,0xae,0xd3,0x10,0x31,0x82,0x24]894v_mfma_scale_f32_32x32x64_f8f6f4 v[0:15], v[16:19], v[24:31], v[32:47], v48, v49 cbsz:4 blgp:1895 896// GFX950: v_mfma_scale_f32_32x32x64_f8f6f4 a[0:15], a[16:19], a[24:31], a[32:47], v48, v49 op_sel_hi:[0,0,0] cbsz:4 blgp:1 ; encoding: [0x00,0x00,0xac,0xd3,0x30,0x63,0x02,0x00,0x00,0x8c,0xae,0xd3,0x10,0x31,0x82,0x3c]897v_mfma_scale_f32_32x32x64_f8f6f4 a[0:15], a[16:19], a[24:31], a[32:47], v48, v49 cbsz:4 blgp:1898 899 900// bf6 x bf8901// GFX950: v_mfma_scale_f32_32x32x64_f8f6f4 v[0:15], v[16:21], v[24:31], v[32:47], v48, v49 op_sel_hi:[0,0,0] cbsz:3 blgp:1 ; encoding: [0x00,0x00,0xac,0xd3,0x30,0x63,0x02,0x00,0x00,0x0b,0xae,0xd3,0x10,0x31,0x82,0x24]902v_mfma_scale_f32_32x32x64_f8f6f4 v[0:15], v[16:21], v[24:31], v[32:47], v48, v49 cbsz:3 blgp:1903 904// GFX950: v_mfma_scale_f32_32x32x64_f8f6f4 a[0:15], a[16:21], a[24:31], a[32:47], v48, v49 op_sel_hi:[0,0,0] cbsz:3 blgp:1 ; encoding: [0x00,0x00,0xac,0xd3,0x30,0x63,0x02,0x00,0x00,0x8b,0xae,0xd3,0x10,0x31,0x82,0x3c]905v_mfma_scale_f32_32x32x64_f8f6f4 a[0:15], a[16:21], a[24:31], a[32:47], v48, v49 cbsz:3 blgp:1906 907 908// bf8 x bf6909// GFX950: v_mfma_scale_f32_32x32x64_f8f6f4 v[0:15], v[16:23], v[24:29], v[32:47], v48, v49 op_sel_hi:[0,0,0] cbsz:1 blgp:3 ; encoding: [0x00,0x00,0xac,0xd3,0x30,0x63,0x02,0x00,0x00,0x09,0xae,0xd3,0x10,0x31,0x82,0x64]910v_mfma_scale_f32_32x32x64_f8f6f4 v[0:15], v[16:23], v[24:29], v[32:47], v48, v49 cbsz:1 blgp:3911 912// GFX950: v_mfma_scale_f32_32x32x64_f8f6f4 a[0:15], a[16:23], a[24:29], a[32:47], v48, v49 op_sel_hi:[0,0,0] cbsz:1 blgp:3 ; encoding: [0x00,0x00,0xac,0xd3,0x30,0x63,0x02,0x00,0x00,0x89,0xae,0xd3,0x10,0x31,0x82,0x7c]913v_mfma_scale_f32_32x32x64_f8f6f4 a[0:15], a[16:23], a[24:29], a[32:47], v48, v49 cbsz:1 blgp:3914 915// bf6 x bf6916// GFX950: v_mfma_scale_f32_32x32x64_f8f6f4 v[0:15], v[16:21], v[24:29], v[32:47], v48, v49 op_sel_hi:[0,0,0] cbsz:3 blgp:3 ; encoding: [0x00,0x00,0xac,0xd3,0x30,0x63,0x02,0x00,0x00,0x0b,0xae,0xd3,0x10,0x31,0x82,0x64]917v_mfma_scale_f32_32x32x64_f8f6f4 v[0:15], v[16:21], v[24:29], v[32:47], v48, v49 cbsz:3 blgp:3918 919// GFX950: v_mfma_scale_f32_32x32x64_f8f6f4 a[0:15], a[16:21], a[24:29], a[32:47], v48, v49 op_sel_hi:[0,0,0] cbsz:3 blgp:3 ; encoding: [0x00,0x00,0xac,0xd3,0x30,0x63,0x02,0x00,0x00,0x8b,0xae,0xd3,0x10,0x31,0x82,0x7c]920v_mfma_scale_f32_32x32x64_f8f6f4 a[0:15], a[16:21], a[24:29], a[32:47], v48, v49 cbsz:3 blgp:3921 922// bf6 x fp6923// GFX950: v_mfma_scale_f32_32x32x64_f8f6f4 v[0:15], v[16:21], v[24:29], v[32:47], v48, v49 op_sel_hi:[0,0,0] cbsz:3 blgp:2 ; encoding: [0x00,0x00,0xac,0xd3,0x30,0x63,0x02,0x00,0x00,0x0b,0xae,0xd3,0x10,0x31,0x82,0x44]924v_mfma_scale_f32_32x32x64_f8f6f4 v[0:15], v[16:21], v[24:29], v[32:47], v48, v49 cbsz:3 blgp:2925 926// GFX950: v_mfma_scale_f32_32x32x64_f8f6f4 a[0:15], a[16:21], a[24:29], a[32:47], v48, v49 op_sel_hi:[0,0,0] cbsz:3 blgp:2 ; encoding: [0x00,0x00,0xac,0xd3,0x30,0x63,0x02,0x00,0x00,0x8b,0xae,0xd3,0x10,0x31,0x82,0x5c]927v_mfma_scale_f32_32x32x64_f8f6f4 a[0:15], a[16:21], a[24:29], a[32:47], v48, v49 cbsz:3 blgp:2928 929 930// fp6 x bf6931// GFX950: v_mfma_scale_f32_32x32x64_f8f6f4 v[0:15], v[16:21], v[24:29], v[32:47], v48, v49 op_sel_hi:[0,0,0] cbsz:2 blgp:3 ; encoding: [0x00,0x00,0xac,0xd3,0x30,0x63,0x02,0x00,0x00,0x0a,0xae,0xd3,0x10,0x31,0x82,0x64]932v_mfma_scale_f32_32x32x64_f8f6f4 v[0:15], v[16:21], v[24:29], v[32:47], v48, v49 cbsz:2 blgp:3933 934// GFX950: v_mfma_scale_f32_32x32x64_f8f6f4 a[0:15], a[16:21], a[24:29], a[32:47], v48, v49 op_sel_hi:[0,0,0] cbsz:2 blgp:3 ; encoding: [0x00,0x00,0xac,0xd3,0x30,0x63,0x02,0x00,0x00,0x8a,0xae,0xd3,0x10,0x31,0x82,0x7c]935v_mfma_scale_f32_32x32x64_f8f6f4 a[0:15], a[16:21], a[24:29], a[32:47], v48, v49 cbsz:2 blgp:3936 937 938// fp6 x fp4939// GFX950: v_mfma_scale_f32_32x32x64_f8f6f4 v[0:15], v[16:21], v[24:27], v[32:47], v48, v49 op_sel_hi:[0,0,0] cbsz:2 blgp:4 ; encoding: [0x00,0x00,0xac,0xd3,0x30,0x63,0x02,0x00,0x00,0x0a,0xae,0xd3,0x10,0x31,0x82,0x84]940v_mfma_scale_f32_32x32x64_f8f6f4 v[0:15], v[16:21], v[24:27], v[32:47], v48, v49 cbsz:2 blgp:4941 942// GFX950: v_mfma_scale_f32_32x32x64_f8f6f4 a[0:15], a[16:21], a[24:27], a[32:47], v48, v49 op_sel_hi:[0,0,0] cbsz:2 blgp:4 ; encoding: [0x00,0x00,0xac,0xd3,0x30,0x63,0x02,0x00,0x00,0x8a,0xae,0xd3,0x10,0x31,0x82,0x9c]943v_mfma_scale_f32_32x32x64_f8f6f4 a[0:15], a[16:21], a[24:27], a[32:47], v48, v49 cbsz:2 blgp:4944 945 946// fp4 x fp6947// GFX950: v_mfma_scale_f32_32x32x64_f8f6f4 v[0:15], v[16:19], v[24:29], v[32:47], v48, v49 op_sel_hi:[0,0,0] cbsz:4 blgp:2 ; encoding: [0x00,0x00,0xac,0xd3,0x30,0x63,0x02,0x00,0x00,0x0c,0xae,0xd3,0x10,0x31,0x82,0x44]948v_mfma_scale_f32_32x32x64_f8f6f4 v[0:15], v[16:19], v[24:29], v[32:47], v48, v49 cbsz:4 blgp:2949 950// GFX950: v_mfma_scale_f32_32x32x64_f8f6f4 a[0:15], a[16:19], a[24:29], a[32:47], v48, v49 op_sel_hi:[0,0,0] cbsz:4 blgp:2 ; encoding: [0x00,0x00,0xac,0xd3,0x30,0x63,0x02,0x00,0x00,0x8c,0xae,0xd3,0x10,0x31,0x82,0x5c]951v_mfma_scale_f32_32x32x64_f8f6f4 a[0:15], a[16:19], a[24:29], a[32:47], v48, v49 cbsz:4 blgp:2952 953 954// fp4 x fp4955// GFX950: v_mfma_scale_f32_32x32x64_f8f6f4 v[0:15], v[16:19], v[24:27], v[32:47], v48, v49 op_sel_hi:[0,0,0] cbsz:4 blgp:4 ; encoding: [0x00,0x00,0xac,0xd3,0x30,0x63,0x02,0x00,0x00,0x0c,0xae,0xd3,0x10,0x31,0x82,0x84]956v_mfma_scale_f32_32x32x64_f8f6f4 v[0:15], v[16:19], v[24:27], v[32:47], v48, v49 cbsz:4 blgp:4957 958// GFX950: v_mfma_scale_f32_32x32x64_f8f6f4 a[0:15], a[16:19], a[24:27], a[32:47], v48, v49 op_sel_hi:[0,0,0] cbsz:4 blgp:4 ; encoding: [0x00,0x00,0xac,0xd3,0x30,0x63,0x02,0x00,0x00,0x8c,0xae,0xd3,0x10,0x31,0x82,0x9c]959v_mfma_scale_f32_32x32x64_f8f6f4 a[0:15], a[16:19], a[24:27], a[32:47], v48, v49 cbsz:4 blgp:4960 961//===----------------------------------------------------------------------===//962// v_mfma_i32_16x16x64_i8963//===----------------------------------------------------------------------===//964 965// GFX950: v_mfma_i32_16x16x64_i8 v[0:3], v[0:3], v[0:3], v[0:3] ; encoding: [0x00,0x00,0xb6,0xd3,0x00,0x01,0x02,0x04]966// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU967v_mfma_i32_16x16x64_i8 v[0:3], v[0:3], v[0:3], v[0:3]968 969// GFX950: v_mfma_i32_16x16x64_i8 v[0:3], v[0:3], v[0:3], v[0:3] ; encoding: [0x00,0x00,0xb6,0xd3,0x00,0x01,0x02,0x04]970// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU971v_mfma_i32_16x16x64i8 v[0:3], v[0:3], v[0:3], v[0:3]972 973// GFX950: v_mfma_i32_16x16x64_i8 a[0:3], a[0:3], a[0:3], a[0:3] ; encoding: [0x00,0x80,0xb6,0xd3,0x00,0x01,0x02,0x1c]974// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU975v_mfma_i32_16x16x64_i8 a[0:3], a[0:3], a[0:3], a[0:3]976 977// GFX950: v_mfma_i32_16x16x64_i8 a[0:3], a[0:3], a[0:3], a[0:3] ; encoding: [0x00,0x80,0xb6,0xd3,0x00,0x01,0x02,0x1c]978// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU979v_mfma_i32_16x16x64i8 a[0:3], a[0:3], a[0:3], a[0:3]980 981// GFX950: v_mfma_i32_16x16x64_i8 v[0:3], a[0:3], v[0:3], 1.0 ; encoding: [0x00,0x00,0xb6,0xd3,0x00,0x01,0xca,0x0b]982// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU983v_mfma_i32_16x16x64_i8 v[0:3], a[0:3], v[0:3], 1.0984 985// GFX950: v_mfma_i32_16x16x64_i8 a[0:3], v[0:3], a[0:3], 1.0 ; encoding: [0x00,0x80,0xb6,0xd3,0x00,0x01,0xca,0x13]986// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU987v_mfma_i32_16x16x64_i8 a[0:3], v[0:3], a[0:3], 1.0988 989// GFX950: v_mfma_i32_16x16x64_i8 v[0:3], v[0:3], v[0:3], v[0:3] blgp:5 ; encoding: [0x00,0x00,0xb6,0xd3,0x00,0x01,0x02,0xa4]990// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU991v_mfma_i32_16x16x64_i8 v[0:3], v[0:3], v[0:3], v[0:3] blgp:5992 993// GFX950: v_mfma_i32_16x16x64_i8 a[0:3], a[0:3], a[0:3], a[0:3] blgp:1 ; encoding: [0x00,0x80,0xb6,0xd3,0x00,0x01,0x02,0x3c]994// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU995v_mfma_i32_16x16x64_i8 a[0:3], a[0:3], a[0:3], a[0:3] blgp:1996 997// GFX950: v_mfma_i32_16x16x64_i8 a[0:3], a[0:3], a[0:3], a[0:3] cbsz:3 ; encoding: [0x00,0x83,0xb6,0xd3,0x00,0x01,0x02,0x1c]998// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU999v_mfma_i32_16x16x64_i8 a[0:3], a[0:3], a[0:3], a[0:3] cbsz:31000 1001// GFX950: v_mfma_i32_16x16x64_i8 a[0:3], a[0:3], a[0:3], a[0:3] abid:1 ; encoding: [0x00,0x88,0xb6,0xd3,0x00,0x01,0x02,0x1c]1002// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU1003v_mfma_i32_16x16x64_i8 a[0:3], a[0:3], a[0:3], a[0:3] abid:11004 1005// GFX950: v_mfma_i32_16x16x64_i8 a[0:3], a[0:3], a[0:3], a[0:3] cbsz:3 abid:1 ; encoding: [0x00,0x8b,0xb6,0xd3,0x00,0x01,0x02,0x1c]1006// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU1007v_mfma_i32_16x16x64_i8 a[0:3], a[0:3], a[0:3], a[0:3] cbsz:3 abid:11008 1009// GFX950: v_mfma_i32_16x16x64_i8 a[0:3], v[0:3], v[0:3], a[4:7] ; encoding: [0x00,0x80,0xb6,0xd3,0x00,0x01,0x12,0x04]1010// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU1011v_mfma_i32_16x16x64_i8 a[0:3], v[0:3], v[0:3], a[4:7]1012 1013// GFX950: v_mfma_i32_16x16x64_i8 v[0:3], a[0:3], a[0:3], v[4:7] ; encoding: [0x00,0x00,0xb6,0xd3,0x00,0x01,0x12,0x1c]1014// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU1015v_mfma_i32_16x16x64_i8 v[0:3], a[0:3], a[0:3], v[4:7]1016 1017//===----------------------------------------------------------------------===//1018// v_mfma_i32_32x32x32_i81019//===----------------------------------------------------------------------===//1020 1021// GFX950: v_mfma_i32_32x32x32_i8 v[0:15], v[0:3], v[0:3], v[0:15] ; encoding: [0x00,0x00,0xb8,0xd3,0x00,0x01,0x02,0x04]1022// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU1023v_mfma_i32_32x32x32_i8 v[0:15], v[0:3], v[0:3], v[0:15]1024 1025// GFX950: v_mfma_i32_32x32x32_i8 a[0:15], a[0:3], a[0:3], a[0:15] ; encoding: [0x00,0x80,0xb8,0xd3,0x00,0x01,0x02,0x1c]1026// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU1027v_mfma_i32_32x32x32_i8 a[0:15], a[0:3], a[0:3], a[0:15]1028 1029// GFX950: v_mfma_i32_32x32x32_i8 v[0:15], v[0:3], v[0:3], v[0:15] ; encoding: [0x00,0x00,0xb8,0xd3,0x00,0x01,0x02,0x04]1030// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU1031v_mfma_i32_32x32x32i8 v[0:15], v[0:3], v[0:3], v[0:15]1032 1033// GFX950: v_mfma_i32_32x32x32_i8 a[0:15], a[0:3], a[0:3], a[0:15] ; encoding: [0x00,0x80,0xb8,0xd3,0x00,0x01,0x02,0x1c]1034// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU1035v_mfma_i32_32x32x32i8 a[0:15], a[0:3], a[0:3], a[0:15]1036 1037// GFX950: v_mfma_i32_32x32x32_i8 v[0:15], v[0:3], v[0:3], 1.0 ; encoding: [0x00,0x00,0xb8,0xd3,0x00,0x01,0xca,0x03]1038// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU1039v_mfma_i32_32x32x32_i8 v[0:15], v[0:3], v[0:3], 1.01040 1041// GFX950: v_mfma_i32_32x32x32_i8 a[0:15], a[0:3], a[0:3], 1.0 ; encoding: [0x00,0x80,0xb8,0xd3,0x00,0x01,0xca,0x1b]1042// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU1043v_mfma_i32_32x32x32_i8 a[0:15], a[0:3], a[0:3], 1.01044 1045// GFX950: v_mfma_i32_32x32x32_i8 v[0:15], v[0:3], v[0:3], v[0:15] blgp:5 ; encoding: [0x00,0x00,0xb8,0xd3,0x00,0x01,0x02,0xa4]1046// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU1047v_mfma_i32_32x32x32_i8 v[0:15], v[0:3], v[0:3], v[0:15] blgp:51048 1049// GFX950: v_mfma_i32_32x32x32_i8 a[0:15], a[0:3], a[0:3], a[0:15] blgp:2 ; encoding: [0x00,0x80,0xb8,0xd3,0x00,0x01,0x02,0x5c]1050// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU1051v_mfma_i32_32x32x32_i8 a[0:15], a[0:3], a[0:3], a[0:15] blgp:21052 1053// GFX950: v_mfma_i32_32x32x32_i8 v[0:15], v[0:3], v[0:3], v[0:15] cbsz:3 ; encoding: [0x00,0x03,0xb8,0xd3,0x00,0x01,0x02,0x04]1054// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU1055v_mfma_i32_32x32x32_i8 v[0:15], v[0:3], v[0:3], v[0:15] cbsz:31056 1057// GFX950: v_mfma_i32_32x32x32_i8 v[0:15], v[0:3], v[0:3], v[0:15] abid:1 ; encoding: [0x00,0x08,0xb8,0xd3,0x00,0x01,0x02,0x04]1058// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU1059v_mfma_i32_32x32x32_i8 v[0:15], v[0:3], v[0:3], v[0:15] abid:11060 1061// GFX950: v_mfma_i32_32x32x32_i8 a[0:15], a[0:3], a[0:3], a[0:15] cbsz:3 abid:1 ; encoding: [0x00,0x8b,0xb8,0xd3,0x00,0x01,0x02,0x1c]1062// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU1063v_mfma_i32_32x32x32_i8 a[0:15], a[0:3], a[0:3], a[0:15] cbsz:3 abid:11064 1065//===----------------------------------------------------------------------===//1066// v_mfma_f32_16x16x32_bf161067//===----------------------------------------------------------------------===//1068 1069// GFX950: v_mfma_f32_16x16x32_bf16 v[0:3], v[0:3], v[0:3], v[0:3] ; encoding: [0x00,0x00,0xb5,0xd3,0x00,0x01,0x02,0x04]1070// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU1071v_mfma_f32_16x16x32_bf16 v[0:3], v[0:3], v[0:3], v[0:3]1072 1073// GFX950: v_mfma_f32_16x16x32_bf16 v[0:3], v[0:3], v[0:3], v[0:3] ; encoding: [0x00,0x00,0xb5,0xd3,0x00,0x01,0x02,0x04]1074// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU1075v_mfma_f32_16x16x32bf16 v[0:3], v[0:3], v[0:3], v[0:3]1076 1077// GFX950: v_mfma_f32_16x16x32_bf16 a[0:3], a[0:3], a[0:3], a[0:3] ; encoding: [0x00,0x80,0xb5,0xd3,0x00,0x01,0x02,0x1c]1078// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU1079v_mfma_f32_16x16x32_bf16 a[0:3], a[0:3], a[0:3], a[0:3]1080 1081// GFX950: v_mfma_f32_16x16x32_bf16 a[0:3], a[0:3], a[0:3], a[0:3] ; encoding: [0x00,0x80,0xb5,0xd3,0x00,0x01,0x02,0x1c]1082// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU1083v_mfma_f32_16x16x32bf16 a[0:3], a[0:3], a[0:3], a[0:3]1084 1085// GFX950: v_mfma_f32_16x16x32_bf16 v[0:3], a[0:3], v[0:3], 1.0 ; encoding: [0x00,0x00,0xb5,0xd3,0x00,0x01,0xca,0x0b]1086// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU1087v_mfma_f32_16x16x32_bf16 v[0:3], a[0:3], v[0:3], 1.01088 1089// GFX950: v_mfma_f32_16x16x32_bf16 a[0:3], v[0:3], a[0:3], 1.0 ; encoding: [0x00,0x80,0xb5,0xd3,0x00,0x01,0xca,0x13]1090// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU1091v_mfma_f32_16x16x32_bf16 a[0:3], v[0:3], a[0:3], 1.01092 1093// GFX950: v_mfma_f32_16x16x32_bf16 v[0:3], v[0:3], v[0:3], v[0:3] blgp:5 ; encoding: [0x00,0x00,0xb5,0xd3,0x00,0x01,0x02,0xa4]1094// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU1095v_mfma_f32_16x16x32_bf16 v[0:3], v[0:3], v[0:3], v[0:3] blgp:51096 1097// GFX950: v_mfma_f32_16x16x32_bf16 a[0:3], a[0:3], a[0:3], a[0:3] blgp:1 ; encoding: [0x00,0x80,0xb5,0xd3,0x00,0x01,0x02,0x3c]1098// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU1099v_mfma_f32_16x16x32_bf16 a[0:3], a[0:3], a[0:3], a[0:3] blgp:11100 1101// GFX950: v_mfma_f32_16x16x32_bf16 a[0:3], a[0:3], a[0:3], a[0:3] cbsz:3 ; encoding: [0x00,0x83,0xb5,0xd3,0x00,0x01,0x02,0x1c]1102// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU1103v_mfma_f32_16x16x32_bf16 a[0:3], a[0:3], a[0:3], a[0:3] cbsz:31104 1105// GFX950: v_mfma_f32_16x16x32_bf16 a[0:3], a[0:3], a[0:3], a[0:3] abid:1 ; encoding: [0x00,0x88,0xb5,0xd3,0x00,0x01,0x02,0x1c]1106// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU1107v_mfma_f32_16x16x32_bf16 a[0:3], a[0:3], a[0:3], a[0:3] abid:11108 1109// GFX950: v_mfma_f32_16x16x32_bf16 a[0:3], a[0:3], a[0:3], a[0:3] cbsz:3 abid:1 ; encoding: [0x00,0x8b,0xb5,0xd3,0x00,0x01,0x02,0x1c]1110// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU1111v_mfma_f32_16x16x32_bf16 a[0:3], a[0:3], a[0:3], a[0:3] cbsz:3 abid:11112 1113// GFX950: v_mfma_f32_16x16x32_bf16 a[0:3], v[0:3], v[0:3], a[4:7] ; encoding: [0x00,0x80,0xb5,0xd3,0x00,0x01,0x12,0x04]1114// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU1115v_mfma_f32_16x16x32_bf16 a[0:3], v[0:3], v[0:3], a[4:7]1116 1117// GFX950: v_mfma_f32_16x16x32_bf16 v[0:3], a[0:3], a[0:3], v[4:7] ; encoding: [0x00,0x00,0xb5,0xd3,0x00,0x01,0x12,0x1c]1118// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU1119v_mfma_f32_16x16x32_bf16 v[0:3], a[0:3], a[0:3], v[4:7]1120 1121 1122//===----------------------------------------------------------------------===//1123// SMFMAC opcodes.1124//===----------------------------------------------------------------------===//1125 1126//===----------------------------------------------------------------------===//1127// v_smfmac_f32_16x16x64_f161128//===----------------------------------------------------------------------===//1129 1130// GFX950: v_smfmac_f32_16x16x64_f16 v[10:13], a[2:5], v[4:11], v3 cbsz:3 abid:1 ; encoding: [0x0a,0x0b,0xda,0xd3,0x02,0x09,0x0e,0x0c]1131// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU1132v_smfmac_f32_16x16x64_f16 v[10:13], a[2:5], v[4:11], v3 cbsz:3 abid:11133 1134// GFX950: v_smfmac_f32_16x16x64_f16 v[10:13], a[2:5], v[4:11], v3 cbsz:3 abid:1 ; encoding: [0x0a,0x0b,0xda,0xd3,0x02,0x09,0x0e,0x0c]1135// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU1136v_smfmac_f32_16x16x64f16 v[10:13], a[2:5], v[4:11], v3 cbsz:3 abid:11137 1138// GFX950: v_smfmac_f32_16x16x64_f16 a[10:13], v[2:5], a[4:11], v1 ; encoding: [0x0a,0x80,0xda,0xd3,0x02,0x09,0x06,0x14]1139// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU1140v_smfmac_f32_16x16x64_f16 a[10:13], v[2:5], a[4:11], v11141 1142// GFX950: v_smfmac_f32_16x16x64_f16 v[10:13], a[2:5], v[4:11], v2 cbsz:3 abid:1 ; encoding: [0x0a,0x0b,0xda,0xd3,0x02,0x09,0x0a,0x0c]1143// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU1144v_smfmac_f32_16x16x64_f16 v[10:13], a[2:5], v[4:11], v2 cbsz:3 abid:11145 1146// GFX950: v_smfmac_f32_16x16x64_f16 a[10:13], v[2:5], a[4:11], v3 ; encoding: [0x0a,0x80,0xda,0xd3,0x02,0x09,0x0e,0x14]1147// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU1148v_smfmac_f32_16x16x64_f16 a[10:13], v[2:5], a[4:11], v31149 1150// GFX950: v_smfmac_f32_16x16x64_f16 v[10:13], v[2:5], v[6:13], v2 cbsz:3 abid:1 ; encoding: [0x0a,0x0b,0xda,0xd3,0x02,0x0d,0x0a,0x04]1151// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU1152v_smfmac_f32_16x16x64_f16 v[10:13], v[2:5], v[6:13], v2 cbsz:3 abid:11153 1154// GFX950: v_smfmac_f32_16x16x64_f16 a[10:13], a[2:5], a[6:13], v2 cbsz:3 abid:1 ; encoding: [0x0a,0x8b,0xda,0xd3,0x02,0x0d,0x0a,0x1c]1155// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU1156v_smfmac_f32_16x16x64_f16 a[10:13], a[2:5], a[6:13], v2 cbsz:3 abid:11157 1158// GFX950: v_smfmac_f32_16x16x64_f16 a[10:13], a[2:5], a[6:13], v3 cbsz:1 abid:3 ; encoding: [0x0a,0x99,0xda,0xd3,0x02,0x0d,0x0e,0x1c]1159// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU1160v_smfmac_f32_16x16x64_f16 a[10:13], a[2:5], a[6:13], v3 cbsz:1 abid:31161 1162//===----------------------------------------------------------------------===//1163// v_smfmac_f32_32x32x32_f161164//===----------------------------------------------------------------------===//1165 1166// GFX950: v_smfmac_f32_32x32x32_f16 v[10:25], a[2:5], v[4:11], v3 cbsz:3 abid:1 ; encoding: [0x0a,0x0b,0xdb,0xd3,0x02,0x09,0x0e,0x0c]1167// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU1168v_smfmac_f32_32x32x32_f16 v[10:25], a[2:5], v[4:11], v3 cbsz:3 abid:11169 1170// GFX950: v_smfmac_f32_32x32x32_f16 v[10:25], a[2:5], v[4:11], v3 cbsz:3 abid:1 ; encoding: [0x0a,0x0b,0xdb,0xd3,0x02,0x09,0x0e,0x0c]1171// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU1172v_smfmac_f32_32x32x32f16 v[10:25], a[2:5], v[4:11], v3 cbsz:3 abid:11173 1174// GFX950: v_smfmac_f32_32x32x32_f16 a[10:25], v[2:5], a[4:11], v1 ; encoding: [0x0a,0x80,0xdb,0xd3,0x02,0x09,0x06,0x14]1175// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU1176v_smfmac_f32_32x32x32_f16 a[10:25], v[2:5], a[4:11], v11177 1178// GFX950: v_smfmac_f32_32x32x32_f16 v[10:25], a[2:5], v[4:11], v2 cbsz:3 abid:1 ; encoding: [0x0a,0x0b,0xdb,0xd3,0x02,0x09,0x0a,0x0c]1179// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU1180v_smfmac_f32_32x32x32_f16 v[10:25], a[2:5], v[4:11], v2 cbsz:3 abid:11181 1182// GFX950: v_smfmac_f32_32x32x32_f16 a[10:25], v[2:5], a[4:11], v3 ; encoding: [0x0a,0x80,0xdb,0xd3,0x02,0x09,0x0e,0x14]1183// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU1184v_smfmac_f32_32x32x32_f16 a[10:25], v[2:5], a[4:11], v31185 1186// GFX950: v_smfmac_f32_32x32x32_f16 v[10:25], v[2:5], v[6:13], v2 cbsz:3 abid:1 ; encoding: [0x0a,0x0b,0xdb,0xd3,0x02,0x0d,0x0a,0x04]1187// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU1188v_smfmac_f32_32x32x32_f16 v[10:25], v[2:5], v[6:13], v2 cbsz:3 abid:11189 1190// GFX950: v_smfmac_f32_32x32x32_f16 a[10:25], a[2:5], a[6:13], v2 cbsz:3 abid:1 ; encoding: [0x0a,0x8b,0xdb,0xd3,0x02,0x0d,0x0a,0x1c]1191// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU1192v_smfmac_f32_32x32x32_f16 a[10:25], a[2:5], a[6:13], v2 cbsz:3 abid:11193 1194// GFX950: v_smfmac_f32_32x32x32_f16 a[10:25], a[2:5], a[6:13], v3 cbsz:1 abid:3 ; encoding: [0x0a,0x99,0xdb,0xd3,0x02,0x0d,0x0e,0x1c]1195// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU1196v_smfmac_f32_32x32x32_f16 a[10:25], a[2:5], a[6:13], v3 cbsz:1 abid:31197 1198//===----------------------------------------------------------------------===//1199// v_smfmac_f32_16x16x64_bf161200//===----------------------------------------------------------------------===//1201 1202// GFX950: v_smfmac_f32_16x16x64_bf16 v[10:13], a[2:5], v[4:11], v3 cbsz:3 abid:1 ; encoding: [0x0a,0x0b,0xb9,0xd3,0x02,0x09,0x0e,0x0c]1203// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU1204v_smfmac_f32_16x16x64_bf16 v[10:13], a[2:5], v[4:11], v3 cbsz:3 abid:11205 1206// GFX950: v_smfmac_f32_16x16x64_bf16 v[10:13], a[2:5], v[4:11], v3 cbsz:3 abid:1 ; encoding: [0x0a,0x0b,0xb9,0xd3,0x02,0x09,0x0e,0x0c]1207// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU1208v_smfmac_f32_16x16x64bf16 v[10:13], a[2:5], v[4:11], v3 cbsz:3 abid:11209 1210// GFX950: v_smfmac_f32_16x16x64_bf16 a[10:13], v[2:5], a[4:11], v1 ; encoding: [0x0a,0x80,0xb9,0xd3,0x02,0x09,0x06,0x14]1211// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU1212v_smfmac_f32_16x16x64_bf16 a[10:13], v[2:5], a[4:11], v11213 1214// GFX950: v_smfmac_f32_16x16x64_bf16 v[10:13], a[2:5], v[4:11], v2 cbsz:3 abid:1 ; encoding: [0x0a,0x0b,0xb9,0xd3,0x02,0x09,0x0a,0x0c]1215// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU1216v_smfmac_f32_16x16x64_bf16 v[10:13], a[2:5], v[4:11], v2 cbsz:3 abid:11217 1218// GFX950: v_smfmac_f32_16x16x64_bf16 a[10:13], v[2:5], a[4:11], v3 ; encoding: [0x0a,0x80,0xb9,0xd3,0x02,0x09,0x0e,0x14]1219// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU1220v_smfmac_f32_16x16x64_bf16 a[10:13], v[2:5], a[4:11], v31221 1222// GFX950: v_smfmac_f32_16x16x64_bf16 v[10:13], v[2:5], v[6:13], v2 cbsz:3 abid:1 ; encoding: [0x0a,0x0b,0xb9,0xd3,0x02,0x0d,0x0a,0x04]1223// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU1224v_smfmac_f32_16x16x64_bf16 v[10:13], v[2:5], v[6:13], v2 cbsz:3 abid:11225 1226// GFX950: v_smfmac_f32_16x16x64_bf16 a[10:13], a[2:5], a[6:13], v2 cbsz:3 abid:1 ; encoding: [0x0a,0x8b,0xb9,0xd3,0x02,0x0d,0x0a,0x1c]1227// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU1228v_smfmac_f32_16x16x64_bf16 a[10:13], a[2:5], a[6:13], v2 cbsz:3 abid:11229 1230// GFX950: v_smfmac_f32_16x16x64_bf16 a[10:13], a[2:5], a[6:13], v3 cbsz:1 abid:3 ; encoding: [0x0a,0x99,0xb9,0xd3,0x02,0x0d,0x0e,0x1c]1231// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU1232v_smfmac_f32_16x16x64_bf16 a[10:13], a[2:5], a[6:13], v3 cbsz:1 abid:31233 1234//===----------------------------------------------------------------------===//1235// v_smfmac_f32_32x32x32_bf161236//===----------------------------------------------------------------------===//1237 1238// GFX950: v_smfmac_f32_32x32x32_bf16 v[10:25], a[2:5], v[4:11], v3 cbsz:3 abid:1 ; encoding: [0x0a,0x0b,0xc6,0xd3,0x02,0x09,0x0e,0x0c]1239// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU1240v_smfmac_f32_32x32x32_bf16 v[10:25], a[2:5], v[4:11], v3 cbsz:3 abid:11241 1242// GFX950: v_smfmac_f32_32x32x32_bf16 v[10:25], a[2:5], v[4:11], v3 cbsz:3 abid:1 ; encoding: [0x0a,0x0b,0xc6,0xd3,0x02,0x09,0x0e,0x0c]1243// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU1244v_smfmac_f32_32x32x32bf16 v[10:25], a[2:5], v[4:11], v3 cbsz:3 abid:11245 1246// GFX950: v_smfmac_f32_32x32x32_bf16 a[10:25], v[2:5], a[4:11], v1 ; encoding: [0x0a,0x80,0xc6,0xd3,0x02,0x09,0x06,0x14]1247// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU1248v_smfmac_f32_32x32x32_bf16 a[10:25], v[2:5], a[4:11], v11249 1250// GFX950: v_smfmac_f32_32x32x32_bf16 v[10:25], a[2:5], v[4:11], v2 cbsz:3 abid:1 ; encoding: [0x0a,0x0b,0xc6,0xd3,0x02,0x09,0x0a,0x0c]1251// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU1252v_smfmac_f32_32x32x32_bf16 v[10:25], a[2:5], v[4:11], v2 cbsz:3 abid:11253 1254// GFX950: v_smfmac_f32_32x32x32_bf16 a[10:25], v[2:5], a[4:11], v3 ; encoding: [0x0a,0x80,0xc6,0xd3,0x02,0x09,0x0e,0x14]1255// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU1256v_smfmac_f32_32x32x32_bf16 a[10:25], v[2:5], a[4:11], v31257 1258// GFX950: v_smfmac_f32_32x32x32_bf16 v[10:25], v[2:5], v[6:13], v2 cbsz:3 abid:1 ; encoding: [0x0a,0x0b,0xc6,0xd3,0x02,0x0d,0x0a,0x04]1259// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU1260v_smfmac_f32_32x32x32_bf16 v[10:25], v[2:5], v[6:13], v2 cbsz:3 abid:11261 1262// GFX950: v_smfmac_f32_32x32x32_bf16 a[10:25], a[2:5], a[6:13], v2 cbsz:3 abid:1 ; encoding: [0x0a,0x8b,0xc6,0xd3,0x02,0x0d,0x0a,0x1c]1263// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU1264v_smfmac_f32_32x32x32_bf16 a[10:25], a[2:5], a[6:13], v2 cbsz:3 abid:11265 1266// GFX950: v_smfmac_f32_32x32x32_bf16 a[10:25], a[2:5], a[6:13], v3 cbsz:1 abid:3 ; encoding: [0x0a,0x99,0xc6,0xd3,0x02,0x0d,0x0e,0x1c]1267// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU1268v_smfmac_f32_32x32x32_bf16 a[10:25], a[2:5], a[6:13], v3 cbsz:1 abid:31269 1270//===----------------------------------------------------------------------===//1271// v_smfmac_i32_16x16x128_i81272//===----------------------------------------------------------------------===//1273 1274// GFX950: v_smfmac_i32_16x16x128_i8 v[10:13], a[2:5], v[4:11], v3 cbsz:3 abid:1 ; encoding: [0x0a,0x0b,0xba,0xd3,0x02,0x09,0x0e,0x0c]1275// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU1276v_smfmac_i32_16x16x128_i8 v[10:13], a[2:5], v[4:11], v3 cbsz:3 abid:11277 1278// GFX950: v_smfmac_i32_16x16x128_i8 v[10:13], a[2:5], v[4:11], v3 cbsz:3 abid:1 ; encoding: [0x0a,0x0b,0xba,0xd3,0x02,0x09,0x0e,0x0c]1279// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU1280v_smfmac_i32_16x16x128i8 v[10:13], a[2:5], v[4:11], v3 cbsz:3 abid:11281 1282// GFX950: v_smfmac_i32_16x16x128_i8 a[10:13], v[2:5], a[4:11], v1 ; encoding: [0x0a,0x80,0xba,0xd3,0x02,0x09,0x06,0x14]1283// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU1284v_smfmac_i32_16x16x128_i8 a[10:13], v[2:5], a[4:11], v11285 1286// GFX950: v_smfmac_i32_16x16x128_i8 v[10:13], a[2:5], v[4:11], v2 cbsz:3 abid:1 ; encoding: [0x0a,0x0b,0xba,0xd3,0x02,0x09,0x0a,0x0c]1287// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU1288v_smfmac_i32_16x16x128_i8 v[10:13], a[2:5], v[4:11], v2 cbsz:3 abid:11289 1290// GFX950: v_smfmac_i32_16x16x128_i8 a[10:13], v[2:5], a[4:11], v3 ; encoding: [0x0a,0x80,0xba,0xd3,0x02,0x09,0x0e,0x14]1291// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU1292v_smfmac_i32_16x16x128_i8 a[10:13], v[2:5], a[4:11], v31293 1294// GFX950: v_smfmac_i32_16x16x128_i8 v[10:13], v[2:5], v[6:13], v2 cbsz:3 abid:1 ; encoding: [0x0a,0x0b,0xba,0xd3,0x02,0x0d,0x0a,0x04]1295// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU1296v_smfmac_i32_16x16x128_i8 v[10:13], v[2:5], v[6:13], v2 cbsz:3 abid:11297 1298// GFX950: v_smfmac_i32_16x16x128_i8 a[10:13], a[2:5], a[6:13], v2 cbsz:3 abid:1 ; encoding: [0x0a,0x8b,0xba,0xd3,0x02,0x0d,0x0a,0x1c]1299// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU1300v_smfmac_i32_16x16x128_i8 a[10:13], a[2:5], a[6:13], v2 cbsz:3 abid:11301 1302// GFX950: v_smfmac_i32_16x16x128_i8 a[10:13], a[2:5], a[6:13], v3 cbsz:1 abid:3 ; encoding: [0x0a,0x99,0xba,0xd3,0x02,0x0d,0x0e,0x1c]1303// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU1304v_smfmac_i32_16x16x128_i8 a[10:13], a[2:5], a[6:13], v3 cbsz:1 abid:31305 1306//===----------------------------------------------------------------------===//1307// v_smfmac_i32_32x32x64_i81308//===----------------------------------------------------------------------===//1309 1310// GFX950: v_smfmac_i32_32x32x64_i8 v[10:25], a[2:5], v[4:11], v3 cbsz:3 abid:1 ; encoding: [0x0a,0x0b,0xc7,0xd3,0x02,0x09,0x0e,0x0c]1311// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU1312v_smfmac_i32_32x32x64_i8 v[10:25], a[2:5], v[4:11], v3 cbsz:3 abid:11313 1314// GFX950: v_smfmac_i32_32x32x64_i8 v[10:25], a[2:5], v[4:11], v3 cbsz:3 abid:1 ; encoding: [0x0a,0x0b,0xc7,0xd3,0x02,0x09,0x0e,0x0c]1315// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU1316v_smfmac_i32_32x32x64i8 v[10:25], a[2:5], v[4:11], v3 cbsz:3 abid:11317 1318// GFX950: v_smfmac_i32_32x32x64_i8 a[10:25], v[2:5], a[4:11], v1 ; encoding: [0x0a,0x80,0xc7,0xd3,0x02,0x09,0x06,0x14]1319// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU1320v_smfmac_i32_32x32x64_i8 a[10:25], v[2:5], a[4:11], v11321 1322// GFX950: v_smfmac_i32_32x32x64_i8 v[10:25], a[2:5], v[4:11], v2 cbsz:3 abid:1 ; encoding: [0x0a,0x0b,0xc7,0xd3,0x02,0x09,0x0a,0x0c]1323// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU1324v_smfmac_i32_32x32x64_i8 v[10:25], a[2:5], v[4:11], v2 cbsz:3 abid:11325 1326// GFX950: v_smfmac_i32_32x32x64_i8 a[10:25], v[2:5], a[4:11], v3 ; encoding: [0x0a,0x80,0xc7,0xd3,0x02,0x09,0x0e,0x14]1327// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU1328v_smfmac_i32_32x32x64_i8 a[10:25], v[2:5], a[4:11], v31329 1330// GFX950: v_smfmac_i32_32x32x64_i8 v[10:25], v[2:5], v[6:13], v2 cbsz:3 abid:1 ; encoding: [0x0a,0x0b,0xc7,0xd3,0x02,0x0d,0x0a,0x04]1331// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU1332v_smfmac_i32_32x32x64_i8 v[10:25], v[2:5], v[6:13], v2 cbsz:3 abid:11333 1334// GFX950: v_smfmac_i32_32x32x64_i8 a[10:25], a[2:5], a[6:13], v2 cbsz:3 abid:1 ; encoding: [0x0a,0x8b,0xc7,0xd3,0x02,0x0d,0x0a,0x1c]1335// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU1336v_smfmac_i32_32x32x64_i8 a[10:25], a[2:5], a[6:13], v2 cbsz:3 abid:11337 1338// GFX950: v_smfmac_i32_32x32x64_i8 a[10:25], a[2:5], a[6:13], v3 cbsz:1 abid:3 ; encoding: [0x0a,0x99,0xc7,0xd3,0x02,0x0d,0x0e,0x1c]1339// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU1340v_smfmac_i32_32x32x64_i8 a[10:25], a[2:5], a[6:13], v3 cbsz:1 abid:31341 1342//===----------------------------------------------------------------------===//1343// v_smfmac_f32_16x16x128_bf8_bf81344//===----------------------------------------------------------------------===//1345 1346// GFX950: v_smfmac_f32_16x16x128_bf8_bf8 v[10:13], a[2:5], v[4:11], v3 cbsz:3 abid:1 ; encoding: [0x0a,0x0b,0xbb,0xd3,0x02,0x09,0x0e,0x0c]1347// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU1348v_smfmac_f32_16x16x128_bf8_bf8 v[10:13], a[2:5], v[4:11], v3 cbsz:3 abid:11349 1350// GFX950: v_smfmac_f32_16x16x128_bf8_bf8 v[10:13], a[2:5], v[4:11], v3 cbsz:3 abid:1 ; encoding: [0x0a,0x0b,0xbb,0xd3,0x02,0x09,0x0e,0x0c]1351// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU1352v_smfmac_f32_16x16x128bf8bf8 v[10:13], a[2:5], v[4:11], v3 cbsz:3 abid:11353 1354// GFX950: v_smfmac_f32_16x16x128_bf8_bf8 a[10:13], v[2:5], a[4:11], v1 ; encoding: [0x0a,0x80,0xbb,0xd3,0x02,0x09,0x06,0x14]1355// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU1356v_smfmac_f32_16x16x128_bf8_bf8 a[10:13], v[2:5], a[4:11], v11357 1358// GFX950: v_smfmac_f32_16x16x128_bf8_bf8 v[10:13], a[2:5], v[4:11], v2 cbsz:3 abid:1 ; encoding: [0x0a,0x0b,0xbb,0xd3,0x02,0x09,0x0a,0x0c]1359// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU1360v_smfmac_f32_16x16x128_bf8_bf8 v[10:13], a[2:5], v[4:11], v2 cbsz:3 abid:11361 1362// GFX950: v_smfmac_f32_16x16x128_bf8_bf8 a[10:13], v[2:5], a[4:11], v3 ; encoding: [0x0a,0x80,0xbb,0xd3,0x02,0x09,0x0e,0x14]1363// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU1364v_smfmac_f32_16x16x128_bf8_bf8 a[10:13], v[2:5], a[4:11], v31365 1366// GFX950: v_smfmac_f32_16x16x128_bf8_bf8 v[10:13], v[2:5], v[6:13], v2 cbsz:3 abid:1 ; encoding: [0x0a,0x0b,0xbb,0xd3,0x02,0x0d,0x0a,0x04]1367// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU1368v_smfmac_f32_16x16x128_bf8_bf8 v[10:13], v[2:5], v[6:13], v2 cbsz:3 abid:11369 1370// GFX950: v_smfmac_f32_16x16x128_bf8_bf8 a[10:13], a[2:5], a[6:13], v2 cbsz:3 abid:1 ; encoding: [0x0a,0x8b,0xbb,0xd3,0x02,0x0d,0x0a,0x1c]1371// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU1372v_smfmac_f32_16x16x128_bf8_bf8 a[10:13], a[2:5], a[6:13], v2 cbsz:3 abid:11373 1374// GFX950: v_smfmac_f32_16x16x128_bf8_bf8 a[10:13], a[2:5], a[6:13], v3 cbsz:1 abid:3 ; encoding: [0x0a,0x99,0xbb,0xd3,0x02,0x0d,0x0e,0x1c]1375// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU1376v_smfmac_f32_16x16x128_bf8_bf8 a[10:13], a[2:5], a[6:13], v3 cbsz:1 abid:31377 1378//===----------------------------------------------------------------------===//1379// v_smfmac_f32_16x16x128_bf8_fp81380//===----------------------------------------------------------------------===//1381 1382// GFX950: v_smfmac_f32_16x16x128_bf8_fp8 v[10:13], a[2:5], v[4:11], v3 cbsz:3 abid:1 ; encoding: [0x0a,0x0b,0xbc,0xd3,0x02,0x09,0x0e,0x0c]1383// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU1384v_smfmac_f32_16x16x128_bf8_fp8 v[10:13], a[2:5], v[4:11], v3 cbsz:3 abid:11385 1386// GFX950: v_smfmac_f32_16x16x128_bf8_fp8 v[10:13], a[2:5], v[4:11], v3 cbsz:3 abid:1 ; encoding: [0x0a,0x0b,0xbc,0xd3,0x02,0x09,0x0e,0x0c]1387// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU1388v_smfmac_f32_16x16x128bf8fp8 v[10:13], a[2:5], v[4:11], v3 cbsz:3 abid:11389 1390// GFX950: v_smfmac_f32_16x16x128_bf8_fp8 a[10:13], v[2:5], a[4:11], v1 ; encoding: [0x0a,0x80,0xbc,0xd3,0x02,0x09,0x06,0x14]1391// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU1392v_smfmac_f32_16x16x128_bf8_fp8 a[10:13], v[2:5], a[4:11], v11393 1394// GFX950: v_smfmac_f32_16x16x128_bf8_fp8 v[10:13], a[2:5], v[4:11], v2 cbsz:3 abid:1 ; encoding: [0x0a,0x0b,0xbc,0xd3,0x02,0x09,0x0a,0x0c]1395// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU1396v_smfmac_f32_16x16x128_bf8_fp8 v[10:13], a[2:5], v[4:11], v2 cbsz:3 abid:11397 1398// GFX950: v_smfmac_f32_16x16x128_bf8_fp8 a[10:13], v[2:5], a[4:11], v3 ; encoding: [0x0a,0x80,0xbc,0xd3,0x02,0x09,0x0e,0x14]1399// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU1400v_smfmac_f32_16x16x128_bf8_fp8 a[10:13], v[2:5], a[4:11], v31401 1402// GFX950: v_smfmac_f32_16x16x128_bf8_fp8 v[10:13], v[2:5], v[6:13], v2 cbsz:3 abid:1 ; encoding: [0x0a,0x0b,0xbc,0xd3,0x02,0x0d,0x0a,0x04]1403// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU1404v_smfmac_f32_16x16x128_bf8_fp8 v[10:13], v[2:5], v[6:13], v2 cbsz:3 abid:11405 1406// GFX950: v_smfmac_f32_16x16x128_bf8_fp8 a[10:13], a[2:5], a[6:13], v2 cbsz:3 abid:1 ; encoding: [0x0a,0x8b,0xbc,0xd3,0x02,0x0d,0x0a,0x1c]1407// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU1408v_smfmac_f32_16x16x128_bf8_fp8 a[10:13], a[2:5], a[6:13], v2 cbsz:3 abid:11409 1410// GFX950: v_smfmac_f32_16x16x128_bf8_fp8 a[10:13], a[2:5], a[6:13], v3 cbsz:1 abid:3 ; encoding: [0x0a,0x99,0xbc,0xd3,0x02,0x0d,0x0e,0x1c]1411// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU1412v_smfmac_f32_16x16x128_bf8_fp8 a[10:13], a[2:5], a[6:13], v3 cbsz:1 abid:31413 1414//===----------------------------------------------------------------------===//1415// v_smfmac_f32_16x16x128_fp8_bf81416//===----------------------------------------------------------------------===//1417 1418// GFX950: v_smfmac_f32_16x16x128_fp8_bf8 v[10:13], a[2:5], v[4:11], v3 cbsz:3 abid:1 ; encoding: [0x0a,0x0b,0xbd,0xd3,0x02,0x09,0x0e,0x0c]1419// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU1420v_smfmac_f32_16x16x128_fp8_bf8 v[10:13], a[2:5], v[4:11], v3 cbsz:3 abid:11421 1422// GFX950: v_smfmac_f32_16x16x128_fp8_bf8 v[10:13], a[2:5], v[4:11], v3 cbsz:3 abid:1 ; encoding: [0x0a,0x0b,0xbd,0xd3,0x02,0x09,0x0e,0x0c]1423// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU1424v_smfmac_f32_16x16x128fp8bf8 v[10:13], a[2:5], v[4:11], v3 cbsz:3 abid:11425 1426// GFX950: v_smfmac_f32_16x16x128_fp8_bf8 a[10:13], v[2:5], a[4:11], v1 ; encoding: [0x0a,0x80,0xbd,0xd3,0x02,0x09,0x06,0x14]1427// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU1428v_smfmac_f32_16x16x128_fp8_bf8 a[10:13], v[2:5], a[4:11], v11429 1430// GFX950: v_smfmac_f32_16x16x128_fp8_bf8 v[10:13], a[2:5], v[4:11], v2 cbsz:3 abid:1 ; encoding: [0x0a,0x0b,0xbd,0xd3,0x02,0x09,0x0a,0x0c]1431// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU1432v_smfmac_f32_16x16x128_fp8_bf8 v[10:13], a[2:5], v[4:11], v2 cbsz:3 abid:11433 1434// GFX950: v_smfmac_f32_16x16x128_fp8_bf8 a[10:13], v[2:5], a[4:11], v3 ; encoding: [0x0a,0x80,0xbd,0xd3,0x02,0x09,0x0e,0x14]1435// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU1436v_smfmac_f32_16x16x128_fp8_bf8 a[10:13], v[2:5], a[4:11], v31437 1438// GFX950: v_smfmac_f32_16x16x128_fp8_bf8 v[10:13], v[2:5], v[6:13], v2 cbsz:3 abid:1 ; encoding: [0x0a,0x0b,0xbd,0xd3,0x02,0x0d,0x0a,0x04]1439// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU1440v_smfmac_f32_16x16x128_fp8_bf8 v[10:13], v[2:5], v[6:13], v2 cbsz:3 abid:11441 1442// GFX950: v_smfmac_f32_16x16x128_fp8_bf8 a[10:13], a[2:5], a[6:13], v2 cbsz:3 abid:1 ; encoding: [0x0a,0x8b,0xbd,0xd3,0x02,0x0d,0x0a,0x1c]1443// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU1444v_smfmac_f32_16x16x128_fp8_bf8 a[10:13], a[2:5], a[6:13], v2 cbsz:3 abid:11445 1446// GFX950: v_smfmac_f32_16x16x128_fp8_bf8 a[10:13], a[2:5], a[6:13], v3 cbsz:1 abid:3 ; encoding: [0x0a,0x99,0xbd,0xd3,0x02,0x0d,0x0e,0x1c]1447// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU1448v_smfmac_f32_16x16x128_fp8_bf8 a[10:13], a[2:5], a[6:13], v3 cbsz:1 abid:31449 1450//===----------------------------------------------------------------------===//1451// v_smfmac_f32_16x16x128_fp8_fp81452//===----------------------------------------------------------------------===//1453 1454// GFX950: v_smfmac_f32_16x16x128_fp8_fp8 v[10:13], a[2:5], v[4:11], v3 cbsz:3 abid:1 ; encoding: [0x0a,0x0b,0xc3,0xd3,0x02,0x09,0x0e,0x0c]1455// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU1456v_smfmac_f32_16x16x128_fp8_fp8 v[10:13], a[2:5], v[4:11], v3 cbsz:3 abid:11457 1458// GFX950: v_smfmac_f32_16x16x128_fp8_fp8 v[10:13], a[2:5], v[4:11], v3 cbsz:3 abid:1 ; encoding: [0x0a,0x0b,0xc3,0xd3,0x02,0x09,0x0e,0x0c]1459// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU1460v_smfmac_f32_16x16x128fp8fp8 v[10:13], a[2:5], v[4:11], v3 cbsz:3 abid:11461 1462// GFX950: v_smfmac_f32_16x16x128_fp8_fp8 a[10:13], v[2:5], a[4:11], v1 ; encoding: [0x0a,0x80,0xc3,0xd3,0x02,0x09,0x06,0x14]1463// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU1464v_smfmac_f32_16x16x128_fp8_fp8 a[10:13], v[2:5], a[4:11], v11465 1466// GFX950: v_smfmac_f32_16x16x128_fp8_fp8 v[10:13], a[2:5], v[4:11], v2 cbsz:3 abid:1 ; encoding: [0x0a,0x0b,0xc3,0xd3,0x02,0x09,0x0a,0x0c]1467// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU1468v_smfmac_f32_16x16x128_fp8_fp8 v[10:13], a[2:5], v[4:11], v2 cbsz:3 abid:11469 1470// GFX950: v_smfmac_f32_16x16x128_fp8_fp8 a[10:13], v[2:5], a[4:11], v3 ; encoding: [0x0a,0x80,0xc3,0xd3,0x02,0x09,0x0e,0x14]1471// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU1472v_smfmac_f32_16x16x128_fp8_fp8 a[10:13], v[2:5], a[4:11], v31473 1474// GFX950: v_smfmac_f32_16x16x128_fp8_fp8 v[10:13], v[2:5], v[6:13], v2 cbsz:3 abid:1 ; encoding: [0x0a,0x0b,0xc3,0xd3,0x02,0x0d,0x0a,0x04]1475// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU1476v_smfmac_f32_16x16x128_fp8_fp8 v[10:13], v[2:5], v[6:13], v2 cbsz:3 abid:11477 1478// GFX950: v_smfmac_f32_16x16x128_fp8_fp8 a[10:13], a[2:5], a[6:13], v2 cbsz:3 abid:1 ; encoding: [0x0a,0x8b,0xc3,0xd3,0x02,0x0d,0x0a,0x1c]1479// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU1480v_smfmac_f32_16x16x128_fp8_fp8 a[10:13], a[2:5], a[6:13], v2 cbsz:3 abid:11481 1482// GFX950: v_smfmac_f32_16x16x128_fp8_fp8 a[10:13], a[2:5], a[6:13], v3 cbsz:1 abid:3 ; encoding: [0x0a,0x99,0xc3,0xd3,0x02,0x0d,0x0e,0x1c]1483// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU1484v_smfmac_f32_16x16x128_fp8_fp8 a[10:13], a[2:5], a[6:13], v3 cbsz:1 abid:31485 1486//===----------------------------------------------------------------------===//1487// v_smfmac_f32_32x32x64_bf8_bf81488//===----------------------------------------------------------------------===//1489 1490// GFX950: v_smfmac_f32_32x32x64_bf8_bf8 v[10:25], a[2:5], v[4:11], v3 cbsz:3 abid:1 ; encoding: [0x0a,0x0b,0xcb,0xd3,0x02,0x09,0x0e,0x0c]1491// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU1492v_smfmac_f32_32x32x64_bf8_bf8 v[10:25], a[2:5], v[4:11], v3 cbsz:3 abid:11493 1494// GFX950: v_smfmac_f32_32x32x64_bf8_bf8 v[10:25], a[2:5], v[4:11], v3 cbsz:3 abid:1 ; encoding: [0x0a,0x0b,0xcb,0xd3,0x02,0x09,0x0e,0x0c]1495// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU1496v_smfmac_f32_32x32x64bf8bf8 v[10:25], a[2:5], v[4:11], v3 cbsz:3 abid:11497 1498// GFX950: v_smfmac_f32_32x32x64_bf8_bf8 a[10:25], v[2:5], a[4:11], v1 ; encoding: [0x0a,0x80,0xcb,0xd3,0x02,0x09,0x06,0x14]1499// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU1500v_smfmac_f32_32x32x64_bf8_bf8 a[10:25], v[2:5], a[4:11], v11501 1502// GFX950: v_smfmac_f32_32x32x64_bf8_bf8 v[10:25], a[2:5], v[4:11], v2 cbsz:3 abid:1 ; encoding: [0x0a,0x0b,0xcb,0xd3,0x02,0x09,0x0a,0x0c]1503// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU1504v_smfmac_f32_32x32x64_bf8_bf8 v[10:25], a[2:5], v[4:11], v2 cbsz:3 abid:11505 1506// GFX950: v_smfmac_f32_32x32x64_bf8_bf8 a[10:25], v[2:5], a[4:11], v3 ; encoding: [0x0a,0x80,0xcb,0xd3,0x02,0x09,0x0e,0x14]1507// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU1508v_smfmac_f32_32x32x64_bf8_bf8 a[10:25], v[2:5], a[4:11], v31509 1510// GFX950: v_smfmac_f32_32x32x64_bf8_bf8 v[10:25], v[2:5], v[6:13], v2 cbsz:3 abid:1 ; encoding: [0x0a,0x0b,0xcb,0xd3,0x02,0x0d,0x0a,0x04]1511// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU1512v_smfmac_f32_32x32x64_bf8_bf8 v[10:25], v[2:5], v[6:13], v2 cbsz:3 abid:11513 1514// GFX950: v_smfmac_f32_32x32x64_bf8_bf8 a[10:25], a[2:5], a[6:13], v2 cbsz:3 abid:1 ; encoding: [0x0a,0x8b,0xcb,0xd3,0x02,0x0d,0x0a,0x1c]1515// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU1516v_smfmac_f32_32x32x64_bf8_bf8 a[10:25], a[2:5], a[6:13], v2 cbsz:3 abid:11517 1518// GFX950: v_smfmac_f32_32x32x64_bf8_bf8 a[10:25], a[2:5], a[6:13], v3 cbsz:1 abid:3 ; encoding: [0x0a,0x99,0xcb,0xd3,0x02,0x0d,0x0e,0x1c]1519// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU1520v_smfmac_f32_32x32x64_bf8_bf8 a[10:25], a[2:5], a[6:13], v3 cbsz:1 abid:31521 1522//===----------------------------------------------------------------------===//1523// v_smfmac_f32_32x32x64_bf8_fp81524//===----------------------------------------------------------------------===//1525 1526// GFX950: v_smfmac_f32_32x32x64_bf8_fp8 v[10:25], a[2:5], v[4:11], v3 cbsz:3 abid:1 ; encoding: [0x0a,0x0b,0xce,0xd3,0x02,0x09,0x0e,0x0c]1527// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU1528v_smfmac_f32_32x32x64_bf8_fp8 v[10:25], a[2:5], v[4:11], v3 cbsz:3 abid:11529 1530// GFX950: v_smfmac_f32_32x32x64_bf8_fp8 v[10:25], a[2:5], v[4:11], v3 cbsz:3 abid:1 ; encoding: [0x0a,0x0b,0xce,0xd3,0x02,0x09,0x0e,0x0c]1531// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU1532v_smfmac_f32_32x32x64bf8fp8 v[10:25], a[2:5], v[4:11], v3 cbsz:3 abid:11533 1534// GFX950: v_smfmac_f32_32x32x64_bf8_fp8 a[10:25], v[2:5], a[4:11], v1 ; encoding: [0x0a,0x80,0xce,0xd3,0x02,0x09,0x06,0x14]1535// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU1536v_smfmac_f32_32x32x64_bf8_fp8 a[10:25], v[2:5], a[4:11], v11537 1538// GFX950: v_smfmac_f32_32x32x64_bf8_fp8 v[10:25], a[2:5], v[4:11], v2 cbsz:3 abid:1 ; encoding: [0x0a,0x0b,0xce,0xd3,0x02,0x09,0x0a,0x0c]1539// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU1540v_smfmac_f32_32x32x64_bf8_fp8 v[10:25], a[2:5], v[4:11], v2 cbsz:3 abid:11541 1542// GFX950: v_smfmac_f32_32x32x64_bf8_fp8 a[10:25], v[2:5], a[4:11], v3 ; encoding: [0x0a,0x80,0xce,0xd3,0x02,0x09,0x0e,0x14]1543// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU1544v_smfmac_f32_32x32x64_bf8_fp8 a[10:25], v[2:5], a[4:11], v31545 1546// GFX950: v_smfmac_f32_32x32x64_bf8_fp8 v[10:25], v[2:5], v[6:13], v2 cbsz:3 abid:1 ; encoding: [0x0a,0x0b,0xce,0xd3,0x02,0x0d,0x0a,0x04]1547// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU1548v_smfmac_f32_32x32x64_bf8_fp8 v[10:25], v[2:5], v[6:13], v2 cbsz:3 abid:11549 1550// GFX950: v_smfmac_f32_32x32x64_bf8_fp8 a[10:25], a[2:5], a[6:13], v2 cbsz:3 abid:1 ; encoding: [0x0a,0x8b,0xce,0xd3,0x02,0x0d,0x0a,0x1c]1551// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU1552v_smfmac_f32_32x32x64_bf8_fp8 a[10:25], a[2:5], a[6:13], v2 cbsz:3 abid:11553 1554// GFX950: v_smfmac_f32_32x32x64_bf8_fp8 a[10:25], a[2:5], a[6:13], v3 cbsz:1 abid:3 ; encoding: [0x0a,0x99,0xce,0xd3,0x02,0x0d,0x0e,0x1c]1555// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU1556v_smfmac_f32_32x32x64_bf8_fp8 a[10:25], a[2:5], a[6:13], v3 cbsz:1 abid:31557 1558//===----------------------------------------------------------------------===//1559// v_smfmac_f32_32x32x64_fp8_bf81560//===----------------------------------------------------------------------===//1561 1562// GFX950: v_smfmac_f32_32x32x64_fp8_bf8 v[10:25], a[2:5], v[4:11], v3 cbsz:3 abid:1 ; encoding: [0x0a,0x0b,0xcf,0xd3,0x02,0x09,0x0e,0x0c]1563// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU1564v_smfmac_f32_32x32x64_fp8_bf8 v[10:25], a[2:5], v[4:11], v3 cbsz:3 abid:11565 1566// GFX950: v_smfmac_f32_32x32x64_fp8_bf8 v[10:25], a[2:5], v[4:11], v3 cbsz:3 abid:1 ; encoding: [0x0a,0x0b,0xcf,0xd3,0x02,0x09,0x0e,0x0c]1567// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU1568v_smfmac_f32_32x32x64fp8bf8 v[10:25], a[2:5], v[4:11], v3 cbsz:3 abid:11569 1570// GFX950: v_smfmac_f32_32x32x64_fp8_bf8 a[10:25], v[2:5], a[4:11], v1 ; encoding: [0x0a,0x80,0xcf,0xd3,0x02,0x09,0x06,0x14]1571// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU1572v_smfmac_f32_32x32x64_fp8_bf8 a[10:25], v[2:5], a[4:11], v11573 1574// GFX950: v_smfmac_f32_32x32x64_fp8_bf8 v[10:25], a[2:5], v[4:11], v2 cbsz:3 abid:1 ; encoding: [0x0a,0x0b,0xcf,0xd3,0x02,0x09,0x0a,0x0c]1575// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU1576v_smfmac_f32_32x32x64_fp8_bf8 v[10:25], a[2:5], v[4:11], v2 cbsz:3 abid:11577 1578// GFX950: v_smfmac_f32_32x32x64_fp8_bf8 a[10:25], v[2:5], a[4:11], v3 ; encoding: [0x0a,0x80,0xcf,0xd3,0x02,0x09,0x0e,0x14]1579// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU1580v_smfmac_f32_32x32x64_fp8_bf8 a[10:25], v[2:5], a[4:11], v31581 1582// GFX950: v_smfmac_f32_32x32x64_fp8_bf8 v[10:25], v[2:5], v[6:13], v2 cbsz:3 abid:1 ; encoding: [0x0a,0x0b,0xcf,0xd3,0x02,0x0d,0x0a,0x04]1583// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU1584v_smfmac_f32_32x32x64_fp8_bf8 v[10:25], v[2:5], v[6:13], v2 cbsz:3 abid:11585 1586// GFX950: v_smfmac_f32_32x32x64_fp8_bf8 a[10:25], a[2:5], a[6:13], v2 cbsz:3 abid:1 ; encoding: [0x0a,0x8b,0xcf,0xd3,0x02,0x0d,0x0a,0x1c]1587// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU1588v_smfmac_f32_32x32x64_fp8_bf8 a[10:25], a[2:5], a[6:13], v2 cbsz:3 abid:11589 1590// GFX950: v_smfmac_f32_32x32x64_fp8_bf8 a[10:25], a[2:5], a[6:13], v3 cbsz:1 abid:3 ; encoding: [0x0a,0x99,0xcf,0xd3,0x02,0x0d,0x0e,0x1c]1591// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU1592v_smfmac_f32_32x32x64_fp8_bf8 a[10:25], a[2:5], a[6:13], v3 cbsz:1 abid:31593 1594//===----------------------------------------------------------------------===//1595// v_smfmac_f32_32x32x64_fp8_fp81596//===----------------------------------------------------------------------===//1597 1598// GFX950: v_smfmac_f32_32x32x64_fp8_fp8 v[10:25], a[2:5], v[4:11], v3 cbsz:3 abid:1 ; encoding: [0x0a,0x0b,0xd3,0xd3,0x02,0x09,0x0e,0x0c]1599// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU1600v_smfmac_f32_32x32x64_fp8_fp8 v[10:25], a[2:5], v[4:11], v3 cbsz:3 abid:11601 1602// GFX950: v_smfmac_f32_32x32x64_fp8_fp8 v[10:25], a[2:5], v[4:11], v3 cbsz:3 abid:1 ; encoding: [0x0a,0x0b,0xd3,0xd3,0x02,0x09,0x0e,0x0c]1603// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU1604v_smfmac_f32_32x32x64fp8fp8 v[10:25], a[2:5], v[4:11], v3 cbsz:3 abid:11605 1606// GFX950: v_smfmac_f32_32x32x64_fp8_fp8 a[10:25], v[2:5], a[4:11], v1 ; encoding: [0x0a,0x80,0xd3,0xd3,0x02,0x09,0x06,0x14]1607// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU1608v_smfmac_f32_32x32x64_fp8_fp8 a[10:25], v[2:5], a[4:11], v11609 1610// GFX950: v_smfmac_f32_32x32x64_fp8_fp8 v[10:25], a[2:5], v[4:11], v2 cbsz:3 abid:1 ; encoding: [0x0a,0x0b,0xd3,0xd3,0x02,0x09,0x0a,0x0c]1611// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU1612v_smfmac_f32_32x32x64_fp8_fp8 v[10:25], a[2:5], v[4:11], v2 cbsz:3 abid:11613 1614// GFX950: v_smfmac_f32_32x32x64_fp8_fp8 a[10:25], v[2:5], a[4:11], v3 ; encoding: [0x0a,0x80,0xd3,0xd3,0x02,0x09,0x0e,0x14]1615// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU1616v_smfmac_f32_32x32x64_fp8_fp8 a[10:25], v[2:5], a[4:11], v31617 1618// GFX950: v_smfmac_f32_32x32x64_fp8_fp8 v[10:25], v[2:5], v[6:13], v2 cbsz:3 abid:1 ; encoding: [0x0a,0x0b,0xd3,0xd3,0x02,0x0d,0x0a,0x04]1619// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU1620v_smfmac_f32_32x32x64_fp8_fp8 v[10:25], v[2:5], v[6:13], v2 cbsz:3 abid:11621 1622// GFX950: v_smfmac_f32_32x32x64_fp8_fp8 a[10:25], a[2:5], a[6:13], v2 cbsz:3 abid:1 ; encoding: [0x0a,0x8b,0xd3,0xd3,0x02,0x0d,0x0a,0x1c]1623// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU1624v_smfmac_f32_32x32x64_fp8_fp8 a[10:25], a[2:5], a[6:13], v2 cbsz:3 abid:11625 1626// GFX950: v_smfmac_f32_32x32x64_fp8_fp8 a[10:25], a[2:5], a[6:13], v3 cbsz:1 abid:3 ; encoding: [0x0a,0x99,0xd3,0xd3,0x02,0x0d,0x0e,0x1c]1627// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU1628v_smfmac_f32_32x32x64_fp8_fp8 a[10:25], a[2:5], a[6:13], v3 cbsz:1 abid:31629