742 lines · plain
1// RUN: llvm-mc -triple=amdgcn -mcpu=gfx942 -show-encoding %s | FileCheck -check-prefix=GFX942 %s2// RUN: not llvm-mc -triple=amdgcn -mcpu=gfx90a %s 2>&1 | FileCheck -check-prefix=GFX90A %s3 4//===----------------------------------------------------------------------===//5// Misc opcodes.6//===----------------------------------------------------------------------===//7 8v_accvgpr_write_b32 a10, s209// GFX942: v_accvgpr_write_b32 a10, s20 ; encoding: [0x0a,0x40,0xd9,0xd3,0x14,0x00,0x00,0x18]10 11//===----------------------------------------------------------------------===//12// MFMA opcodes.13//===----------------------------------------------------------------------===//14 15v_mfma_f64_4x4x4_4b_f64 a[0:1], v[0:1], a[2:3], a[2:3]16// GFX942: v_mfma_f64_4x4x4_4b_f64 a[0:1], v[0:1], a[2:3], a[2:3] ; encoding: [0x00,0x80,0xef,0xd3,0x00,0x05,0x0a,0x14]17// GFX90A: :[[@LINE-2]]:{{[0-9]+}}: error: instruction not supported on this GPU18 19v_mfma_f64_4x4x4_4b_f64 v[0:1], v[0:1], a[2:3], v[2:3]20// GFX942: v_mfma_f64_4x4x4_4b_f64 v[0:1], v[0:1], a[2:3], v[2:3] ; encoding: [0x00,0x00,0xef,0xd3,0x00,0x05,0x0a,0x14]21// GFX90A: :[[@LINE-2]]:{{[0-9]+}}: error: instruction not supported on this GPU22 23v_mfma_f64_4x4x4f64 a[0:1], v[0:1], a[2:3], a[2:3]24// GFX942: v_mfma_f64_4x4x4_4b_f64 a[0:1], v[0:1], a[2:3], a[2:3] ; encoding: [0x00,0x80,0xef,0xd3,0x00,0x05,0x0a,0x14]25 26v_mfma_f64_4x4x4f64 v[0:1], v[0:1], a[2:3], v[2:3]27// GFX942: v_mfma_f64_4x4x4_4b_f64 v[0:1], v[0:1], a[2:3], v[2:3] ; encoding: [0x00,0x00,0xef,0xd3,0x00,0x05,0x0a,0x14]28 29v_mfma_f64_4x4x4_4b_f64 a[0:1], v[0:1], a[2:3], a[2:3] neg:[1,0,0]30// GFX942: v_mfma_f64_4x4x4_4b_f64 a[0:1], v[0:1], a[2:3], a[2:3] neg:[1,0,0] ; encoding: [0x00,0x80,0xef,0xd3,0x00,0x05,0x0a,0x34]31// GFX90A: :[[@LINE-2]]:{{[0-9]+}}: error: instruction not supported on this GPU32 33v_mfma_f64_4x4x4_4b_f64 a[0:1], v[0:1], a[2:3], a[2:3] neg:[0,1,0]34// GFX942: v_mfma_f64_4x4x4_4b_f64 a[0:1], v[0:1], a[2:3], a[2:3] neg:[0,1,0] ; encoding: [0x00,0x80,0xef,0xd3,0x00,0x05,0x0a,0x54]35// GFX90A: :[[@LINE-2]]:{{[0-9]+}}: error: instruction not supported on this GPU36 37v_mfma_f64_4x4x4_4b_f64 a[0:1], v[0:1], a[2:3], a[2:3] neg:[0,0,1]38// GFX942: v_mfma_f64_4x4x4_4b_f64 a[0:1], v[0:1], a[2:3], a[2:3] neg:[0,0,1] ; encoding: [0x00,0x80,0xef,0xd3,0x00,0x05,0x0a,0x94]39// GFX90A: :[[@LINE-2]]:{{[0-9]+}}: error: instruction not supported on this GPU40 41v_mfma_f64_4x4x4_4b_f64 v[0:1], v[0:1], a[2:3], v[2:3] neg:[1,1,1]42// GFX942: v_mfma_f64_4x4x4_4b_f64 v[0:1], v[0:1], a[2:3], v[2:3] neg:[1,1,1] ; encoding: [0x00,0x00,0xef,0xd3,0x00,0x05,0x0a,0xf4]43// GFX90A: :[[@LINE-2]]:{{[0-9]+}}: error: instruction not supported on this GPU44 45v_mfma_f64_4x4x4f64 a[0:1], v[0:1], a[2:3], a[2:3] neg:[1,0,0]46// GFX942: v_mfma_f64_4x4x4_4b_f64 a[0:1], v[0:1], a[2:3], a[2:3] neg:[1,0,0] ; encoding: [0x00,0x80,0xef,0xd3,0x00,0x05,0x0a,0x34]47// GFX90A: :[[@LINE-2]]:{{[0-9]+}}: error: invalid modifier: neg is not supported48 49v_mfma_f64_4x4x4f64 v[0:1], v[0:1], a[2:3], v[2:3] neg:[1,0,0]50// GFX942: v_mfma_f64_4x4x4_4b_f64 v[0:1], v[0:1], a[2:3], v[2:3] neg:[1,0,0] ; encoding: [0x00,0x00,0xef,0xd3,0x00,0x05,0x0a,0x34]51// GFX90A: :[[@LINE-2]]:{{[0-9]+}}: error: invalid modifier: neg is not supported52 53v_mfma_f64_16x16x4_f64 a[0:7], v[0:1], v[2:3], a[0:7]54// GFX942: v_mfma_f64_16x16x4_f64 a[0:7], v[0:1], v[2:3], a[0:7] ; encoding: [0x00,0x80,0xee,0xd3,0x00,0x05,0x02,0x04]55// GFX90A: :[[@LINE-2]]:{{[0-9]+}}: error: instruction not supported on this GPU56 57v_mfma_f64_16x16x4_f64 v[0:7], v[0:1], v[2:3], v[0:7]58// GFX942: v_mfma_f64_16x16x4_f64 v[0:7], v[0:1], v[2:3], v[0:7] ; encoding: [0x00,0x00,0xee,0xd3,0x00,0x05,0x02,0x04]59// GFX90A: :[[@LINE-2]]:{{[0-9]+}}: error: instruction not supported on this GPU60 61v_mfma_f64_16x16x4f64 a[0:7], v[0:1], v[2:3], a[0:7]62// GFX942: v_mfma_f64_16x16x4_f64 a[0:7], v[0:1], v[2:3], a[0:7] ; encoding: [0x00,0x80,0xee,0xd3,0x00,0x05,0x02,0x04]63 64v_mfma_f64_16x16x4f64 v[0:7], v[0:1], v[2:3], v[0:7]65// GFX942: v_mfma_f64_16x16x4_f64 v[0:7], v[0:1], v[2:3], v[0:7] ; encoding: [0x00,0x00,0xee,0xd3,0x00,0x05,0x02,0x04]66 67v_mfma_f64_16x16x4_f64 a[0:7], v[0:1], v[2:3], a[0:7] neg:[1,1,1]68// GFX942: v_mfma_f64_16x16x4_f64 a[0:7], v[0:1], v[2:3], a[0:7] neg:[1,1,1] ; encoding: [0x00,0x80,0xee,0xd3,0x00,0x05,0x02,0xe4]69// GFX90A: :[[@LINE-2]]:{{[0-9]+}}: error: instruction not supported on this GPU70 71v_mfma_f64_16x16x4_f64 v[0:7], v[0:1], v[2:3], v[0:7] neg:[1,1,1]72// GFX942: v_mfma_f64_16x16x4_f64 v[0:7], v[0:1], v[2:3], v[0:7] neg:[1,1,1] ; encoding: [0x00,0x00,0xee,0xd3,0x00,0x05,0x02,0xe4]73// GFX90A: :[[@LINE-2]]:{{[0-9]+}}: error: instruction not supported on this GPU74 75v_mfma_f64_16x16x4f64 a[0:7], v[0:1], v[2:3], a[0:7] neg:[1,0,0]76// GFX942: v_mfma_f64_16x16x4_f64 a[0:7], v[0:1], v[2:3], a[0:7] neg:[1,0,0] ; encoding: [0x00,0x80,0xee,0xd3,0x00,0x05,0x02,0x24]77// GFX90A: :[[@LINE-2]]:{{[0-9]+}}: error: invalid modifier: neg is not supported78 79v_mfma_f64_16x16x4f64 v[0:7], v[0:1], v[2:3], v[0:7] neg:[1,0,0]80// GFX942: v_mfma_f64_16x16x4_f64 v[0:7], v[0:1], v[2:3], v[0:7] neg:[1,0,0] ; encoding: [0x00,0x00,0xee,0xd3,0x00,0x05,0x02,0x24]81// GFX90A: :[[@LINE-2]]:{{[0-9]+}}: error: invalid modifier: neg is not supported82 83v_mfma_f32_16x16x1_4b_f32 a[0:15], v0, v1, a[18:33]84// GFX942: v_mfma_f32_16x16x1_4b_f32 a[0:15], v0, v1, a[18:33] ; encoding: [0x00,0x80,0xc1,0xd3,0x00,0x03,0x4a,0x04]85// GFX90A: :[[@LINE-2]]:{{[0-9]+}}: error: instruction not supported on this GPU86 87v_mfma_f32_16x16x1_4b_f32 v[0:15], v0, v1, v[18:33]88// GFX942: v_mfma_f32_16x16x1_4b_f32 v[0:15], v0, v1, v[18:33] ; encoding: [0x00,0x00,0xc1,0xd3,0x00,0x03,0x4a,0x04]89// GFX90A: :[[@LINE-2]]:{{[0-9]+}}: error: instruction not supported on this GPU90 91v_mfma_f32_16x16x1f32 a[0:15], v0, v1, a[18:33]92// GFX942: v_mfma_f32_16x16x1_4b_f32 a[0:15], v0, v1, a[18:33] ; encoding: [0x00,0x80,0xc1,0xd3,0x00,0x03,0x4a,0x04]93 94v_mfma_f32_16x16x1f32 v[0:15], v0, v1, v[18:33]95// GFX942: v_mfma_f32_16x16x1_4b_f32 v[0:15], v0, v1, v[18:33] ; encoding: [0x00,0x00,0xc1,0xd3,0x00,0x03,0x4a,0x04]96 97v_mfma_f32_4x4x1_16b_f32 a[0:3], v0, v1, a[2:5]98// GFX942: v_mfma_f32_4x4x1_16b_f32 a[0:3], v0, v1, a[2:5] ; encoding: [0x00,0x80,0xc2,0xd3,0x00,0x03,0x0a,0x04]99// GFX90A: :[[@LINE-2]]:{{[0-9]+}}: error: instruction not supported on this GPU100 101v_mfma_f32_4x4x1_16b_f32 v[0:3], v0, v1, v[2:5]102// GFX942: v_mfma_f32_4x4x1_16b_f32 v[0:3], v0, v1, v[2:5] ; encoding: [0x00,0x00,0xc2,0xd3,0x00,0x03,0x0a,0x04]103// GFX90A: :[[@LINE-2]]:{{[0-9]+}}: error: instruction not supported on this GPU104 105v_mfma_f32_4x4x1f32 a[0:3], v0, v1, a[2:5]106// GFX942: v_mfma_f32_4x4x1_16b_f32 a[0:3], v0, v1, a[2:5] ; encoding: [0x00,0x80,0xc2,0xd3,0x00,0x03,0x0a,0x04]107 108v_mfma_f32_4x4x1f32 v[0:3], v0, v1, v[2:5]109// GFX942: v_mfma_f32_4x4x1_16b_f32 v[0:3], v0, v1, v[2:5] ; encoding: [0x00,0x00,0xc2,0xd3,0x00,0x03,0x0a,0x04]110 111v_mfma_f32_32x32x2_f32 a[0:15], v0, v1, a[18:33]112// GFX942: v_mfma_f32_32x32x2_f32 a[0:15], v0, v1, a[18:33] ; encoding: [0x00,0x80,0xc4,0xd3,0x00,0x03,0x4a,0x04]113// GFX90A: :[[@LINE-2]]:{{[0-9]+}}: error: instruction not supported on this GPU114 115v_mfma_f32_32x32x2_f32 v[0:15], v0, v1, v[18:33]116// GFX942: v_mfma_f32_32x32x2_f32 v[0:15], v0, v1, v[18:33] ; encoding: [0x00,0x00,0xc4,0xd3,0x00,0x03,0x4a,0x04]117// GFX90A: :[[@LINE-2]]:{{[0-9]+}}: error: instruction not supported on this GPU118 119v_mfma_f32_32x32x2f32 a[0:15], v0, v1, a[18:33]120// GFX942: v_mfma_f32_32x32x2_f32 a[0:15], v0, v1, a[18:33] ; encoding: [0x00,0x80,0xc4,0xd3,0x00,0x03,0x4a,0x04]121 122v_mfma_f32_32x32x2f32 v[0:15], v0, v1, v[18:33]123// GFX942: v_mfma_f32_32x32x2_f32 v[0:15], v0, v1, v[18:33] ; encoding: [0x00,0x00,0xc4,0xd3,0x00,0x03,0x4a,0x04]124 125v_mfma_f32_16x16x4_f32 a[0:3], v0, v1, a[2:5]126// GFX942: v_mfma_f32_16x16x4_f32 a[0:3], v0, v1, a[2:5] ; encoding: [0x00,0x80,0xc5,0xd3,0x00,0x03,0x0a,0x04]127// GFX90A: :[[@LINE-2]]:{{[0-9]+}}: error: instruction not supported on this GPU128 129v_mfma_f32_16x16x4_f32 v[0:3], v0, v1, v[2:5]130// GFX942: v_mfma_f32_16x16x4_f32 v[0:3], v0, v1, v[2:5] ; encoding: [0x00,0x00,0xc5,0xd3,0x00,0x03,0x0a,0x04]131// GFX90A: :[[@LINE-2]]:{{[0-9]+}}: error: instruction not supported on this GPU132 133v_mfma_f32_16x16x4f32 a[0:3], v0, v1, a[2:5]134// GFX942: v_mfma_f32_16x16x4_f32 a[0:3], v0, v1, a[2:5] ; encoding: [0x00,0x80,0xc5,0xd3,0x00,0x03,0x0a,0x04]135 136v_mfma_f32_16x16x4f32 v[0:3], v0, v1, v[2:5]137// GFX942: v_mfma_f32_16x16x4_f32 v[0:3], v0, v1, v[2:5] ; encoding: [0x00,0x00,0xc5,0xd3,0x00,0x03,0x0a,0x04]138 139v_mfma_f32_32x32x4_2b_f16 a[0:31], v[0:1], v[2:3], a[34:65]140// GFX942: v_mfma_f32_32x32x4_2b_f16 a[0:31], v[0:1], v[2:3], a[34:65] ; encoding: [0x00,0x80,0xc8,0xd3,0x00,0x05,0x8a,0x04]141// GFX90A: :[[@LINE-2]]:{{[0-9]+}}: error: instruction not supported on this GPU142 143v_mfma_f32_32x32x4_2b_f16 v[0:31], v[0:1], v[2:3], v[34:65]144// GFX942: v_mfma_f32_32x32x4_2b_f16 v[0:31], v[0:1], v[2:3], v[34:65] ; encoding: [0x00,0x00,0xc8,0xd3,0x00,0x05,0x8a,0x04]145// GFX90A: :[[@LINE-2]]:{{[0-9]+}}: error: instruction not supported on this GPU146 147v_mfma_f32_32x32x4f16 a[0:31], v[0:1], v[2:3], a[34:65]148// GFX942: v_mfma_f32_32x32x4_2b_f16 a[0:31], v[0:1], v[2:3], a[34:65] ; encoding: [0x00,0x80,0xc8,0xd3,0x00,0x05,0x8a,0x04]149 150v_mfma_f32_32x32x4f16 v[0:31], v[0:1], v[2:3], v[34:65]151// GFX942: v_mfma_f32_32x32x4_2b_f16 v[0:31], v[0:1], v[2:3], v[34:65] ; encoding: [0x00,0x00,0xc8,0xd3,0x00,0x05,0x8a,0x04]152 153v_mfma_f32_16x16x4_4b_f16 a[0:15], v[0:1], v[2:3], a[18:33]154// GFX942: v_mfma_f32_16x16x4_4b_f16 a[0:15], v[0:1], v[2:3], a[18:33] ; encoding: [0x00,0x80,0xc9,0xd3,0x00,0x05,0x4a,0x04]155// GFX90A: :[[@LINE-2]]:{{[0-9]+}}: error: instruction not supported on this GPU156 157v_mfma_f32_16x16x4_4b_f16 v[0:15], v[0:1], v[2:3], v[18:33]158// GFX942: v_mfma_f32_16x16x4_4b_f16 v[0:15], v[0:1], v[2:3], v[18:33] ; encoding: [0x00,0x00,0xc9,0xd3,0x00,0x05,0x4a,0x04]159// GFX90A: :[[@LINE-2]]:{{[0-9]+}}: error: instruction not supported on this GPU160 161v_mfma_f32_16x16x4f16 a[0:15], v[0:1], v[2:3], a[18:33]162// GFX942: v_mfma_f32_16x16x4_4b_f16 a[0:15], v[0:1], v[2:3], a[18:33] ; encoding: [0x00,0x80,0xc9,0xd3,0x00,0x05,0x4a,0x04]163 164v_mfma_f32_16x16x4f16 v[0:15], v[0:1], v[2:3], v[18:33]165// GFX942: v_mfma_f32_16x16x4_4b_f16 v[0:15], v[0:1], v[2:3], v[18:33] ; encoding: [0x00,0x00,0xc9,0xd3,0x00,0x05,0x4a,0x04]166 167v_mfma_f32_4x4x4_16b_f16 a[0:3], v[0:1], v[2:3], a[2:5]168// GFX942: v_mfma_f32_4x4x4_16b_f16 a[0:3], v[0:1], v[2:3], a[2:5] ; encoding: [0x00,0x80,0xca,0xd3,0x00,0x05,0x0a,0x04]169// GFX90A: :[[@LINE-2]]:{{[0-9]+}}: error: instruction not supported on this GPU170 171v_mfma_f32_4x4x4_16b_f16 v[0:3], v[0:1], v[2:3], v[2:5]172// GFX942: v_mfma_f32_4x4x4_16b_f16 v[0:3], v[0:1], v[2:3], v[2:5] ; encoding: [0x00,0x00,0xca,0xd3,0x00,0x05,0x0a,0x04]173// GFX90A: :[[@LINE-2]]:{{[0-9]+}}: error: instruction not supported on this GPU174 175v_mfma_f32_4x4x4f16 a[0:3], v[0:1], v[2:3], a[2:5]176// GFX942: v_mfma_f32_4x4x4_16b_f16 a[0:3], v[0:1], v[2:3], a[2:5] ; encoding: [0x00,0x80,0xca,0xd3,0x00,0x05,0x0a,0x04]177 178v_mfma_f32_4x4x4f16 v[0:3], v[0:1], v[2:3], v[2:5]179// GFX942: v_mfma_f32_4x4x4_16b_f16 v[0:3], v[0:1], v[2:3], v[2:5] ; encoding: [0x00,0x00,0xca,0xd3,0x00,0x05,0x0a,0x04]180 181v_mfma_f32_32x32x8_f16 a[0:15], v[0:1], v[2:3], a[18:33]182// GFX942: v_mfma_f32_32x32x8_f16 a[0:15], v[0:1], v[2:3], a[18:33] ; encoding: [0x00,0x80,0xcc,0xd3,0x00,0x05,0x4a,0x04]183// GFX90A: :[[@LINE-2]]:{{[0-9]+}}: error: instruction not supported on this GPU184 185v_mfma_f32_32x32x8_f16 v[0:15], v[0:1], v[2:3], v[18:33]186// GFX942: v_mfma_f32_32x32x8_f16 v[0:15], v[0:1], v[2:3], v[18:33] ; encoding: [0x00,0x00,0xcc,0xd3,0x00,0x05,0x4a,0x04]187// GFX90A: :[[@LINE-2]]:{{[0-9]+}}: error: instruction not supported on this GPU188 189v_mfma_f32_32x32x8f16 a[0:15], v[0:1], v[2:3], a[18:33]190// GFX942: v_mfma_f32_32x32x8_f16 a[0:15], v[0:1], v[2:3], a[18:33] ; encoding: [0x00,0x80,0xcc,0xd3,0x00,0x05,0x4a,0x04]191 192v_mfma_f32_32x32x8f16 v[0:15], v[0:1], v[2:3], v[18:33]193// GFX942: v_mfma_f32_32x32x8_f16 v[0:15], v[0:1], v[2:3], v[18:33] ; encoding: [0x00,0x00,0xcc,0xd3,0x00,0x05,0x4a,0x04]194 195v_mfma_f32_16x16x16_f16 a[0:3], v[0:1], v[2:3], a[2:5]196// GFX942: v_mfma_f32_16x16x16_f16 a[0:3], v[0:1], v[2:3], a[2:5] ; encoding: [0x00,0x80,0xcd,0xd3,0x00,0x05,0x0a,0x04]197// GFX90A: :[[@LINE-2]]:{{[0-9]+}}: error: instruction not supported on this GPU198 199v_mfma_f32_16x16x16_f16 v[0:3], v[0:1], v[2:3], v[2:5]200// GFX942: v_mfma_f32_16x16x16_f16 v[0:3], v[0:1], v[2:3], v[2:5] ; encoding: [0x00,0x00,0xcd,0xd3,0x00,0x05,0x0a,0x04]201// GFX90A: :[[@LINE-2]]:{{[0-9]+}}: error: instruction not supported on this GPU202 203v_mfma_f32_16x16x16f16 a[0:3], v[0:1], v[2:3], a[2:5]204// GFX942: v_mfma_f32_16x16x16_f16 a[0:3], v[0:1], v[2:3], a[2:5] ; encoding: [0x00,0x80,0xcd,0xd3,0x00,0x05,0x0a,0x04]205 206v_mfma_f32_16x16x16f16 v[0:3], v[0:1], v[2:3], v[2:5]207// GFX942: v_mfma_f32_16x16x16_f16 v[0:3], v[0:1], v[2:3], v[2:5] ; encoding: [0x00,0x00,0xcd,0xd3,0x00,0x05,0x0a,0x04]208 209v_mfma_i32_32x32x4_2b_i8 a[0:31], v0, v1, a[34:65]210// GFX942: v_mfma_i32_32x32x4_2b_i8 a[0:31], v0, v1, a[34:65] ; encoding: [0x00,0x80,0xd0,0xd3,0x00,0x03,0x8a,0x04]211// GFX90A: :[[@LINE-2]]:{{[0-9]+}}: error: instruction not supported on this GPU212 213v_mfma_i32_32x32x4_2b_i8 v[0:31], v0, a1, v[34:65]214// GFX942: v_mfma_i32_32x32x4_2b_i8 v[0:31], v0, a1, v[34:65] ; encoding: [0x00,0x00,0xd0,0xd3,0x00,0x03,0x8a,0x14]215// GFX90A: :[[@LINE-2]]:{{[0-9]+}}: error: instruction not supported on this GPU216 217v_mfma_i32_32x32x4i8 a[0:31], v0, v1, a[34:65]218// GFX942: v_mfma_i32_32x32x4_2b_i8 a[0:31], v0, v1, a[34:65] ; encoding: [0x00,0x80,0xd0,0xd3,0x00,0x03,0x8a,0x04]219 220v_mfma_i32_32x32x4i8 v[0:31], v0, a1, v[34:65]221// GFX942: v_mfma_i32_32x32x4_2b_i8 v[0:31], v0, a1, v[34:65] ; encoding: [0x00,0x00,0xd0,0xd3,0x00,0x03,0x8a,0x14]222 223v_mfma_i32_16x16x4_4b_i8 a[0:15], v0, v1, a[18:33]224// GFX942: v_mfma_i32_16x16x4_4b_i8 a[0:15], v0, v1, a[18:33] ; encoding: [0x00,0x80,0xd1,0xd3,0x00,0x03,0x4a,0x04]225// GFX90A: :[[@LINE-2]]:{{[0-9]+}}: error: instruction not supported on this GPU226 227v_mfma_i32_16x16x4_4b_i8 v[0:15], v0, v1, v[18:33]228// GFX942: v_mfma_i32_16x16x4_4b_i8 v[0:15], v0, v1, v[18:33] ; encoding: [0x00,0x00,0xd1,0xd3,0x00,0x03,0x4a,0x04]229// GFX90A: :[[@LINE-2]]:{{[0-9]+}}: error: instruction not supported on this GPU230 231v_mfma_i32_16x16x4i8 a[0:15], v0, v1, a[18:33]232// GFX942: v_mfma_i32_16x16x4_4b_i8 a[0:15], v0, v1, a[18:33] ; encoding: [0x00,0x80,0xd1,0xd3,0x00,0x03,0x4a,0x04]233 234v_mfma_i32_16x16x4i8 v[0:15], v0, v1, v[18:33]235// GFX942: v_mfma_i32_16x16x4_4b_i8 v[0:15], v0, v1, v[18:33] ; encoding: [0x00,0x00,0xd1,0xd3,0x00,0x03,0x4a,0x04]236 237v_mfma_i32_4x4x4_16b_i8 a[0:3], v0, v1, a[2:5]238// GFX942: v_mfma_i32_4x4x4_16b_i8 a[0:3], v0, v1, a[2:5] ; encoding: [0x00,0x80,0xd2,0xd3,0x00,0x03,0x0a,0x04]239// GFX90A: :[[@LINE-2]]:{{[0-9]+}}: error: instruction not supported on this GPU240 241v_mfma_i32_4x4x4_16b_i8 v[0:3], v0, v1, v[2:5]242// GFX942: v_mfma_i32_4x4x4_16b_i8 v[0:3], v0, v1, v[2:5] ; encoding: [0x00,0x00,0xd2,0xd3,0x00,0x03,0x0a,0x04]243// GFX90A: :[[@LINE-2]]:{{[0-9]+}}: error: instruction not supported on this GPU244 245v_mfma_i32_4x4x4i8 a[0:3], v0, v1, a[2:5]246// GFX942: v_mfma_i32_4x4x4_16b_i8 a[0:3], v0, v1, a[2:5] ; encoding: [0x00,0x80,0xd2,0xd3,0x00,0x03,0x0a,0x04]247 248v_mfma_i32_4x4x4i8 v[0:3], v0, v1, v[2:5]249// GFX942: v_mfma_i32_4x4x4_16b_i8 v[0:3], v0, v1, v[2:5] ; encoding: [0x00,0x00,0xd2,0xd3,0x00,0x03,0x0a,0x04]250 251v_mfma_f32_32x32x1_2b_f32 a[0:31], v0, v1, a[34:65] blgp:7252// GFX942: v_mfma_f32_32x32x1_2b_f32 a[0:31], v0, v1, a[34:65] blgp:7 ; encoding: [0x00,0x80,0xc0,0xd3,0x00,0x03,0x8a,0xe4]253// GFX90A: :[[@LINE-2]]:{{[0-9]+}}: error: instruction not supported on this GPU254 255v_mfma_f32_32x32x1_2b_f32 v[0:31], v0, v1, v[34:65] blgp:7256// GFX942: v_mfma_f32_32x32x1_2b_f32 v[0:31], v0, v1, v[34:65] blgp:7 ; encoding: [0x00,0x00,0xc0,0xd3,0x00,0x03,0x8a,0xe4]257// GFX90A: :[[@LINE-2]]:{{[0-9]+}}: error: instruction not supported on this GPU258 259v_mfma_f32_32x32x1_2b_f32 a[0:31], v0, v1, a[34:65] blgp:7260// GFX942: v_mfma_f32_32x32x1_2b_f32 a[0:31], v0, v1, a[34:65] blgp:7 ; encoding: [0x00,0x80,0xc0,0xd3,0x00,0x03,0x8a,0xe4]261// GFX90A: :[[@LINE-2]]:{{[0-9]+}}: error: instruction not supported on this GPU262 263v_mfma_f32_32x32x1_2b_f32 v[0:31], v0, v1, v[34:65] blgp:7264// GFX942: v_mfma_f32_32x32x1_2b_f32 v[0:31], v0, v1, v[34:65] blgp:7 ; encoding: [0x00,0x00,0xc0,0xd3,0x00,0x03,0x8a,0xe4]265// GFX90A: :[[@LINE-2]]:{{[0-9]+}}: error: instruction not supported on this GPU266 267v_mfma_f32_32x32x1f32 a[0:31], v0, v1, a[34:65] blgp:7268// GFX942: v_mfma_f32_32x32x1_2b_f32 a[0:31], v0, v1, a[34:65] blgp:7 ; encoding: [0x00,0x80,0xc0,0xd3,0x00,0x03,0x8a,0xe4]269 270v_mfma_f32_32x32x1f32 v[0:31], v0, v1, v[34:65] blgp:7271// GFX942: v_mfma_f32_32x32x1_2b_f32 v[0:31], v0, v1, v[34:65] blgp:7 ; encoding: [0x00,0x00,0xc0,0xd3,0x00,0x03,0x8a,0xe4]272 273v_mfma_i32_32x32x16_i8 v[0:15], v[2:3], v[4:5], v[0:15]274// GFX942: v_mfma_i32_32x32x16_i8 v[0:15], v[2:3], v[4:5], v[0:15] ; encoding: [0x00,0x00,0xd6,0xd3,0x02,0x09,0x02,0x04]275// GFX90A: :[[@LINE-2]]:{{[0-9]+}}: error: instruction not supported on this GPU276 277v_mfma_i32_32x32x16_i8 a[0:15], v[2:3], v[4:5], a[0:15]278// GFX942: v_mfma_i32_32x32x16_i8 a[0:15], v[2:3], v[4:5], a[0:15] ; encoding: [0x00,0x80,0xd6,0xd3,0x02,0x09,0x02,0x04]279// GFX90A: :[[@LINE-2]]:{{[0-9]+}}: error: instruction not supported on this GPU280 281v_mfma_i32_32x32x16_i8 v[0:15], v[2:3], v[4:5], v[0:15]282// GFX942: v_mfma_i32_32x32x16_i8 v[0:15], v[2:3], v[4:5], v[0:15] ; encoding: [0x00,0x00,0xd6,0xd3,0x02,0x09,0x02,0x04]283// GFX90A: :[[@LINE-2]]:{{[0-9]+}}: error: instruction not supported on this GPU284 285v_mfma_i32_32x32x16_i8 a[0:15], v[2:3], v[4:5], a[0:15]286// GFX942: v_mfma_i32_32x32x16_i8 a[0:15], v[2:3], v[4:5], a[0:15] ; encoding: [0x00,0x80,0xd6,0xd3,0x02,0x09,0x02,0x04]287// GFX90A: :[[@LINE-2]]:{{[0-9]+}}: error: instruction not supported on this GPU288 289v_mfma_i32_32x32x16_i8 a[0:15], v[2:3], v[4:5], a[0:15] blgp:5290// GFX942: v_mfma_i32_32x32x16_i8 a[0:15], v[2:3], v[4:5], a[0:15] blgp:5 ; encoding: [0x00,0x80,0xd6,0xd3,0x02,0x09,0x02,0xa4]291// GFX90A: :[[@LINE-2]]:{{[0-9]+}}: error: instruction not supported on this GPU292 293v_mfma_i32_32x32x16i8 a[0:15], v[2:3], v[4:5], a[0:15] blgp:5294// GFX942: v_mfma_i32_32x32x16_i8 a[0:15], v[2:3], v[4:5], a[0:15] blgp:5 ; encoding: [0x00,0x80,0xd6,0xd3,0x02,0x09,0x02,0xa4]295// GFX90A: :[[@LINE-2]]:{{[0-9]+}}: error: instruction not supported on this GPU296 297v_mfma_i32_32x32x16i8 v[0:15], v[2:3], v[4:5], v[0:15] blgp:5298// GFX942: v_mfma_i32_32x32x16_i8 v[0:15], v[2:3], v[4:5], v[0:15] blgp:5 ; encoding: [0x00,0x00,0xd6,0xd3,0x02,0x09,0x02,0xa4]299// GFX90A: :[[@LINE-2]]:{{[0-9]+}}: error: instruction not supported on this GPU300 301v_mfma_i32_16x16x32_i8 v[0:3], v[2:3], v[4:5], v[0:3]302// GFX942: v_mfma_i32_16x16x32_i8 v[0:3], v[2:3], v[4:5], v[0:3] ; encoding: [0x00,0x00,0xd7,0xd3,0x02,0x09,0x02,0x04]303// GFX90A: :[[@LINE-2]]:{{[0-9]+}}: error: instruction not supported on this GPU304 305v_mfma_i32_16x16x32_i8 a[0:3], v[2:3], v[4:5], a[0:3]306// GFX942: v_mfma_i32_16x16x32_i8 a[0:3], v[2:3], v[4:5], a[0:3] ; encoding: [0x00,0x80,0xd7,0xd3,0x02,0x09,0x02,0x04]307// GFX90A: :[[@LINE-2]]:{{[0-9]+}}: error: instruction not supported on this GPU308 309v_mfma_i32_16x16x32_i8 a[0:3], v[2:3], v[4:5], a[0:3] blgp:5310// GFX942: v_mfma_i32_16x16x32_i8 a[0:3], v[2:3], v[4:5], a[0:3] blgp:5 ; encoding: [0x00,0x80,0xd7,0xd3,0x02,0x09,0x02,0xa4]311// GFX90A: :[[@LINE-2]]:{{[0-9]+}}: error: instruction not supported on this GPU312 313v_mfma_i32_16x16x32i8 v[0:3], v[2:3], v[4:5], v[0:3] blgp:5314// GFX942: v_mfma_i32_16x16x32_i8 v[0:3], v[2:3], v[4:5], v[0:3] blgp:5 ; encoding: [0x00,0x00,0xd7,0xd3,0x02,0x09,0x02,0xa4]315// GFX90A: :[[@LINE-2]]:{{[0-9]+}}: error: instruction not supported on this GPU316 317v_mfma_i32_16x16x32i8 a[0:3], v[2:3], v[4:5], a[0:3] blgp:5318// GFX942: v_mfma_i32_16x16x32_i8 a[0:3], v[2:3], v[4:5], a[0:3] blgp:5 ; encoding: [0x00,0x80,0xd7,0xd3,0x02,0x09,0x02,0xa4]319// GFX90A: :[[@LINE-2]]:{{[0-9]+}}: error: instruction not supported on this GPU320 321v_mfma_f32_32x32x4_2b_bf16 v[0:31], v[2:3], v[4:5], v[34:65]322// GFX942: v_mfma_f32_32x32x4_2b_bf16 v[0:31], v[2:3], v[4:5], v[34:65] ; encoding: [0x00,0x00,0xdd,0xd3,0x02,0x09,0x8a,0x04]323// GFX90A: :[[@LINE-2]]:{{[0-9]+}}: error: instruction not supported on this GPU324 325v_mfma_f32_32x32x4_2b_bf16 a[0:31], v[2:3], v[4:5], a[34:65]326// GFX942: v_mfma_f32_32x32x4_2b_bf16 a[0:31], v[2:3], v[4:5], a[34:65] ; encoding: [0x00,0x80,0xdd,0xd3,0x02,0x09,0x8a,0x04]327// GFX90A: :[[@LINE-2]]:{{[0-9]+}}: error: instruction not supported on this GPU328 329v_mfma_f32_32x32x4_2b_bf16 v[0:31], v[2:3], v[4:5], v[34:65]330// GFX942: v_mfma_f32_32x32x4_2b_bf16 v[0:31], v[2:3], v[4:5], v[34:65] ; encoding: [0x00,0x00,0xdd,0xd3,0x02,0x09,0x8a,0x04]331// GFX90A: :[[@LINE-2]]:{{[0-9]+}}: error: instruction not supported on this GPU332 333v_mfma_f32_32x32x4_2b_bf16 a[0:31], v[2:3], v[4:5], a[34:65]334// GFX942: v_mfma_f32_32x32x4_2b_bf16 a[0:31], v[2:3], v[4:5], a[34:65] ; encoding: [0x00,0x80,0xdd,0xd3,0x02,0x09,0x8a,0x04]335// GFX90A: :[[@LINE-2]]:{{[0-9]+}}: error: instruction not supported on this GPU336 337v_mfma_f32_32x32x4bf16 v[0:31], v[2:3], v[4:5], v[34:65] blgp:5338// GFX942: v_mfma_f32_32x32x4_2b_bf16 v[0:31], v[2:3], v[4:5], v[34:65] blgp:5 ; encoding: [0x00,0x00,0xdd,0xd3,0x02,0x09,0x8a,0xa4]339// GFX90A: :[[@LINE-2]]:{{[0-9]+}}: error: operands are not valid for this GPU or mode340 341v_mfma_f32_32x32x4bf16 a[0:31], v[2:3], v[4:5], a[34:65] blgp:5342// GFX942: v_mfma_f32_32x32x4_2b_bf16 a[0:31], v[2:3], v[4:5], a[34:65] blgp:5 ; encoding: [0x00,0x80,0xdd,0xd3,0x02,0x09,0x8a,0xa4]343// GFX90A: :[[@LINE-2]]:{{[0-9]+}}: error: operands are not valid for this GPU or mode344 345v_mfma_f32_32x32x4bf16_1k v[0:31], v[2:3], v[4:5], v[34:65] blgp:5346// GFX942: v_mfma_f32_32x32x4_2b_bf16 v[0:31], v[2:3], v[4:5], v[34:65] blgp:5 ; encoding: [0x00,0x00,0xdd,0xd3,0x02,0x09,0x8a,0xa4]347 348v_mfma_f32_32x32x4bf16_1k a[0:31], v[2:3], v[4:5], a[34:65] blgp:5349// GFX942: v_mfma_f32_32x32x4_2b_bf16 a[0:31], v[2:3], v[4:5], a[34:65] blgp:5 ; encoding: [0x00,0x80,0xdd,0xd3,0x02,0x09,0x8a,0xa4]350 351v_mfma_f32_16x16x4_4b_bf16 v[0:15], v[2:3], v[4:5], v[18:33]352// GFX942: v_mfma_f32_16x16x4_4b_bf16 v[0:15], v[2:3], v[4:5], v[18:33] ; encoding: [0x00,0x00,0xde,0xd3,0x02,0x09,0x4a,0x04]353// GFX90A: :[[@LINE-2]]:{{[0-9]+}}: error: instruction not supported on this GPU354 355v_mfma_f32_16x16x4_4b_bf16 a[0:15], v[2:3], v[4:5], a[18:33]356// GFX942: v_mfma_f32_16x16x4_4b_bf16 a[0:15], v[2:3], v[4:5], a[18:33] ; encoding: [0x00,0x80,0xde,0xd3,0x02,0x09,0x4a,0x04]357// GFX90A: :[[@LINE-2]]:{{[0-9]+}}: error: instruction not supported on this GPU358 359v_mfma_f32_16x16x4bf16 v[0:15], v[2:3], v[4:5], v[18:33] blgp:5360// GFX942: v_mfma_f32_16x16x4_4b_bf16 v[0:15], v[2:3], v[4:5], v[18:33] blgp:5 ; encoding: [0x00,0x00,0xde,0xd3,0x02,0x09,0x4a,0xa4]361// GFX90A: :[[@LINE-2]]:{{[0-9]+}}: error: instruction not supported on this GPU362 363v_mfma_f32_16x16x4bf16 a[0:15], v[2:3], v[4:5], a[18:33] blgp:5364// GFX942: v_mfma_f32_16x16x4_4b_bf16 a[0:15], v[2:3], v[4:5], a[18:33] blgp:5 ; encoding: [0x00,0x80,0xde,0xd3,0x02,0x09,0x4a,0xa4]365// GFX90A: :[[@LINE-2]]:{{[0-9]+}}: error: instruction not supported on this GPU366 367v_mfma_f32_16x16x4bf16_1k v[0:15], v[2:3], v[4:5], v[18:33] blgp:5368// GFX942: v_mfma_f32_16x16x4_4b_bf16 v[0:15], v[2:3], v[4:5], v[18:33] blgp:5 ; encoding: [0x00,0x00,0xde,0xd3,0x02,0x09,0x4a,0xa4]369 370v_mfma_f32_16x16x4bf16_1k a[0:15], v[2:3], v[4:5], a[18:33] blgp:5371// GFX942: v_mfma_f32_16x16x4_4b_bf16 a[0:15], v[2:3], v[4:5], a[18:33] blgp:5 ; encoding: [0x00,0x80,0xde,0xd3,0x02,0x09,0x4a,0xa4]372 373v_mfma_f32_4x4x4_16b_bf16 v[0:3], v[2:3], v[4:5], v[2:5]374// GFX942: v_mfma_f32_4x4x4_16b_bf16 v[0:3], v[2:3], v[4:5], v[2:5] ; encoding: [0x00,0x00,0xdf,0xd3,0x02,0x09,0x0a,0x04]375// GFX90A: :[[@LINE-2]]:{{[0-9]+}}: error: instruction not supported on this GPU376 377v_mfma_f32_4x4x4_16b_bf16 a[0:3], v[2:3], v[4:5], a[2:5]378// GFX942: v_mfma_f32_4x4x4_16b_bf16 a[0:3], v[2:3], v[4:5], a[2:5] ; encoding: [0x00,0x80,0xdf,0xd3,0x02,0x09,0x0a,0x04]379// GFX90A: :[[@LINE-2]]:{{[0-9]+}}: error: instruction not supported on this GPU380 381v_mfma_f32_4x4x4bf16 v[0:3], v[2:3], v[4:5], v[2:5]382// GFX942: v_mfma_f32_4x4x4_16b_bf16 v[0:3], v[2:3], v[4:5], v[2:5] ; encoding: [0x00,0x00,0xdf,0xd3,0x02,0x09,0x0a,0x04]383// GFX90A: :[[@LINE-2]]:{{[0-9]+}}: error: instruction not supported on this GPU384 385v_mfma_f32_4x4x4bf16 a[0:3], v[2:3], v[4:5], a[2:5]386// GFX942: v_mfma_f32_4x4x4_16b_bf16 a[0:3], v[2:3], v[4:5], a[2:5] ; encoding: [0x00,0x80,0xdf,0xd3,0x02,0x09,0x0a,0x04]387// GFX90A: :[[@LINE-2]]:{{[0-9]+}}: error: instruction not supported on this GPU388 389v_mfma_f32_4x4x4bf16_1k v[0:3], v[2:3], v[4:5], v[2:5]390// GFX942: v_mfma_f32_4x4x4_16b_bf16 v[0:3], v[2:3], v[4:5], v[2:5] ; encoding: [0x00,0x00,0xdf,0xd3,0x02,0x09,0x0a,0x04]391 392v_mfma_f32_4x4x4bf16_1k a[0:3], v[2:3], v[4:5], a[2:5]393// GFX942: v_mfma_f32_4x4x4_16b_bf16 a[0:3], v[2:3], v[4:5], a[2:5] ; encoding: [0x00,0x80,0xdf,0xd3,0x02,0x09,0x0a,0x04]394 395v_mfma_f32_32x32x8_bf16 v[0:15], v[2:3], v[4:5], v[18:33]396// GFX942: v_mfma_f32_32x32x8_bf16 v[0:15], v[2:3], v[4:5], v[18:33] ; encoding: [0x00,0x00,0xe0,0xd3,0x02,0x09,0x4a,0x04]397// GFX90A: :[[@LINE-2]]:{{[0-9]+}}: error: instruction not supported on this GPU398 399v_mfma_f32_32x32x8_bf16 a[0:15], v[2:3], v[4:5], a[18:33]400// GFX942: v_mfma_f32_32x32x8_bf16 a[0:15], v[2:3], v[4:5], a[18:33] ; encoding: [0x00,0x80,0xe0,0xd3,0x02,0x09,0x4a,0x04]401// GFX90A: :[[@LINE-2]]:{{[0-9]+}}: error: instruction not supported on this GPU402 403v_mfma_f32_32x32x8bf16 v[0:15], v[2:3], v[4:5], v[18:33]404// GFX942: v_mfma_f32_32x32x8_bf16 v[0:15], v[2:3], v[4:5], v[18:33] ; encoding: [0x00,0x00,0xe0,0xd3,0x02,0x09,0x4a,0x04]405// GFX90A: :[[@LINE-2]]:{{[0-9]+}}: error: instruction not supported on this GPU406 407v_mfma_f32_32x32x8bf16 a[0:15], v[2:3], v[4:5], a[18:33]408// GFX942: v_mfma_f32_32x32x8_bf16 a[0:15], v[2:3], v[4:5], a[18:33] ; encoding: [0x00,0x80,0xe0,0xd3,0x02,0x09,0x4a,0x04]409// GFX90A: :[[@LINE-2]]:{{[0-9]+}}: error: instruction not supported on this GPU410 411v_mfma_f32_32x32x8bf16_1k v[0:15], v[2:3], v[4:5], v[18:33]412// GFX942: v_mfma_f32_32x32x8_bf16 v[0:15], v[2:3], v[4:5], v[18:33] ; encoding: [0x00,0x00,0xe0,0xd3,0x02,0x09,0x4a,0x04]413 414v_mfma_f32_32x32x8bf16_1k a[0:15], v[2:3], v[4:5], a[18:33]415// GFX942: v_mfma_f32_32x32x8_bf16 a[0:15], v[2:3], v[4:5], a[18:33] ; encoding: [0x00,0x80,0xe0,0xd3,0x02,0x09,0x4a,0x04]416 417v_mfma_f32_16x16x16_bf16 v[0:3], v[2:3], v[4:5], v[2:5]418// GFX942: v_mfma_f32_16x16x16_bf16 v[0:3], v[2:3], v[4:5], v[2:5] ; encoding: [0x00,0x00,0xe1,0xd3,0x02,0x09,0x0a,0x04]419// GFX90A: :[[@LINE-2]]:{{[0-9]+}}: error: instruction not supported on this GPU420 421v_mfma_f32_16x16x16_bf16 a[0:3], v[2:3], v[4:5], a[2:5]422// GFX942: v_mfma_f32_16x16x16_bf16 a[0:3], v[2:3], v[4:5], a[2:5] ; encoding: [0x00,0x80,0xe1,0xd3,0x02,0x09,0x0a,0x04]423// GFX90A: :[[@LINE-2]]:{{[0-9]+}}: error: instruction not supported on this GPU424 425v_mfma_f32_16x16x16bf16 v[0:3], v[2:3], v[4:5], v[2:5]426// GFX942: v_mfma_f32_16x16x16_bf16 v[0:3], v[2:3], v[4:5], v[2:5] ; encoding: [0x00,0x00,0xe1,0xd3,0x02,0x09,0x0a,0x04]427// GFX90A: :[[@LINE-2]]:{{[0-9]+}}: error: instruction not supported on this GPU428 429v_mfma_f32_16x16x16bf16 a[0:3], v[2:3], v[4:5], a[2:5]430// GFX942: v_mfma_f32_16x16x16_bf16 a[0:3], v[2:3], v[4:5], a[2:5] ; encoding: [0x00,0x80,0xe1,0xd3,0x02,0x09,0x0a,0x04]431// GFX90A: :[[@LINE-2]]:{{[0-9]+}}: error: instruction not supported on this GPU432 433v_mfma_f32_16x16x16bf16_1k v[0:3], v[2:3], v[4:5], v[2:5]434// GFX942: v_mfma_f32_16x16x16_bf16 v[0:3], v[2:3], v[4:5], v[2:5] ; encoding: [0x00,0x00,0xe1,0xd3,0x02,0x09,0x0a,0x04]435 436v_mfma_f32_16x16x16bf16_1k a[0:3], v[2:3], v[4:5], a[2:5]437// GFX942: v_mfma_f32_16x16x16_bf16 a[0:3], v[2:3], v[4:5], a[2:5] ; encoding: [0x00,0x80,0xe1,0xd3,0x02,0x09,0x0a,0x04]438 439v_mfma_f32_16x16x8_xf32 a[0:3], v[2:3], v[4:5], a[2:5]440// GFX942: v_mfma_f32_16x16x8_xf32 a[0:3], v[2:3], v[4:5], a[2:5] ; encoding: [0x00,0x80,0xbe,0xd3,0x02,0x09,0x0a,0x04]441// GFX90A: :[[@LINE-2]]:{{[0-9]+}}: error: instruction not supported on this GPU442 443v_mfma_f32_16x16x8_xf32 v[0:3], v[2:3], v[4:5], v[2:5]444// GFX942: v_mfma_f32_16x16x8_xf32 v[0:3], v[2:3], v[4:5], v[2:5] ; encoding: [0x00,0x00,0xbe,0xd3,0x02,0x09,0x0a,0x04]445// GFX90A: :[[@LINE-2]]:{{[0-9]+}}: error: instruction not supported on this GPU446 447v_mfma_f32_16x16x8xf32 a[0:3], v[2:3], v[4:5], a[2:5]448// GFX942: v_mfma_f32_16x16x8_xf32 a[0:3], v[2:3], v[4:5], a[2:5] ; encoding: [0x00,0x80,0xbe,0xd3,0x02,0x09,0x0a,0x04]449// GFX90A: :[[@LINE-2]]:{{[0-9]+}}: error: instruction not supported on this GPU450 451v_mfma_f32_16x16x8xf32 v[0:3], v[2:3], v[4:5], v[2:5]452// GFX942: v_mfma_f32_16x16x8_xf32 v[0:3], v[2:3], v[4:5], v[2:5] ; encoding: [0x00,0x00,0xbe,0xd3,0x02,0x09,0x0a,0x04]453// GFX90A: :[[@LINE-2]]:{{[0-9]+}}: error: instruction not supported on this GPU454 455v_mfma_f32_32x32x4_xf32 v[0:15], v[2:3], v[4:5], v[18:33]456// GFX942: v_mfma_f32_32x32x4_xf32 v[0:15], v[2:3], v[4:5], v[18:33] ; encoding: [0x00,0x00,0xbf,0xd3,0x02,0x09,0x4a,0x04]457// GFX90A: :[[@LINE-2]]:{{[0-9]+}}: error: instruction not supported on this GPU458 459v_mfma_f32_32x32x4_xf32 a[0:15], v[2:3], v[4:5], a[18:33]460// GFX942: v_mfma_f32_32x32x4_xf32 a[0:15], v[2:3], v[4:5], a[18:33] ; encoding: [0x00,0x80,0xbf,0xd3,0x02,0x09,0x4a,0x04]461// GFX90A: :[[@LINE-2]]:{{[0-9]+}}: error: instruction not supported on this GPU462 463v_mfma_f32_32x32x4xf32 v[0:15], v[2:3], v[4:5], v[18:33]464// GFX942: v_mfma_f32_32x32x4_xf32 v[0:15], v[2:3], v[4:5], v[18:33] ; encoding: [0x00,0x00,0xbf,0xd3,0x02,0x09,0x4a,0x04]465// GFX90A: :[[@LINE-2]]:{{[0-9]+}}: error: instruction not supported on this GPU466 467v_mfma_f32_32x32x4xf32 a[0:15], v[2:3], v[4:5], a[18:33]468// GFX942: v_mfma_f32_32x32x4_xf32 a[0:15], v[2:3], v[4:5], a[18:33] ; encoding: [0x00,0x80,0xbf,0xd3,0x02,0x09,0x4a,0x04]469// GFX90A: :[[@LINE-2]]:{{[0-9]+}}: error: instruction not supported on this GPU470 471v_mfma_f32_16x16x32_bf8_bf8 v[0:3], v[2:3], v[4:5], v[0:3]472// GFX942: v_mfma_f32_16x16x32_bf8_bf8 v[0:3], v[2:3], v[4:5], v[0:3] ; encoding: [0x00,0x00,0xf0,0xd3,0x02,0x09,0x02,0x04]473// GFX90A: :[[@LINE-2]]:{{[0-9]+}}: error: instruction not supported on this GPU474 475v_mfma_f32_16x16x32_bf8_bf8 a[0:3], v[2:3], v[4:5], a[0:3]476// GFX942: v_mfma_f32_16x16x32_bf8_bf8 a[0:3], v[2:3], v[4:5], a[0:3] ; encoding: [0x00,0x80,0xf0,0xd3,0x02,0x09,0x02,0x04]477// GFX90A: :[[@LINE-2]]:{{[0-9]+}}: error: instruction not supported on this GPU478 479v_mfma_f32_16x16x32_bf8_bf8 a[0:3], v[2:3], v[4:5], a[0:3] blgp:5480// GFX942: v_mfma_f32_16x16x32_bf8_bf8 a[0:3], v[2:3], v[4:5], a[0:3] blgp:5 ; encoding: [0x00,0x80,0xf0,0xd3,0x02,0x09,0x02,0xa4]481// GFX90A: :[[@LINE-2]]:{{[0-9]+}}: error: instruction not supported on this GPU482 483v_mfma_f32_16x16x32_bf8_fp8 v[0:3], v[2:3], v[4:5], v[0:3]484// GFX942: v_mfma_f32_16x16x32_bf8_fp8 v[0:3], v[2:3], v[4:5], v[0:3] ; encoding: [0x00,0x00,0xf1,0xd3,0x02,0x09,0x02,0x04]485// GFX90A: :[[@LINE-2]]:{{[0-9]+}}: error: instruction not supported on this GPU486 487v_mfma_f32_16x16x32_bf8_fp8 a[0:3], v[2:3], v[4:5], a[0:3]488// GFX942: v_mfma_f32_16x16x32_bf8_fp8 a[0:3], v[2:3], v[4:5], a[0:3] ; encoding: [0x00,0x80,0xf1,0xd3,0x02,0x09,0x02,0x04]489// GFX90A: :[[@LINE-2]]:{{[0-9]+}}: error: instruction not supported on this GPU490 491v_mfma_f32_16x16x32_bf8_fp8 a[0:3], v[2:3], v[4:5], a[0:3] blgp:5492// GFX942: v_mfma_f32_16x16x32_bf8_fp8 a[0:3], v[2:3], v[4:5], a[0:3] blgp:5 ; encoding: [0x00,0x80,0xf1,0xd3,0x02,0x09,0x02,0xa4]493// GFX90A: :[[@LINE-2]]:{{[0-9]+}}: error: instruction not supported on this GPU494 495v_mfma_f32_16x16x32_fp8_bf8 v[0:3], v[2:3], v[4:5], v[0:3]496// GFX942: v_mfma_f32_16x16x32_fp8_bf8 v[0:3], v[2:3], v[4:5], v[0:3] ; encoding: [0x00,0x00,0xf2,0xd3,0x02,0x09,0x02,0x04]497// GFX90A: :[[@LINE-2]]:{{[0-9]+}}: error: instruction not supported on this GPU498 499v_mfma_f32_16x16x32_fp8_bf8 a[0:3], v[2:3], v[4:5], a[0:3]500// GFX942: v_mfma_f32_16x16x32_fp8_bf8 a[0:3], v[2:3], v[4:5], a[0:3] ; encoding: [0x00,0x80,0xf2,0xd3,0x02,0x09,0x02,0x04]501// GFX90A: :[[@LINE-2]]:{{[0-9]+}}: error: instruction not supported on this GPU502 503v_mfma_f32_16x16x32_fp8_bf8 a[0:3], v[2:3], v[4:5], a[0:3] blgp:5504// GFX942: v_mfma_f32_16x16x32_fp8_bf8 a[0:3], v[2:3], v[4:5], a[0:3] blgp:5 ; encoding: [0x00,0x80,0xf2,0xd3,0x02,0x09,0x02,0xa4]505// GFX90A: :[[@LINE-2]]:{{[0-9]+}}: error: instruction not supported on this GPU506 507v_mfma_f32_16x16x32_fp8_fp8 v[0:3], v[2:3], v[4:5], v[0:3]508// GFX942: v_mfma_f32_16x16x32_fp8_fp8 v[0:3], v[2:3], v[4:5], v[0:3] ; encoding: [0x00,0x00,0xf3,0xd3,0x02,0x09,0x02,0x04]509// GFX90A: :[[@LINE-2]]:{{[0-9]+}}: error: instruction not supported on this GPU510 511v_mfma_f32_16x16x32_fp8_fp8 a[0:3], v[2:3], v[4:5], a[0:3]512// GFX942: v_mfma_f32_16x16x32_fp8_fp8 a[0:3], v[2:3], v[4:5], a[0:3] ; encoding: [0x00,0x80,0xf3,0xd3,0x02,0x09,0x02,0x04]513// GFX90A: :[[@LINE-2]]:{{[0-9]+}}: error: instruction not supported on this GPU514 515v_mfma_f32_16x16x32_fp8_fp8 a[0:3], v[2:3], v[4:5], a[0:3] blgp:5516// GFX942: v_mfma_f32_16x16x32_fp8_fp8 a[0:3], v[2:3], v[4:5], a[0:3] blgp:5 ; encoding: [0x00,0x80,0xf3,0xd3,0x02,0x09,0x02,0xa4]517// GFX90A: :[[@LINE-2]]:{{[0-9]+}}: error: instruction not supported on this GPU518 519v_mfma_f32_32x32x16_bf8_bf8 v[0:15], v[2:3], v[4:5], v[0:15]520// GFX942: v_mfma_f32_32x32x16_bf8_bf8 v[0:15], v[2:3], v[4:5], v[0:15] ; encoding: [0x00,0x00,0xf4,0xd3,0x02,0x09,0x02,0x04]521// GFX90A: :[[@LINE-2]]:{{[0-9]+}}: error: instruction not supported on this GPU522 523v_mfma_f32_32x32x16_bf8_bf8 a[0:15], v[2:3], v[4:5], a[0:15]524// GFX942: v_mfma_f32_32x32x16_bf8_bf8 a[0:15], v[2:3], v[4:5], a[0:15] ; encoding: [0x00,0x80,0xf4,0xd3,0x02,0x09,0x02,0x04]525// GFX90A: :[[@LINE-2]]:{{[0-9]+}}: error: instruction not supported on this GPU526 527v_mfma_f32_32x32x16_bf8_bf8 a[0:15], v[2:3], v[4:5], a[0:15] blgp:5528// GFX942: v_mfma_f32_32x32x16_bf8_bf8 a[0:15], v[2:3], v[4:5], a[0:15] blgp:5 ; encoding: [0x00,0x80,0xf4,0xd3,0x02,0x09,0x02,0xa4]529// GFX90A: :[[@LINE-2]]:{{[0-9]+}}: error: instruction not supported on this GPU530 531v_mfma_f32_32x32x16_bf8_fp8 v[0:15], v[2:3], v[4:5], v[0:15]532// GFX942: v_mfma_f32_32x32x16_bf8_fp8 v[0:15], v[2:3], v[4:5], v[0:15] ; encoding: [0x00,0x00,0xf5,0xd3,0x02,0x09,0x02,0x04]533// GFX90A: :[[@LINE-2]]:{{[0-9]+}}: error: instruction not supported on this GPU534 535v_mfma_f32_32x32x16_bf8_fp8 a[0:15], v[2:3], v[4:5], a[0:15]536// GFX942: v_mfma_f32_32x32x16_bf8_fp8 a[0:15], v[2:3], v[4:5], a[0:15] ; encoding: [0x00,0x80,0xf5,0xd3,0x02,0x09,0x02,0x04]537// GFX90A: :[[@LINE-2]]:{{[0-9]+}}: error: instruction not supported on this GPU538 539v_mfma_f32_32x32x16_bf8_fp8 a[0:15], v[2:3], v[4:5], a[0:15] blgp:5540// GFX942: v_mfma_f32_32x32x16_bf8_fp8 a[0:15], v[2:3], v[4:5], a[0:15] blgp:5 ; encoding: [0x00,0x80,0xf5,0xd3,0x02,0x09,0x02,0xa4]541// GFX90A: :[[@LINE-2]]:{{[0-9]+}}: error: instruction not supported on this GPU542 543v_mfma_f32_32x32x16_fp8_bf8 v[0:15], v[2:3], v[4:5], v[0:15]544// GFX942: v_mfma_f32_32x32x16_fp8_bf8 v[0:15], v[2:3], v[4:5], v[0:15] ; encoding: [0x00,0x00,0xf6,0xd3,0x02,0x09,0x02,0x04]545// GFX90A: :[[@LINE-2]]:{{[0-9]+}}: error: instruction not supported on this GPU546 547v_mfma_f32_32x32x16_fp8_bf8 a[0:15], v[2:3], v[4:5], a[0:15]548// GFX942: v_mfma_f32_32x32x16_fp8_bf8 a[0:15], v[2:3], v[4:5], a[0:15] ; encoding: [0x00,0x80,0xf6,0xd3,0x02,0x09,0x02,0x04]549// GFX90A: :[[@LINE-2]]:{{[0-9]+}}: error: instruction not supported on this GPU550 551v_mfma_f32_32x32x16_fp8_bf8 a[0:15], v[2:3], v[4:5], a[0:15] blgp:5552// GFX942: v_mfma_f32_32x32x16_fp8_bf8 a[0:15], v[2:3], v[4:5], a[0:15] blgp:5 ; encoding: [0x00,0x80,0xf6,0xd3,0x02,0x09,0x02,0xa4]553// GFX90A: :[[@LINE-2]]:{{[0-9]+}}: error: instruction not supported on this GPU554 555v_mfma_f32_32x32x16_fp8_fp8 v[0:15], v[2:3], v[4:5], v[0:15]556// GFX942: v_mfma_f32_32x32x16_fp8_fp8 v[0:15], v[2:3], v[4:5], v[0:15] ; encoding: [0x00,0x00,0xf7,0xd3,0x02,0x09,0x02,0x04]557// GFX90A: :[[@LINE-2]]:{{[0-9]+}}: error: instruction not supported on this GPU558 559v_mfma_f32_32x32x16_fp8_fp8 a[0:15], v[2:3], v[4:5], a[0:15]560// GFX942: v_mfma_f32_32x32x16_fp8_fp8 a[0:15], v[2:3], v[4:5], a[0:15] ; encoding: [0x00,0x80,0xf7,0xd3,0x02,0x09,0x02,0x04]561// GFX90A: :[[@LINE-2]]:{{[0-9]+}}: error: instruction not supported on this GPU562 563v_mfma_f32_32x32x16_fp8_fp8 a[0:15], v[2:3], v[4:5], a[0:15] blgp:5564// GFX942: v_mfma_f32_32x32x16_fp8_fp8 a[0:15], v[2:3], v[4:5], a[0:15] blgp:5 ; encoding: [0x00,0x80,0xf7,0xd3,0x02,0x09,0x02,0xa4]565// GFX90A: :[[@LINE-2]]:{{[0-9]+}}: error: instruction not supported on this GPU566 567//===----------------------------------------------------------------------===//568// SMFMAC opcodes.569//===----------------------------------------------------------------------===//570 571v_smfmac_f32_16x16x32_f16 v[10:13], a[2:3], v[4:7], v0 cbsz:3 abid:1572// GFX942: v_smfmac_f32_16x16x32_f16 v[10:13], a[2:3], v[4:7], v0 cbsz:3 abid:1 ; encoding: [0x0a,0x0b,0xe2,0xd3,0x02,0x09,0x02,0x0c]573// GFX90A: :[[@LINE-2]]:{{[0-9]+}}: error: instruction not supported on this GPU574 575v_smfmac_f32_16x16x32_f16 a[10:13], v[2:3], a[4:7], v1576// GFX942: v_smfmac_f32_16x16x32_f16 a[10:13], v[2:3], a[4:7], v1 ; encoding: [0x0a,0x80,0xe2,0xd3,0x02,0x09,0x06,0x14]577// GFX90A: :[[@LINE-2]]:{{[0-9]+}}: error: instruction not supported on this GPU578 579v_smfmac_f32_32x32x16_f16 v[10:25], a[2:3], v[4:7], v2 cbsz:3 abid:1580// GFX942: v_smfmac_f32_32x32x16_f16 v[10:25], a[2:3], v[4:7], v2 cbsz:3 abid:1 ; encoding: [0x0a,0x0b,0xe4,0xd3,0x02,0x09,0x0a,0x0c]581// GFX90A: :[[@LINE-2]]:{{[0-9]+}}: error: instruction not supported on this GPU582 583v_smfmac_f32_32x32x16_f16 a[10:25], v[2:3], a[4:7], v3584// GFX942: v_smfmac_f32_32x32x16_f16 a[10:25], v[2:3], a[4:7], v3 ; encoding: [0x0a,0x80,0xe4,0xd3,0x02,0x09,0x0e,0x14]585// GFX90A: :[[@LINE-2]]:{{[0-9]+}}: error: instruction not supported on this GPU586 587v_smfmac_f32_16x16x32_bf16 v[10:13], a[2:3], v[4:7], v4 cbsz:3 abid:1588// GFX942: v_smfmac_f32_16x16x32_bf16 v[10:13], a[2:3], v[4:7], v4 cbsz:3 abid:1 ; encoding: [0x0a,0x0b,0xe6,0xd3,0x02,0x09,0x12,0x0c]589// GFX90A: :[[@LINE-2]]:{{[0-9]+}}: error: instruction not supported on this GPU590 591v_smfmac_f32_16x16x32_bf16 a[10:13], v[2:3], a[4:7], v5592// GFX942: v_smfmac_f32_16x16x32_bf16 a[10:13], v[2:3], a[4:7], v5 ; encoding: [0x0a,0x80,0xe6,0xd3,0x02,0x09,0x16,0x14]593// GFX90A: :[[@LINE-2]]:{{[0-9]+}}: error: instruction not supported on this GPU594 595v_smfmac_f32_32x32x16_bf16 v[10:25], a[2:3], v[4:7], v6 cbsz:3 abid:1596// GFX942: v_smfmac_f32_32x32x16_bf16 v[10:25], a[2:3], v[4:7], v6 cbsz:3 abid:1 ; encoding: [0x0a,0x0b,0xe8,0xd3,0x02,0x09,0x1a,0x0c]597// GFX90A: :[[@LINE-2]]:{{[0-9]+}}: error: instruction not supported on this GPU598 599v_smfmac_f32_32x32x16_bf16 a[10:25], v[2:3], a[4:7], v7600// GFX942: v_smfmac_f32_32x32x16_bf16 a[10:25], v[2:3], a[4:7], v7 ; encoding: [0x0a,0x80,0xe8,0xd3,0x02,0x09,0x1e,0x14]601// GFX90A: :[[@LINE-2]]:{{[0-9]+}}: error: instruction not supported on this GPU602 603v_smfmac_i32_16x16x64_i8 v[10:13], a[2:3], v[4:7], v8 cbsz:3 abid:1604// GFX942: v_smfmac_i32_16x16x64_i8 v[10:13], a[2:3], v[4:7], v8 cbsz:3 abid:1 ; encoding: [0x0a,0x0b,0xea,0xd3,0x02,0x09,0x22,0x0c]605// GFX90A: :[[@LINE-2]]:{{[0-9]+}}: error: instruction not supported on this GPU606 607v_smfmac_i32_16x16x64_i8 a[10:13], v[2:3], a[4:7], v9608// GFX942: v_smfmac_i32_16x16x64_i8 a[10:13], v[2:3], a[4:7], v9 ; encoding: [0x0a,0x80,0xea,0xd3,0x02,0x09,0x26,0x14]609// GFX90A: :[[@LINE-2]]:{{[0-9]+}}: error: instruction not supported on this GPU610 611v_smfmac_i32_32x32x32_i8 v[10:25], a[2:3], v[4:7], v10 cbsz:3 abid:1612// GFX942: v_smfmac_i32_32x32x32_i8 v[10:25], a[2:3], v[4:7], v10 cbsz:3 abid:1 ; encoding: [0x0a,0x0b,0xec,0xd3,0x02,0x09,0x2a,0x0c]613// GFX90A: :[[@LINE-2]]:{{[0-9]+}}: error: instruction not supported on this GPU614 615v_smfmac_i32_32x32x32_i8 a[10:25], v[2:3], a[4:7], v11616// GFX942: v_smfmac_i32_32x32x32_i8 a[10:25], v[2:3], a[4:7], v11 ; encoding: [0x0a,0x80,0xec,0xd3,0x02,0x09,0x2e,0x14]617// GFX90A: :[[@LINE-2]]:{{[0-9]+}}: error: instruction not supported on this GPU618 619v_smfmac_f32_16x16x64_bf8_bf8 v[0:3], a[2:3], v[4:7], v1 cbsz:3 abid:1620// GFX942: v_smfmac_f32_16x16x64_bf8_bf8 v[0:3], a[2:3], v[4:7], v1 cbsz:3 abid:1 ; encoding: [0x00,0x0b,0xf8,0xd3,0x02,0x09,0x06,0x0c]621// GFX90A: :[[@LINE-2]]:{{[0-9]+}}: error: instruction not supported on this GPU622 623v_smfmac_f32_16x16x64_bf8_bf8 a[0:3], v[2:3], a[4:7], v1624// GFX942: v_smfmac_f32_16x16x64_bf8_bf8 a[0:3], v[2:3], a[4:7], v1 ; encoding: [0x00,0x80,0xf8,0xd3,0x02,0x09,0x06,0x14]625// GFX90A: :[[@LINE-2]]:{{[0-9]+}}: error: instruction not supported on this GPU626 627v_smfmac_f32_16x16x64_bf8_fp8 v[0:3], a[2:3], v[4:7], v1 cbsz:3 abid:1628// GFX942: v_smfmac_f32_16x16x64_bf8_fp8 v[0:3], a[2:3], v[4:7], v1 cbsz:3 abid:1 ; encoding: [0x00,0x0b,0xf9,0xd3,0x02,0x09,0x06,0x0c]629// GFX90A: :[[@LINE-2]]:{{[0-9]+}}: error: instruction not supported on this GPU630 631v_smfmac_f32_16x16x64_bf8_fp8 a[0:3], v[2:3], a[4:7], v1632// GFX942: v_smfmac_f32_16x16x64_bf8_fp8 a[0:3], v[2:3], a[4:7], v1 ; encoding: [0x00,0x80,0xf9,0xd3,0x02,0x09,0x06,0x14]633// GFX90A: :[[@LINE-2]]:{{[0-9]+}}: error: instruction not supported on this GPU634 635v_smfmac_f32_16x16x64_fp8_bf8 v[0:3], a[2:3], v[4:7], v1 cbsz:3 abid:1636// GFX942: v_smfmac_f32_16x16x64_fp8_bf8 v[0:3], a[2:3], v[4:7], v1 cbsz:3 abid:1 ; encoding: [0x00,0x0b,0xfa,0xd3,0x02,0x09,0x06,0x0c]637// GFX90A: :[[@LINE-2]]:{{[0-9]+}}: error: instruction not supported on this GPU638 639v_smfmac_f32_16x16x64_fp8_bf8 a[0:3], v[2:3], a[4:7], v1640// GFX942: v_smfmac_f32_16x16x64_fp8_bf8 a[0:3], v[2:3], a[4:7], v1 ; encoding: [0x00,0x80,0xfa,0xd3,0x02,0x09,0x06,0x14]641// GFX90A: :[[@LINE-2]]:{{[0-9]+}}: error: instruction not supported on this GPU642 643v_smfmac_f32_16x16x64_fp8_fp8 v[0:3], a[2:3], v[4:7], v1 cbsz:3 abid:1644// GFX942: v_smfmac_f32_16x16x64_fp8_fp8 v[0:3], a[2:3], v[4:7], v1 cbsz:3 abid:1 ; encoding: [0x00,0x0b,0xfb,0xd3,0x02,0x09,0x06,0x0c]645// GFX90A: :[[@LINE-2]]:{{[0-9]+}}: error: instruction not supported on this GPU646 647v_smfmac_f32_16x16x64_fp8_fp8 a[0:3], v[2:3], a[4:7], v1648// GFX942: v_smfmac_f32_16x16x64_fp8_fp8 a[0:3], v[2:3], a[4:7], v1 ; encoding: [0x00,0x80,0xfb,0xd3,0x02,0x09,0x06,0x14]649// GFX90A: :[[@LINE-2]]:{{[0-9]+}}: error: instruction not supported on this GPU650 651v_smfmac_f32_32x32x32_bf8_bf8 v[0:15], a[2:3], v[4:7], v1 cbsz:3 abid:1652// GFX942: v_smfmac_f32_32x32x32_bf8_bf8 v[0:15], a[2:3], v[4:7], v1 cbsz:3 abid:1 ; encoding: [0x00,0x0b,0xfc,0xd3,0x02,0x09,0x06,0x0c]653// GFX90A: :[[@LINE-2]]:{{[0-9]+}}: error: instruction not supported on this GPU654 655v_smfmac_f32_32x32x32_bf8_bf8 a[0:15], v[2:3], a[4:7], v1656// GFX942: v_smfmac_f32_32x32x32_bf8_bf8 a[0:15], v[2:3], a[4:7], v1 ; encoding: [0x00,0x80,0xfc,0xd3,0x02,0x09,0x06,0x14]657// GFX90A: :[[@LINE-2]]:{{[0-9]+}}: error: instruction not supported on this GPU658 659v_smfmac_f32_32x32x32_bf8_fp8 v[0:15], a[2:3], v[4:7], v1 cbsz:3 abid:1660// GFX942: v_smfmac_f32_32x32x32_bf8_fp8 v[0:15], a[2:3], v[4:7], v1 cbsz:3 abid:1 ; encoding: [0x00,0x0b,0xfd,0xd3,0x02,0x09,0x06,0x0c]661// GFX90A: :[[@LINE-2]]:{{[0-9]+}}: error: instruction not supported on this GPU662 663v_smfmac_f32_32x32x32_bf8_fp8 a[0:15], v[2:3], a[4:7], v1664// GFX942: v_smfmac_f32_32x32x32_bf8_fp8 a[0:15], v[2:3], a[4:7], v1 ; encoding: [0x00,0x80,0xfd,0xd3,0x02,0x09,0x06,0x14]665// GFX90A: :[[@LINE-2]]:{{[0-9]+}}: error: instruction not supported on this GPU666 667v_smfmac_f32_32x32x32_fp8_bf8 v[0:15], a[2:3], v[4:7], v1 cbsz:3 abid:1668// GFX942: v_smfmac_f32_32x32x32_fp8_bf8 v[0:15], a[2:3], v[4:7], v1 cbsz:3 abid:1 ; encoding: [0x00,0x0b,0xfe,0xd3,0x02,0x09,0x06,0x0c]669// GFX90A: :[[@LINE-2]]:{{[0-9]+}}: error: instruction not supported on this GPU670 671v_smfmac_f32_32x32x32_fp8_bf8 a[0:15], v[2:3], a[4:7], v1672// GFX942: v_smfmac_f32_32x32x32_fp8_bf8 a[0:15], v[2:3], a[4:7], v1 ; encoding: [0x00,0x80,0xfe,0xd3,0x02,0x09,0x06,0x14]673// GFX90A: :[[@LINE-2]]:{{[0-9]+}}: error: instruction not supported on this GPU674 675v_smfmac_f32_32x32x32_fp8_fp8 v[0:15], a[2:3], v[4:7], v1 cbsz:3 abid:1676// GFX942: v_smfmac_f32_32x32x32_fp8_fp8 v[0:15], a[2:3], v[4:7], v1 cbsz:3 abid:1 ; encoding: [0x00,0x0b,0xff,0xd3,0x02,0x09,0x06,0x0c]677// GFX90A: :[[@LINE-2]]:{{[0-9]+}}: error: instruction not supported on this GPU678 679v_smfmac_f32_32x32x32_fp8_fp8 a[0:15], v[2:3], a[4:7], v1680// GFX942: v_smfmac_f32_32x32x32_fp8_fp8 a[0:15], v[2:3], a[4:7], v1 ; encoding: [0x00,0x80,0xff,0xd3,0x02,0x09,0x06,0x14]681// GFX90A: :[[@LINE-2]]:{{[0-9]+}}: error: instruction not supported on this GPU682 683//===----------------------------------------------------------------------===//684// SMFMAC aliases.685//===----------------------------------------------------------------------===//686 687v_smfmac_f32_16x16x32f16 v[10:13], a[2:3], v[4:7], v0 cbsz:3 abid:1688// GFX942: v_smfmac_f32_16x16x32_f16 v[10:13], a[2:3], v[4:7], v0 cbsz:3 abid:1 ; encoding: [0x0a,0x0b,0xe2,0xd3,0x02,0x09,0x02,0x0c]689// GFX90A: :[[@LINE-2]]:{{[0-9]+}}: error: instruction not supported on this GPU690 691v_smfmac_f32_32x32x16f16 v[10:25], a[2:3], v[4:7], v2 cbsz:3 abid:1692// GFX942: v_smfmac_f32_32x32x16_f16 v[10:25], a[2:3], v[4:7], v2 cbsz:3 abid:1 ; encoding: [0x0a,0x0b,0xe4,0xd3,0x02,0x09,0x0a,0x0c]693// GFX90A: :[[@LINE-2]]:{{[0-9]+}}: error: instruction not supported on this GPU694 695v_smfmac_f32_16x16x32bf16 v[10:13], a[2:3], v[4:7], v4 cbsz:3 abid:1696// GFX942: v_smfmac_f32_16x16x32_bf16 v[10:13], a[2:3], v[4:7], v4 cbsz:3 abid:1 ; encoding: [0x0a,0x0b,0xe6,0xd3,0x02,0x09,0x12,0x0c]697// GFX90A: :[[@LINE-2]]:{{[0-9]+}}: error: instruction not supported on this GPU698 699v_smfmac_f32_32x32x16bf16 v[10:25], a[2:3], v[4:7], v6 cbsz:3 abid:1700// GFX942: v_smfmac_f32_32x32x16_bf16 v[10:25], a[2:3], v[4:7], v6 cbsz:3 abid:1 ; encoding: [0x0a,0x0b,0xe8,0xd3,0x02,0x09,0x1a,0x0c]701// GFX90A: :[[@LINE-2]]:{{[0-9]+}}: error: instruction not supported on this GPU702 703v_smfmac_i32_16x16x64i8 v[10:13], a[2:3], v[4:7], v8 cbsz:3 abid:1704// GFX942: v_smfmac_i32_16x16x64_i8 v[10:13], a[2:3], v[4:7], v8 cbsz:3 abid:1 ; encoding: [0x0a,0x0b,0xea,0xd3,0x02,0x09,0x22,0x0c]705// GFX90A: :[[@LINE-2]]:{{[0-9]+}}: error: instruction not supported on this GPU706 707v_smfmac_i32_32x32x32i8 a[10:25], v[2:3], a[4:7], v11708// GFX942: v_smfmac_i32_32x32x32_i8 a[10:25], v[2:3], a[4:7], v11 ; encoding: [0x0a,0x80,0xec,0xd3,0x02,0x09,0x2e,0x14]709// GFX90A: :[[@LINE-2]]:{{[0-9]+}}: error: instruction not supported on this GPU710 711v_smfmac_f32_16x16x64bf8bf8 v[0:3], a[2:3], v[4:7], v1 cbsz:3 abid:1712// GFX942: v_smfmac_f32_16x16x64_bf8_bf8 v[0:3], a[2:3], v[4:7], v1 cbsz:3 abid:1 ; encoding: [0x00,0x0b,0xf8,0xd3,0x02,0x09,0x06,0x0c]713// GFX90A: :[[@LINE-2]]:{{[0-9]+}}: error: instruction not supported on this GPU714 715v_smfmac_f32_16x16x64bf8fp8 v[0:3], a[2:3], v[4:7], v1 cbsz:3 abid:1716// GFX942: v_smfmac_f32_16x16x64_bf8_fp8 v[0:3], a[2:3], v[4:7], v1 cbsz:3 abid:1 ; encoding: [0x00,0x0b,0xf9,0xd3,0x02,0x09,0x06,0x0c]717// GFX90A: :[[@LINE-2]]:{{[0-9]+}}: error: instruction not supported on this GPU718 719v_smfmac_f32_16x16x64fp8bf8 v[0:3], a[2:3], v[4:7], v1 cbsz:3 abid:1720// GFX942: v_smfmac_f32_16x16x64_fp8_bf8 v[0:3], a[2:3], v[4:7], v1 cbsz:3 abid:1 ; encoding: [0x00,0x0b,0xfa,0xd3,0x02,0x09,0x06,0x0c]721// GFX90A: :[[@LINE-2]]:{{[0-9]+}}: error: instruction not supported on this GPU722 723v_smfmac_f32_16x16x64fp8fp8 v[0:3], a[2:3], v[4:7], v1 cbsz:3 abid:1724// GFX942: v_smfmac_f32_16x16x64_fp8_fp8 v[0:3], a[2:3], v[4:7], v1 cbsz:3 abid:1 ; encoding: [0x00,0x0b,0xfb,0xd3,0x02,0x09,0x06,0x0c]725// GFX90A: :[[@LINE-2]]:{{[0-9]+}}: error: instruction not supported on this GPU726 727v_smfmac_f32_32x32x32bf8bf8 v[0:15], a[2:3], v[4:7], v1 cbsz:3 abid:1728// GFX942: v_smfmac_f32_32x32x32_bf8_bf8 v[0:15], a[2:3], v[4:7], v1 cbsz:3 abid:1 ; encoding: [0x00,0x0b,0xfc,0xd3,0x02,0x09,0x06,0x0c]729// GFX90A: :[[@LINE-2]]:{{[0-9]+}}: error: instruction not supported on this GPU730 731v_smfmac_f32_32x32x32bf8fp8 v[0:15], a[2:3], v[4:7], v1 cbsz:3 abid:1732// GFX942: v_smfmac_f32_32x32x32_bf8_fp8 v[0:15], a[2:3], v[4:7], v1 cbsz:3 abid:1 ; encoding: [0x00,0x0b,0xfd,0xd3,0x02,0x09,0x06,0x0c]733// GFX90A: :[[@LINE-2]]:{{[0-9]+}}: error: instruction not supported on this GPU734 735v_smfmac_f32_32x32x32fp8bf8 v[0:15], a[2:3], v[4:7], v1 cbsz:3 abid:1736// GFX942: v_smfmac_f32_32x32x32_fp8_bf8 v[0:15], a[2:3], v[4:7], v1 cbsz:3 abid:1 ; encoding: [0x00,0x0b,0xfe,0xd3,0x02,0x09,0x06,0x0c]737// GFX90A: :[[@LINE-2]]:{{[0-9]+}}: error: instruction not supported on this GPU738 739v_smfmac_f32_32x32x32fp8fp8 v[0:15], a[2:3], v[4:7], v1 cbsz:3 abid:1740// GFX942: v_smfmac_f32_32x32x32_fp8_fp8 v[0:15], a[2:3], v[4:7], v1 cbsz:3 abid:1 ; encoding: [0x00,0x0b,0xff,0xd3,0x02,0x09,0x06,0x0c]741// GFX90A: :[[@LINE-2]]:{{[0-9]+}}: error: instruction not supported on this GPU742