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