531 lines · plain
1// RUN: not llvm-mc -triple=amdgcn -mcpu=gfx908 -show-encoding %s | FileCheck -check-prefix=GFX908 %s2// RUN: not llvm-mc -triple=amdgcn -mcpu=gfx908 %s 2>&1 | FileCheck -check-prefix=NOGFX908 --implicit-check-not=error: %s3 4v_accvgpr_read_b32 v2, a05// GFX908: v_accvgpr_read_b32 v2, a0 ; encoding: [0x02,0x40,0xd8,0xd3,0x00,0x01,0x00,0x18]6 7v_accvgpr_read_b32 v2, a18// GFX908: v_accvgpr_read_b32 v2, a1 ; encoding: [0x02,0x40,0xd8,0xd3,0x01,0x01,0x00,0x18]9 10v_accvgpr_read_b32 v2, a25511// GFX908: v_accvgpr_read_b32 v2, a255 ; encoding: [0x02,0x40,0xd8,0xd3,0xff,0x01,0x00,0x18]12 13v_accvgpr_read v2, a1014// GFX908: v_accvgpr_read_b32 v2, a10 ; encoding: [0x02,0x40,0xd8,0xd3,0x0a,0x01,0x00,0x18]15 16v_accvgpr_write_b32 a2, -2.017// GFX908: v_accvgpr_write_b32 a2, -2.0 ; encoding: [0x02,0x40,0xd9,0xd3,0xf5,0x00,0x00,0x18]18 19v_accvgpr_write_b32 a2, -220// GFX908: v_accvgpr_write_b32 a2, -2 ; encoding: [0x02,0x40,0xd9,0xd3,0xc2,0x00,0x00,0x18]21 22v_accvgpr_write_b32 a2, v123// GFX908: v_accvgpr_write_b32 a2, v1 ; encoding: [0x02,0x40,0xd9,0xd3,0x01,0x01,0x00,0x18]24 25v_accvgpr_write a2, v25526// GFX908: v_accvgpr_write_b32 a2, v255 ; encoding: [0x02,0x40,0xd9,0xd3,0xff,0x01,0x00,0x18]27 28v_accvgpr_write a2, 10029// NOGFX908: :[[@LINE-1]]:{{[0-9]+}}: error: invalid operand for instruction30 31v_accvgpr_write a2, execz32// NOGFX908: :[[@LINE-1]]:{{[0-9]+}}: error: source operand must be either a VGPR or an inline constant33 34v_accvgpr_write a2, vccz35// NOGFX908: :[[@LINE-1]]:{{[0-9]+}}: error: source operand must be either a VGPR or an inline constant36 37v_accvgpr_write a2, scc38// NOGFX908: :[[@LINE-1]]:{{[0-9]+}}: error: source operand must be either a VGPR or an inline constant39 40v_accvgpr_write a2, shared_base41// NOGFX908: :[[@LINE-1]]:{{[0-9]+}}: error: source operand must be either a VGPR or an inline constant42 43v_accvgpr_write a2, pops_exiting_wave_id44// NOGFX908: :[[@LINE-1]]:{{[0-9]+}}: error: source operand must be either a VGPR or an inline constant45 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:753// 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:759// 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:765// 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:771// 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:777// 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:783// 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:789// 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:795// 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:7104// 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:7110// 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:7116// 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:7122// 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:7128// 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:7134// 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:7140// 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:7146// 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:7152// 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:7158// 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:7164// 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:7170// 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:7176// 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:7182// 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:7188// 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:7194// 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:7200// 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:7206// 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:7212// 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:7218// 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:7224// 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:7230// 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:7236// 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:7242// 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:7248// 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:7254// 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:7260// 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:7266// 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:7272// 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:7278// 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:7284// 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:7290// 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:7296// 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:7302// 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:7308// 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:7314// 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:7320// 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:7326// 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:7332// 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:7338// 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:7344// 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:7350// 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:7356// 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:7362// 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:7368// 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:7374// 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:7380// 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:7386// 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:7392// 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:7398// 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:7404// 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:7410// 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:7416// 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:7422// 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:7428// 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:7434// 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:7440// 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:7446// 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:7452// 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:7458// 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:7464// 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:7470// 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:7476// 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:7482// 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:7488// 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:7494// 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:7500// 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:7506// 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:7512// 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:7518// 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:7524// 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:7530// 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