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