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