xref: /llvm-project/llvm/test/MC/AMDGPU/mai-gfx950.s (revision 7d544c64e3b6ea014c59e230dcf65ac4f9d60f2b) 
1// RUN: llvm-mc -triple=amdgcn -mcpu=gfx950 -show-encoding %s | FileCheck -check-prefix=GFX950 %s
2// RUN: not llvm-mc -triple=amdgcn -mcpu=gfx940 %s 2>&1 | FileCheck -check-prefix=ERR %s
3// RUN: not llvm-mc -triple=amdgcn -mcpu=gfx90a %s 2>&1 | FileCheck -check-prefix=ERR %s
4// RUN: not llvm-mc -triple=amdgcn -mcpu=gfx1010 %s 2>&1 | FileCheck -check-prefix=ERR %s
5
6// cbsz = SrcA
7// blgp = SrcB
8
9// 0 = fp8. 1 = bf8
10// 2 = fp6, 3 = bf6
11// 4 = fp4
12
13//===----------------------------------------------------------------------===//
14// MFMA opcodes.
15//===----------------------------------------------------------------------===//
16
17
18//===----------------------------------------------------------------------===//
19// v_mfma_f32_16x16x32_f16
20//===----------------------------------------------------------------------===//
21
22// GFX950: v_mfma_f32_16x16x32_f16 v[0:3], v[0:3], v[0:3], v[0:3] ; encoding: [0x00,0x00,0xd4,0xd3,0x00,0x01,0x02,0x04]
23// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU
24v_mfma_f32_16x16x32_f16 v[0:3], v[0:3], v[0:3], v[0:3]
25
26// GFX950: v_mfma_f32_16x16x32_f16 v[0:3], v[0:3], v[0:3], v[0:3] ; encoding: [0x00,0x00,0xd4,0xd3,0x00,0x01,0x02,0x04]
27// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU
28v_mfma_f32_16x16x32f16 v[0:3], v[0:3], v[0:3], v[0:3]
29
30// GFX950: v_mfma_f32_16x16x32_f16 a[0:3], a[0:3], a[0:3], a[0:3] ; encoding: [0x00,0x80,0xd4,0xd3,0x00,0x01,0x02,0x1c]
31// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU
32v_mfma_f32_16x16x32_f16 a[0:3], a[0:3], a[0:3], a[0:3]
33
34// GFX950: v_mfma_f32_16x16x32_f16 a[0:3], a[0:3], a[0:3], a[0:3] ; encoding: [0x00,0x80,0xd4,0xd3,0x00,0x01,0x02,0x1c]
35// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU
36v_mfma_f32_16x16x32f16 a[0:3], a[0:3], a[0:3], a[0:3]
37
38// GFX950: v_mfma_f32_16x16x32_f16 v[0:3], a[0:3], v[0:3], 1.0 ; encoding: [0x00,0x00,0xd4,0xd3,0x00,0x01,0xca,0x0b]
39// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU
40v_mfma_f32_16x16x32_f16 v[0:3], a[0:3], v[0:3], 1.0
41
42// GFX950: v_mfma_f32_16x16x32_f16 a[0:3], v[0:3], a[0:3], 1.0 ; encoding: [0x00,0x80,0xd4,0xd3,0x00,0x01,0xca,0x13]
43// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU
44v_mfma_f32_16x16x32_f16 a[0:3], v[0:3], a[0:3], 1.0
45
46// GFX950: v_mfma_f32_16x16x32_f16 v[0:3], v[0:3], v[0:3], v[0:3] blgp:5 ; encoding: [0x00,0x00,0xd4,0xd3,0x00,0x01,0x02,0xa4]
47// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU
48v_mfma_f32_16x16x32_f16 v[0:3], v[0:3], v[0:3], v[0:3] blgp:5
49
50// GFX950: v_mfma_f32_16x16x32_f16 a[0:3], a[0:3], a[0:3], a[0:3] blgp:1 ; encoding: [0x00,0x80,0xd4,0xd3,0x00,0x01,0x02,0x3c]
51// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU
52v_mfma_f32_16x16x32_f16 a[0:3], a[0:3], a[0:3], a[0:3] blgp:1
53
54// GFX950: v_mfma_f32_16x16x32_f16 a[0:3], a[0:3], a[0:3], a[0:3] cbsz:3 ; encoding: [0x00,0x83,0xd4,0xd3,0x00,0x01,0x02,0x1c]
55// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU
56v_mfma_f32_16x16x32_f16 a[0:3], a[0:3], a[0:3], a[0:3] cbsz:3
57
58// GFX950:  v_mfma_f32_16x16x32_f16 a[0:3], a[0:3], a[0:3], a[0:3] abid:1 ; encoding: [0x00,0x88,0xd4,0xd3,0x00,0x01,0x02,0x1c]
59// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU
60v_mfma_f32_16x16x32_f16 a[0:3], a[0:3], a[0:3], a[0:3] abid:1
61
62// GFX950:  v_mfma_f32_16x16x32_f16 a[0:3], a[0:3], a[0:3], a[0:3] cbsz:3 abid:1 ; encoding: [0x00,0x8b,0xd4,0xd3,0x00,0x01,0x02,0x1c]
63// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU
64v_mfma_f32_16x16x32_f16 a[0:3], a[0:3], a[0:3], a[0:3] cbsz:3 abid:1
65
66// GFX950: v_mfma_f32_16x16x32_f16 a[0:3], v[0:3], v[0:3], a[4:7] ; encoding: [0x00,0x80,0xd4,0xd3,0x00,0x01,0x12,0x04]
67// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU
68v_mfma_f32_16x16x32_f16 a[0:3], v[0:3], v[0:3], a[4:7]
69
70// GFX950: v_mfma_f32_16x16x32_f16 v[0:3], a[0:3], a[0:3], v[4:7] ; encoding: [0x00,0x00,0xd4,0xd3,0x00,0x01,0x12,0x1c]
71// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU
72v_mfma_f32_16x16x32_f16 v[0:3], a[0:3], a[0:3], v[4:7]
73
74//===----------------------------------------------------------------------===//
75// v_mfma_f32_32x32x16_f16
76//===----------------------------------------------------------------------===//
77
78// GFX950: v_mfma_f32_32x32x16_f16 v[0:15], v[0:3], v[0:3], v[0:15] ; encoding: [0x00,0x00,0xd5,0xd3,0x00,0x01,0x02,0x04]
79// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU
80v_mfma_f32_32x32x16_f16 v[0:15], v[0:3], v[0:3], v[0:15]
81
82// GFX950: v_mfma_f32_32x32x16_f16 a[0:15], a[0:3], a[0:3], a[0:15] ; encoding: [0x00,0x80,0xd5,0xd3,0x00,0x01,0x02,0x1c]
83// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU
84v_mfma_f32_32x32x16_f16 a[0:15], a[0:3], a[0:3], a[0:15]
85
86// GFX950:  v_mfma_f32_32x32x16_f16 v[0:15], v[0:3], v[0:3], v[0:15] ; encoding: [0x00,0x00,0xd5,0xd3,0x00,0x01,0x02,0x04]
87// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU
88v_mfma_f32_32x32x16f16 v[0:15], v[0:3], v[0:3], v[0:15]
89
90// GFX950: v_mfma_f32_32x32x16_f16 a[0:15], a[0:3], a[0:3], a[0:15] ; encoding: [0x00,0x80,0xd5,0xd3,0x00,0x01,0x02,0x1c]
91// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU
92v_mfma_f32_32x32x16f16 a[0:15], a[0:3], a[0:3], a[0:15]
93
94// GFX950: v_mfma_f32_32x32x16_f16 v[0:15], v[0:3], v[0:3], 1.0 ; encoding: [0x00,0x00,0xd5,0xd3,0x00,0x01,0xca,0x03]
95// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU
96v_mfma_f32_32x32x16_f16 v[0:15], v[0:3], v[0:3], 1.0
97
98// GFX950: v_mfma_f32_32x32x16_f16 a[0:15], a[0:3], a[0:3], 1.0 ; encoding: [0x00,0x80,0xd5,0xd3,0x00,0x01,0xca,0x1b]
99// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU
100v_mfma_f32_32x32x16_f16 a[0:15], a[0:3], a[0:3], 1.0
101
102// GFX950:  v_mfma_f32_32x32x16_f16 v[0:15], v[0:3], v[0:3], v[0:15] blgp:5 ; encoding: [0x00,0x00,0xd5,0xd3,0x00,0x01,0x02,0xa4]
103// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU
104v_mfma_f32_32x32x16_f16 v[0:15], v[0:3], v[0:3], v[0:15] blgp:5
105
106// GFX950: v_mfma_f32_32x32x16_f16 a[0:15], a[0:3], a[0:3], a[0:15] blgp:2 ; encoding: [0x00,0x80,0xd5,0xd3,0x00,0x01,0x02,0x5c]
107// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU
108v_mfma_f32_32x32x16_f16 a[0:15], a[0:3], a[0:3], a[0:15] blgp:2
109
110// GFX950: v_mfma_f32_32x32x16_f16 v[0:15], v[0:3], v[0:3], v[0:15] cbsz:3 ; encoding: [0x00,0x03,0xd5,0xd3,0x00,0x01,0x02,0x04]
111// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU
112v_mfma_f32_32x32x16_f16 v[0:15], v[0:3], v[0:3], v[0:15] cbsz:3
113
114// GFX950: v_mfma_f32_32x32x16_f16 v[0:15], v[0:3], v[0:3], v[0:15] abid:1 ; encoding: [0x00,0x08,0xd5,0xd3,0x00,0x01,0x02,0x04]
115// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU
116v_mfma_f32_32x32x16_f16 v[0:15], v[0:3], v[0:3], v[0:15] abid:1
117
118// GFX950: v_mfma_f32_32x32x16_f16 a[0:15], a[0:3], a[0:3], a[0:15] cbsz:3 abid:1 ; encoding: [0x00,0x8b,0xd5,0xd3,0x00,0x01,0x02,0x1c]
119// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU
120v_mfma_f32_32x32x16_f16 a[0:15], a[0:3], a[0:3], a[0:15] cbsz:3 abid:1
121
122//===----------------------------------------------------------------------===//
123// v_mfma_f32_32x32x16_bf16
124//===----------------------------------------------------------------------===//
125
126// GFX950: v_mfma_f32_32x32x16_bf16 v[0:15], v[0:3], v[0:3], v[0:15] ; encoding: [0x00,0x00,0xb7,0xd3,0x00,0x01,0x02,0x04]
127// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU
128v_mfma_f32_32x32x16_bf16 v[0:15], v[0:3], v[0:3], v[0:15]
129
130// GFX950: v_mfma_f32_32x32x16_bf16 a[0:15], a[0:3], a[0:3], a[0:15] ; encoding: [0x00,0x80,0xb7,0xd3,0x00,0x01,0x02,0x1c]
131// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU
132v_mfma_f32_32x32x16_bf16 a[0:15], a[0:3], a[0:3], a[0:15]
133
134// GFX950:  v_mfma_f32_32x32x16_bf16 v[0:15], v[0:3], v[0:3], v[0:15] ; encoding: [0x00,0x00,0xb7,0xd3,0x00,0x01,0x02,0x04]
135// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU
136v_mfma_f32_32x32x16bf16 v[0:15], v[0:3], v[0:3], v[0:15]
137
138// GFX950: v_mfma_f32_32x32x16_bf16 a[0:15], a[0:3], a[0:3], a[0:15] ; encoding: [0x00,0x80,0xb7,0xd3,0x00,0x01,0x02,0x1c]
139// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU
140v_mfma_f32_32x32x16bf16 a[0:15], a[0:3], a[0:3], a[0:15]
141
142// GFX950: v_mfma_f32_32x32x16_bf16 v[0:15], v[0:3], v[0:3], 1.0 ; encoding: [0x00,0x00,0xb7,0xd3,0x00,0x01,0xca,0x03]
143// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU
144v_mfma_f32_32x32x16_bf16 v[0:15], v[0:3], v[0:3], 1.0
145
146// GFX950: v_mfma_f32_32x32x16_bf16 a[0:15], a[0:3], a[0:3], 1.0 ; encoding: [0x00,0x80,0xb7,0xd3,0x00,0x01,0xca,0x1b]
147// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU
148v_mfma_f32_32x32x16_bf16 a[0:15], a[0:3], a[0:3], 1.0
149
150// GFX950:  v_mfma_f32_32x32x16_bf16 v[0:15], v[0:3], v[0:3], v[0:15] blgp:5 ; encoding: [0x00,0x00,0xb7,0xd3,0x00,0x01,0x02,0xa4]
151// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU
152v_mfma_f32_32x32x16_bf16 v[0:15], v[0:3], v[0:3], v[0:15] blgp:5
153
154// GFX950: v_mfma_f32_32x32x16_bf16 a[0:15], a[0:3], a[0:3], a[0:15] blgp:2 ; encoding: [0x00,0x80,0xb7,0xd3,0x00,0x01,0x02,0x5c]
155// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU
156v_mfma_f32_32x32x16_bf16 a[0:15], a[0:3], a[0:3], a[0:15] blgp:2
157
158// GFX950: v_mfma_f32_32x32x16_bf16 v[0:15], v[0:3], v[0:3], v[0:15] cbsz:3 ; encoding: [0x00,0x03,0xb7,0xd3,0x00,0x01,0x02,0x04]
159// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU
160v_mfma_f32_32x32x16_bf16 v[0:15], v[0:3], v[0:3], v[0:15] cbsz:3
161
162// GFX950: v_mfma_f32_32x32x16_bf16 v[0:15], v[0:3], v[0:3], v[0:15] abid:1 ; encoding: [0x00,0x08,0xb7,0xd3,0x00,0x01,0x02,0x04]
163// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU
164v_mfma_f32_32x32x16_bf16 v[0:15], v[0:3], v[0:3], v[0:15] abid:1
165
166// GFX950: v_mfma_f32_32x32x16_bf16 a[0:15], a[0:3], a[0:3], a[0:15] cbsz:3 abid:1 ; encoding: [0x00,0x8b,0xb7,0xd3,0x00,0x01,0x02,0x1c]
167// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU
168v_mfma_f32_32x32x16_bf16 a[0:15], a[0:3], a[0:3], a[0:15] cbsz:3 abid:1
169
170//===----------------------------------------------------------------------===//
171// v_mfma_ld_scale_b32
172//===----------------------------------------------------------------------===//
173
174// GFX950: v_mfma_ld_scale_b32 v0, 64             ; encoding: [0x00,0x40,0xac,0xd3,0x00,0x81,0x01,0x18]
175// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU
176v_mfma_ld_scale_b32 v0, 64
177
178// GFX950: v_mfma_ld_scale_b32 64, v0             ; encoding: [0x00,0x40,0xac,0xd3,0xc0,0x00,0x02,0x18]
179// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU
180v_mfma_ld_scale_b32 64, v0
181
182// GFX950: v_mfma_ld_scale_b32 64, 64             ; encoding: [0x00,0x40,0xac,0xd3,0xc0,0x80,0x01,0x18]
183// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU
184v_mfma_ld_scale_b32 64, 64
185
186// GFX950: v_mfma_ld_scale_b32 s0, s0             ; encoding: [0x00,0x40,0xac,0xd3,0x00,0x00,0x00,0x18]
187// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU
188v_mfma_ld_scale_b32 s0, s0
189
190// GFX950: v_mfma_ld_scale_b32 s0, v0             ; encoding: [0x00,0x40,0xac,0xd3,0x00,0x00,0x02,0x18]
191// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU
192v_mfma_ld_scale_b32 s0, v0
193
194// GFX950: v_mfma_ld_scale_b32 v0, s0             ; encoding: [0x00,0x40,0xac,0xd3,0x00,0x01,0x00,0x18]
195// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU
196v_mfma_ld_scale_b32 v0, s0
197
198// GFX950: v_mfma_ld_scale_b32 vcc_lo, vcc_lo     ; encoding: [0x00,0x40,0xac,0xd3,0x6a,0xd4,0x00,0x18]
199// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU
200v_mfma_ld_scale_b32 vcc_lo, vcc_lo
201
202// GFX950: v_mfma_ld_scale_b32 m0, m0             ; encoding: [0x00,0x40,0xac,0xd3,0x7c,0xf8,0x00,0x18]
203// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU
204v_mfma_ld_scale_b32 m0, m0
205
206// GFX950: v_mfma_ld_scale_b32 src_vccz, src_vccz ; encoding: [0x00,0x40,0xac,0xd3,0xfb,0xf6,0x01,0x18]
207// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU
208v_mfma_ld_scale_b32 vccz, vccz
209
210// GFX950: v_mfma_ld_scale_b32 src_execz, src_execz ; encoding: [0x00,0x40,0xac,0xd3,0xfc,0xf8,0x01,0x18]
211// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU
212v_mfma_ld_scale_b32 execz, execz
213
214// GFX950:  v_mfma_ld_scale_b32 v0, v0 ; encoding: [0x00,0x40,0xac,0xd3,0x00,0x01,0x02,0x18]
215// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU
216v_mfma_ld_scale_b32 v0, v0
217
218// GFX950: v_mfma_ld_scale_b32 v1, v1 ; encoding: [0x00,0x40,0xac,0xd3,0x01,0x03,0x02,0x18]
219// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU
220v_mfma_ld_scale_b32 v1, v1
221
222// GFX950: v_mfma_ld_scale_b32 0, 0 ; encoding: [0x00,0x40,0xac,0xd3,0x80,0x00,0x01,0x18]
223// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU
224v_mfma_ld_scale_b32 0, 0
225
226// GFX950: v_mfma_ld_scale_b32 1, 0               ; encoding: [0x00,0x40,0xac,0xd3,0x81,0x00,0x01,0x18]
227// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU
228v_mfma_ld_scale_b32 1, 0
229
230// GFX950: v_mfma_ld_scale_b32 v1, v1 op_sel:[1,0] ; encoding: [0x00,0x48,0xac,0xd3,0x01,0x03,0x02,0x18]
231// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU
232v_mfma_ld_scale_b32 v1, v1 op_sel:[1, 0]
233
234// GFX950: v_mfma_ld_scale_b32 v1, v1 op_sel:[0,1] ; encoding: [0x00,0x50,0xac,0xd3,0x01,0x03,0x02,0x18]
235// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU
236v_mfma_ld_scale_b32 v1, v1 op_sel:[0, 1]
237
238// GFX950: v_mfma_ld_scale_b32 v1, v1 op_sel:[1,1] ; encoding: [0x00,0x58,0xac,0xd3,0x01,0x03,0x02,0x18]
239// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU
240v_mfma_ld_scale_b32 v1, v1 op_sel:[1, 1]
241
242// GFX950: v_mfma_ld_scale_b32 v1, v1 op_sel_hi:[1,0] ; encoding: [0x00,0x40,0xac,0xd3,0x01,0x03,0x02,0x08]
243// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU
244v_mfma_ld_scale_b32 v1, v1 op_sel_hi:[1, 0]
245
246// GFX950: v_mfma_ld_scale_b32 v1, v1 op_sel_hi:[0,1] ; encoding: [0x00,0x40,0xac,0xd3,0x01,0x03,0x02,0x10]
247// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU
248v_mfma_ld_scale_b32 v1, v1 op_sel_hi:[0, 1]
249
250// GFX950: v_mfma_ld_scale_b32 v1, v1             ; encoding: [0x00,0x40,0xac,0xd3,0x01,0x03,0x02,0x18]
251// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU
252v_mfma_ld_scale_b32 v1, v1 op_sel_hi:[1, 1]
253
254// GFX950: v_mfma_ld_scale_b32 v1, v1 op_sel_hi:[0,0] ; encoding: [0x00,0x40,0xac,0xd3,0x01,0x03,0x02,0x00]
255// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU
256v_mfma_ld_scale_b32 v1, v1 op_sel:[0,0] op_sel_hi:[0,0]
257
258// GFX950: v_mfma_ld_scale_b32 v1, v1 op_sel:[1,0] op_sel_hi:[1,0] ; encoding: [0x00,0x48,0xac,0xd3,0x01,0x03,0x02,0x08]
259// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU
260v_mfma_ld_scale_b32 v1, v1 op_sel:[1,0] op_sel_hi:[1,0]
261
262// GFX950: v_mfma_ld_scale_b32 v1, v1 op_sel:[0,1] op_sel_hi:[0,1] ; encoding: [0x00,0x50,0xac,0xd3,0x01,0x03,0x02,0x10]
263// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU
264v_mfma_ld_scale_b32 v1, v1 op_sel:[0,1] op_sel_hi:[0,1]
265
266// GFX950: v_mfma_ld_scale_b32 v1, v1 op_sel:[0,1] ; encoding: [0x00,0x50,0xac,0xd3,0x01,0x03,0x02,0x18]
267// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU
268v_mfma_ld_scale_b32 v1, v1 op_sel:[0,1] op_sel_hi:[1,1]
269
270// GFX950: v_mfma_ld_scale_b32 v1, v1 op_sel:[0,1] ; encoding: [0x00,0x50,0xac,0xd3,0x01,0x03,0x02,0x18]
271// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU
272v_mfma_ld_scale_b32 v1, v1 op_sel:[0,1] op_sel_hi:[1,1]
273
274// GFX950: v_mfma_ld_scale_b32 v1, v1 op_sel:[1,1] ; encoding: [0x00,0x58,0xac,0xd3,0x01,0x03,0x02,0x18]
275// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU
276v_mfma_ld_scale_b32 v1, v1 op_sel:[1,1] op_sel_hi:[1,1]
277
278// GFX950: v_mfma_ld_scale_b32 v1, v1 op_sel:[1,0] op_sel_hi:[0,1] ; encoding: [0x00,0x48,0xac,0xd3,0x01,0x03,0x02,0x10]
279// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU
280v_mfma_ld_scale_b32 v1, v1 op_sel:[1,0] op_sel_hi:[0,1]
281
282// GFX950: v_mfma_ld_scale_b32 v1, v1 op_sel:[0,1] op_sel_hi:[1,0] ; encoding: [0x00,0x50,0xac,0xd3,0x01,0x03,0x02,0x08]
283// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU
284v_mfma_ld_scale_b32 v1, v1 op_sel:[0,1] op_sel_hi:[1,0]
285
286
287//===----------------------------------------------------------------------===//
288// v_mfma_f32_16x16x128_f8f6f4
289//===----------------------------------------------------------------------===//
290
291// GFX950: v_mfma_f32_16x16x128_f8f6f4 v[0:3], v[4:11], v[4:11], v[0:3] ; encoding: [0x00,0x00,0xad,0xd3,0x04,0x09,0x02,0x04]
292// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU
293v_mfma_f32_16x16x128_f8f6f4 v[0:3], v[4:11], v[4:11], v[0:3]
294
295// GFX950: v_mfma_f32_16x16x128_f8f6f4 v[0:3], v[4:11], v[4:11], v[0:3] blgp:1 ; encoding: [0x00,0x00,0xad,0xd3,0x04,0x09,0x02,0x24]
296// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU
297v_mfma_f32_16x16x128_f8f6f4 v[0:3], v[4:11], v[4:11], v[0:3] blgp:1
298
299// GFX950: v_mfma_f32_16x16x128_f8f6f4 v[0:3], v[4:9], v[4:11], v[0:3] cbsz:3 ; encoding: [0x00,0x03,0xad,0xd3,0x04,0x09,0x02,0x04]
300// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU
301v_mfma_f32_16x16x128_f8f6f4 v[0:3], v[4:9], v[4:11], v[0:3] cbsz:3
302
303// GFX950: v_mfma_f32_16x16x128_f8f6f4 v[0:3], v[4:11], v[4:11], v[0:3] ; encoding: [0x00,0x00,0xad,0xd3,0x04,0x09,0x02,0x04]
304// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU
305v_mfma_f32_16x16x128_f8f6f4 v[0:3], v[4:11], v[4:11], v[0:3]
306
307// GFX950: v_mfma_f32_16x16x128_f8f6f4 v[0:3], v[4:9], v[4:11], v[0:3] cbsz:3 blgp:1 ; encoding: [0x00,0x03,0xad,0xd3,0x04,0x09,0x02,0x24]
308// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU
309v_mfma_f32_16x16x128_f8f6f4 v[0:3], v[4:9], v[4:11], v[0:3] cbsz:3 blgp:1
310
311// GFX950: v_mfma_f32_16x16x128_f8f6f4 a[0:3], a[4:11], a[4:11], a[0:3] ; encoding: [0x00,0x80,0xad,0xd3,0x04,0x09,0x02,0x1c]
312// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU
313v_mfma_f32_16x16x128_f8f6f4 a[0:3], a[4:11], a[4:11], a[0:3]
314
315// GFX950: v_mfma_f32_16x16x128_f8f6f4 a[0:3], a[4:11], a[4:11], a[0:3] blgp:1 ; encoding: [0x00,0x80,0xad,0xd3,0x04,0x09,0x02,0x3c]
316// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU
317v_mfma_f32_16x16x128_f8f6f4 a[0:3], a[4:11], a[4:11], a[0:3] blgp:1
318
319// GFX950: v_mfma_f32_16x16x128_f8f6f4 a[0:3], a[4:9], a[4:11], a[0:3] cbsz:3 ; encoding: [0x00,0x83,0xad,0xd3,0x04,0x09,0x02,0x1c]
320// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU
321v_mfma_f32_16x16x128_f8f6f4 a[0:3], a[4:9], a[4:11], a[0:3] cbsz:3
322
323// GFX950: v_mfma_f32_16x16x128_f8f6f4 a[0:3], a[4:11], a[4:11], a[0:3] ; encoding: [0x00,0x80,0xad,0xd3,0x04,0x09,0x02,0x1c]
324// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU
325v_mfma_f32_16x16x128_f8f6f4 a[0:3], a[4:11], a[4:11], a[0:3]
326
327// GFX950: v_mfma_f32_16x16x128_f8f6f4 a[0:3], a[4:11], a[4:9], a[0:3] cbsz:1 blgp:3 ; encoding: [0x00,0x81,0xad,0xd3,0x04,0x09,0x02,0x7c]
328// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU
329v_mfma_f32_16x16x128_f8f6f4 a[0:3], a[4:11], a[4:9], a[0:3] cbsz:1 blgp:3
330
331//===----------------------------------------------------------------------===//
332// v_mfma_f32_32x32x64_f8f6f4
333//===----------------------------------------------------------------------===//
334
335// GFX950: v_mfma_f32_32x32x64_f8f6f4 v[0:15], v[4:11], v[4:11], v[0:15] ; encoding: [0x00,0x00,0xae,0xd3,0x04,0x09,0x02,0x04]
336// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU
337v_mfma_f32_32x32x64_f8f6f4 v[0:15], v[4:11], v[4:11], v[0:15]
338
339// GFX950: v_mfma_f32_32x32x64_f8f6f4 v[0:15], v[4:11], v[4:11], v[0:15] blgp:1 ; encoding: [0x00,0x00,0xae,0xd3,0x04,0x09,0x02,0x24]
340// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU
341v_mfma_f32_32x32x64_f8f6f4 v[0:15], v[4:11], v[4:11], v[0:15] blgp:1
342
343// GFX950: v_mfma_f32_32x32x64_f8f6f4 v[0:15], v[4:9], v[4:11], v[0:15] cbsz:3 ; encoding: [0x00,0x03,0xae,0xd3,0x04,0x09,0x02,0x04]
344// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU
345v_mfma_f32_32x32x64_f8f6f4 v[0:15], v[4:9], v[4:11], v[0:15] cbsz:3
346
347// GFX950: v_mfma_f32_32x32x64_f8f6f4 v[0:15], v[4:11], v[4:11], v[0:15] ; encoding: [0x00,0x00,0xae,0xd3,0x04,0x09,0x02,0x04]
348// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU
349v_mfma_f32_32x32x64_f8f6f4 v[0:15], v[4:11], v[4:11], v[0:15]
350
351// GFX950: v_mfma_f32_32x32x64_f8f6f4 v[0:15], v[4:9], v[4:11], v[0:15] cbsz:3 blgp:1 ; encoding: [0x00,0x03,0xae,0xd3,0x04,0x09,0x02,0x24]
352// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU
353v_mfma_f32_32x32x64_f8f6f4 v[0:15], v[4:9], v[4:11], v[0:15] cbsz:3 blgp:1
354
355// GFX950: v_mfma_f32_32x32x64_f8f6f4 a[0:15], a[4:11], a[4:11], a[0:15] ; encoding: [0x00,0x80,0xae,0xd3,0x04,0x09,0x02,0x1c]
356// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU
357v_mfma_f32_32x32x64_f8f6f4 a[0:15], a[4:11], a[4:11], a[0:15]
358
359// GFX950: v_mfma_f32_32x32x64_f8f6f4 a[0:15], a[4:11], a[4:11], a[0:15] blgp:1 ; encoding: [0x00,0x80,0xae,0xd3,0x04,0x09,0x02,0x3c]
360// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU
361v_mfma_f32_32x32x64_f8f6f4 a[0:15], a[4:11], a[4:11], a[0:15] blgp:1
362
363// GFX950: v_mfma_f32_32x32x64_f8f6f4 a[0:15], a[4:9], a[4:11], a[0:15] cbsz:3 ; encoding: [0x00,0x83,0xae,0xd3,0x04,0x09,0x02,0x1c]
364// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU
365v_mfma_f32_32x32x64_f8f6f4 a[0:15], a[4:9], a[4:11], a[0:15] cbsz:3
366
367// GFX950: v_mfma_f32_32x32x64_f8f6f4 a[0:15], a[4:11], a[4:11], a[0:15] ; encoding: [0x00,0x80,0xae,0xd3,0x04,0x09,0x02,0x1c]
368// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU
369v_mfma_f32_32x32x64_f8f6f4 a[0:15], a[4:11], a[4:11], a[0:15]
370
371// GFX950: v_mfma_f32_32x32x64_f8f6f4 a[0:15], a[4:11], a[4:9], a[0:15] cbsz:1 blgp:3 ; encoding: [0x00,0x81,0xae,0xd3,0x04,0x09,0x02,0x7c]
372// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU
373v_mfma_f32_32x32x64_f8f6f4 a[0:15], a[4:11], a[4:9], a[0:15] cbsz:1 blgp:3
374
375//===----------------------------------------------------------------------===//
376// v_mfma_scale_f32_16x16x128_f8f6f4
377//===----------------------------------------------------------------------===//
378// FIXME: Test op_sel, neg, clamp
379
380// GFX950: v_mfma_scale_f32_16x16x128_f8f6f4 v[0:3], v[4:11], v[12:19], v[20:23], v24, v25 op_sel_hi:[0,0,0] ; encoding: [0x00,0x00,0xac,0xd3,0x18,0x33,0x02,0x00,0x00,0x08,0xad,0xd3,0x04,0x19,0x52,0x04]
381// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU
382v_mfma_scale_f32_16x16x128_f8f6f4 v[0:3], v[4:11], v[12:19], v[20:23], v24, v25
383
384// GFX950: v_mfma_scale_f32_16x16x128_f8f6f4 a[0:3], a[4:11], a[12:19], a[20:23], v24, v25 op_sel_hi:[0,0,0] ; encoding: [0x00,0x00,0xac,0xd3,0x18,0x33,0x02,0x00,0x00,0x88,0xad,0xd3,0x04,0x19,0x52,0x1c]
385// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU
386v_mfma_scale_f32_16x16x128_f8f6f4 a[0:3], a[4:11], a[12:19], a[20:23], v24, v25
387
388// GFX950: v_mfma_scale_f32_16x16x128_f8f6f4 a[0:3], v[4:11], v[12:19], a[20:23], v24, v25 op_sel_hi:[0,0,0] ; encoding: [0x00,0x00,0xac,0xd3,0x18,0x33,0x02,0x00,0x00,0x88,0xad,0xd3,0x04,0x19,0x52,0x04]
389// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU
390v_mfma_scale_f32_16x16x128_f8f6f4 a[0:3], v[4:11], v[12:19], a[20:23], v24, v25
391
392// GFX950: v_mfma_scale_f32_16x16x128_f8f6f4 v[0:3], a[4:11], a[12:19], v[20:23], v24, v25 op_sel_hi:[0,0,0] ; encoding: [0x00,0x00,0xac,0xd3,0x18,0x33,0x02,0x00,0x00,0x08,0xad,0xd3,0x04,0x19,0x52,0x1c]
393// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU
394v_mfma_scale_f32_16x16x128_f8f6f4 v[0:3], a[4:11], a[12:19], v[20:23], v24, v25
395
396// GFX950: v_mfma_scale_f32_16x16x128_f8f6f4 v[0:3], v[4:11], a[12:19], v[20:23], v24, v25 op_sel_hi:[0,0,0] ; encoding: [0x00,0x00,0xac,0xd3,0x18,0x33,0x02,0x00,0x00,0x08,0xad,0xd3,0x04,0x19,0x52,0x14]
397// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU
398v_mfma_scale_f32_16x16x128_f8f6f4 v[0:3], v[4:11], a[12:19], v[20:23], v24, v25
399
400// GFX950: v_mfma_scale_f32_16x16x128_f8f6f4 v[0:3], a[4:11], v[12:19], v[20:23], v24, v25 op_sel_hi:[0,0,0] ; encoding: [0x00,0x00,0xac,0xd3,0x18,0x33,0x02,0x00,0x00,0x08,0xad,0xd3,0x04,0x19,0x52,0x0c]
401// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU
402v_mfma_scale_f32_16x16x128_f8f6f4 v[0:3], a[4:11], v[12:19], v[20:23], v24, v25
403
404// GFX950: v_mfma_scale_f32_16x16x128_f8f6f4 v[50:53], v[4:11], v[12:19], v[20:23], v24, v25 op_sel_hi:[0,0,0] ; encoding: [0x00,0x00,0xac,0xd3,0x18,0x33,0x02,0x00,0x32,0x08,0xad,0xd3,0x04,0x19,0x52,0x04]
405// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU
406v_mfma_scale_f32_16x16x128_f8f6f4 v[50:53], v[4:11], v[12:19], v[20:23], v24, v25
407
408// GFX950: v_mfma_scale_f32_16x16x128_f8f6f4 v[0:3], v[4:11], v[12:19], v[20:23], v44, s24 op_sel_hi:[0,0,0] ; encoding: [0x00,0x00,0xac,0xd3,0x2c,0x31,0x00,0x00,0x00,0x08,0xad,0xd3,0x04,0x19,0x52,0x04]
409// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU
410v_mfma_scale_f32_16x16x128_f8f6f4 v[0:3], v[4:11], v[12:19], v[20:23], v44, s24
411
412// GFX950: v_mfma_scale_f32_16x16x128_f8f6f4 v[0:3], v[4:11], v[12:19], v[20:23], s24, s24 op_sel_hi:[0,0,0] ; encoding: [0x00,0x00,0xac,0xd3,0x18,0x30,0x00,0x00,0x00,0x08,0xad,0xd3,0x04,0x19,0x52,0x04]
413// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU
414v_mfma_scale_f32_16x16x128_f8f6f4 v[0:3], v[4:11], v[12:19], v[20:23], s24, s24
415
416// GFX950: v_mfma_scale_f32_16x16x128_f8f6f4 v[0:3], v[4:11], v[12:19], v[20:23], s24, v44 op_sel_hi:[0,0,0] ; encoding: [0x00,0x00,0xac,0xd3,0x18,0x58,0x02,0x00,0x00,0x08,0xad,0xd3,0x04,0x19,0x52,0x04]
417// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU
418v_mfma_scale_f32_16x16x128_f8f6f4 v[0:3], v[4:11], v[12:19], v[20:23], s24, v44
419
420// GFX950: v_mfma_scale_f32_16x16x128_f8f6f4 v[0:3], v[4:11], v[12:19], v[20:23], m0, m0 op_sel_hi:[0,0,0]  ; encoding: [0x00,0x00,0xac,0xd3,0x7c,0xf8,0x00,0x00,0x00,0x08,0xad,0xd3,0x04,0x19,0x52,0x04]
421// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU
422v_mfma_scale_f32_16x16x128_f8f6f4 v[0:3], v[4:11], v[12:19], v[20:23], m0, m0
423
424// GFX950: v_mfma_scale_f32_16x16x128_f8f6f4 v[0:3], v[4:11], v[12:19], v[20:23], vcc_lo, v2 op_sel_hi:[0,0,0] ; encoding: [0x00,0x00,0xac,0xd3,0x6a,0x04,0x02,0x00,0x00,0x08,0xad,0xd3,0x04,0x19,0x52,0x04]
425// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU
426v_mfma_scale_f32_16x16x128_f8f6f4 v[0:3], v[4:11], v[12:19], v[20:23], vcc_lo, v2
427
428// GFX950: v_mfma_scale_f32_16x16x128_f8f6f4 v[0:3], v[4:11], v[12:19], v[20:23], 9, v2 op_sel_hi:[0,0,0] ; encoding: [0x00,0x00,0xac,0xd3,0x89,0x04,0x02,0x00,0x00,0x08,0xad,0xd3,0x04,0x19,0x52,0x04]
429// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU
430v_mfma_scale_f32_16x16x128_f8f6f4 v[0:3], v[4:11], v[12:19], v[20:23], 9, v2
431
432// GFX950: v_mfma_scale_f32_16x16x128_f8f6f4 v[0:3], v[4:11], v[12:19], v[20:23], v2, 9 op_sel_hi:[0,0,0] ; encoding: [0x00,0x00,0xac,0xd3,0x02,0x13,0x01,0x00,0x00,0x08,0xad,0xd3,0x04,0x19,0x52,0x04]
433// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU
434v_mfma_scale_f32_16x16x128_f8f6f4 v[0:3], v[4:11], v[12:19], v[20:23], v2, 9
435
436// GFX950: v_mfma_scale_f32_16x16x128_f8f6f4 v[0:3], v[4:11], v[12:19], v[20:23], s20, 9 op_sel_hi:[0,0,0] ; encoding: [0x00,0x00,0xac,0xd3,0x14,0x12,0x01,0x00,0x00,0x08,0xad,0xd3,0x04,0x19,0x52,0x04]
437// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU
438v_mfma_scale_f32_16x16x128_f8f6f4 v[0:3], v[4:11], v[12:19], v[20:23], s20, 9
439
440// GFX950: v_mfma_scale_f32_16x16x128_f8f6f4 v[0:3], v[4:11], v[12:19], v[20:23], 33, 9 op_sel_hi:[0,0,0] ; encoding: [0x00,0x00,0xac,0xd3,0xa1,0x12,0x01,0x00,0x00,0x08,0xad,0xd3,0x04,0x19,0x52,0x04]
441// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU
442v_mfma_scale_f32_16x16x128_f8f6f4 v[0:3], v[4:11], v[12:19], v[20:23], 33, 9
443
444// GFX950: v_mfma_scale_f32_16x16x128_f8f6f4 v[0:3], v[4:9], v[12:19], v[20:23], v24, v25 op_sel_hi:[0,0,0] cbsz:3 blgp:1 ; encoding: [0x00,0x00,0xac,0xd3,0x18,0x33,0x02,0x00,0x00,0x0b,0xad,0xd3,0x04,0x19,0x52,0x24]
445// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU
446v_mfma_scale_f32_16x16x128_f8f6f4 v[0:3], v[4:9], v[12:19], v[20:23], v24, v25 cbsz:3 blgp:1
447
448// GFX950: v_mfma_scale_f32_16x16x128_f8f6f4 v[0:3], v[4:9], v[12:19], v[20:23], v24, v25 op_sel_hi:[0,0,0] cbsz:3 blgp:1 ; encoding: [0x00,0x00,0xac,0xd3,0x18,0x33,0x02,0x08,0x00,0x0b,0xad,0xd3,0x04,0x19,0x52,0x24]
449// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU
450v_mfma_scale_f32_16x16x128_f8f6f4 v[0:3], v[4:9], v[12:19], v[20:23], v24, v25 op_sel:[0,1,0] cbsz:3 blgp:1
451
452// GFX950: v_mfma_scale_f32_16x16x128_f8f6f4 v[0:3], v[4:9], v[12:19], v[20:23], v24, v25 op_sel_hi:[0,0,0] cbsz:3 blgp:1 ; encoding: [0x00,0x00,0xac,0xd3,0x18,0x33,0x02,0x18,0x00,0x0b,0xad,0xd3,0x04,0x19,0x52,0x24]
453// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU
454v_mfma_scale_f32_16x16x128_f8f6f4 v[0:3], v[4:9], v[12:19], v[20:23], v24, v25 op_sel:[0,1,0] op_sel_hi:[0,1,0] cbsz:3 blgp:1
455
456// GFX950: v_mfma_scale_f32_16x16x128_f8f6f4 v[0:3], v[4:11], v[12:17], v[20:23], v24, v25 op_sel_hi:[0,0,0] blgp:2 ; encoding: [0x00,0x00,0xac,0xd3,0x18,0x33,0x02,0x18,0x00,0x08,0xad,0xd3,0x04,0x19,0x52,0x44]
457// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU
458v_mfma_scale_f32_16x16x128_f8f6f4 v[0:3], v[4:11], v[12:17], v[20:23], v24, v25 op_sel:[0,1,0] op_sel_hi:[0,1,0] blgp:2
459
460// GFX950: v_mfma_scale_f32_16x16x128_f8f6f4 v[0:3], v[4:11], v[12:19], v[20:23], v24, v25 op_sel_hi:[0,0,0] ; encoding: [0x00,0x00,0xac,0xd3,0x18,0x33,0x02,0x18,0x00,0x08,0xad,0xd3,0x04,0x19,0x52,0x04]
461// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU
462v_mfma_scale_f32_16x16x128_f8f6f4 v[0:3], v[4:11], v[12:19], v[20:23], v24, v25 op_sel:[0,1,0] op_sel_hi:[0,1,0]
463
464// GFX950: v_mfma_scale_f32_16x16x128_f8f6f4 v[0:3], v[4:9], v[12:19], v[20:23], v24, v25 op_sel_hi:[0,0,0] cbsz:3 ; encoding: [0x00,0x00,0xac,0xd3,0x18,0x33,0x02,0x18,0x00,0x0b,0xad,0xd3,0x04,0x19,0x52,0x04]
465// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU
466v_mfma_scale_f32_16x16x128_f8f6f4 v[0:3], v[4:9], v[12:19], v[20:23], v24, v25 op_sel:[0,1,0] op_sel_hi:[0,1,0] cbsz:3
467
468//===----------------------------------------------------------------------===//
469// v_mfma_scale_f32_32x32x64_f8f6f4
470//===----------------------------------------------------------------------===//
471// FIXME: Test op_sel, neg, clamp
472
473// GFX950: v_mfma_scale_f32_32x32x64_f8f6f4 v[0:15], v[16:23], v[24:31], v[32:47], v48, v49 op_sel_hi:[0,0,0] ; encoding: [0x00,0x00,0xac,0xd3,0x30,0x63,0x02,0x00,0x00,0x08,0xae,0xd3,0x10,0x31,0x82,0x04]
474// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU
475v_mfma_scale_f32_32x32x64_f8f6f4 v[0:15], v[16:23], v[24:31], v[32:47], v48, v49
476
477// GFX950: v_mfma_scale_f32_32x32x64_f8f6f4 a[0:15], a[16:23], a[24:31], a[32:47], v48, v49 op_sel_hi:[0,0,0] ; encoding: [0x00,0x00,0xac,0xd3,0x30,0x63,0x02,0x00,0x00,0x88,0xae,0xd3,0x10,0x31,0x82,0x1c]
478// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU
479v_mfma_scale_f32_32x32x64_f8f6f4 a[0:15], a[16:23], a[24:31], a[32:47], v48, v49
480
481// GFX950: v_mfma_scale_f32_32x32x64_f8f6f4 a[0:15], v[16:23], v[24:31], a[32:47], v48, v49 op_sel_hi:[0,0,0] ; encoding: [0x00,0x00,0xac,0xd3,0x30,0x63,0x02,0x00,0x00,0x88,0xae,0xd3,0x10,0x31,0x82,0x04]
482// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU
483v_mfma_scale_f32_32x32x64_f8f6f4 a[0:15], v[16:23], v[24:31], a[32:47], v48, v49
484
485// GFX950: v_mfma_scale_f32_32x32x64_f8f6f4 v[0:15], a[16:23], a[24:31], v[32:47], v48, v49 op_sel_hi:[0,0,0] ; encoding: [0x00,0x00,0xac,0xd3,0x30,0x63,0x02,0x00,0x00,0x08,0xae,0xd3,0x10,0x31,0x82,0x1c]
486// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU
487v_mfma_scale_f32_32x32x64_f8f6f4 v[0:15], a[16:23], a[24:31], v[32:47], v48, v49
488
489// GFX950: v_mfma_scale_f32_32x32x64_f8f6f4 v[0:15], v[16:23], a[24:31], v[32:47], v48, v49 op_sel_hi:[0,0,0] ; encoding: [0x00,0x00,0xac,0xd3,0x30,0x63,0x02,0x00,0x00,0x08,0xae,0xd3,0x10,0x31,0x82,0x14]
490// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU
491v_mfma_scale_f32_32x32x64_f8f6f4 v[0:15], v[16:23], a[24:31], v[32:47], v48, v49
492
493// GFX950: v_mfma_scale_f32_32x32x64_f8f6f4 a[0:15], a[16:23], v[24:31], a[32:47], v48, v49 op_sel_hi:[0,0,0] ; encoding: [0x00,0x00,0xac,0xd3,0x30,0x63,0x02,0x00,0x00,0x88,0xae,0xd3,0x10,0x31,0x82,0x0c]
494// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU
495v_mfma_scale_f32_32x32x64_f8f6f4 a[0:15], a[16:23], v[24:31], a[32:47], v48, v49
496
497// GFX950: v_mfma_scale_f32_32x32x64_f8f6f4 v[50:65], v[16:23], v[24:31], v[32:47], v48, v49  op_sel_hi:[0,0,0] ; encoding: [0x00,0x00,0xac,0xd3,0x30,0x63,0x02,0x00,0x32,0x08,0xae,0xd3,0x10,0x31,0x82,0x04]
498// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU
499v_mfma_scale_f32_32x32x64_f8f6f4 v[50:65], v[16:23], v[24:31], v[32:47], v48, v49
500
501// GFX950: v_mfma_scale_f32_32x32x64_f8f6f4 v[0:15], v[16:21], v[24:31], v[32:47], v48, v49 op_sel_hi:[0,0,0] cbsz:3 blgp:1 ; encoding: [0x00,0x00,0xac,0xd3,0x30,0x63,0x02,0x00,0x00,0x0b,0xae,0xd3,0x10,0x31,0x82,0x24]
502// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU
503v_mfma_scale_f32_32x32x64_f8f6f4 v[0:15], v[16:21], v[24:31], v[32:47], v48, v49 cbsz:3 blgp:1
504
505// GFX950: 	v_mfma_scale_f32_32x32x64_f8f6f4 v[0:15], v[16:21], v[24:31], v[32:47], v48, v49 op_sel_hi:[0,0,0] cbsz:2 ; encoding: [0x00,0x00,0xac,0xd3,0x30,0x63,0x02,0x00,0x00,0x0a,0xae,0xd3,0x10,0x31,0x82,0x04]
506// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU
507v_mfma_scale_f32_32x32x64_f8f6f4 v[0:15], v[16:21], v[24:31], v[32:47], v48, v49 cbsz:2
508
509// GFX950: v_mfma_scale_f32_32x32x64_f8f6f4 v[0:15], v[16:23], v[24:31], v[32:47], v48, v49 op_sel_hi:[0,0,0] ; encoding: [0x00,0x00,0xac,0xd3,0x30,0x63,0x02,0x00,0x00,0x08,0xae,0xd3,0x10,0x31,0x82,0x04]
510// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU
511v_mfma_scale_f32_32x32x64_f8f6f4 v[0:15], v[16:23], v[24:31], v[32:47], v48, v49
512
513// GFX950: v_mfma_scale_f32_32x32x64_f8f6f4 v[0:15], v[16:23], v[24:29], v[32:47], v48, v49 op_sel_hi:[0,0,0] blgp:2 ; encoding: [0x00,0x00,0xac,0xd3,0x30,0x63,0x02,0x00,0x00,0x08,0xae,0xd3,0x10,0x31,0x82,0x44]
514// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU
515v_mfma_scale_f32_32x32x64_f8f6f4 v[0:15], v[16:23], v[24:29], v[32:47], v48, v49 blgp:2
516
517// GFX950: v_mfma_scale_f32_32x32x64_f8f6f4 v[0:15], v[16:21], v[24:29], v[32:47], v48, v49 op_sel_hi:[0,0,0] cbsz:2 blgp:3 ; encoding: [0x00,0x00,0xac,0xd3,0x30,0x63,0x02,0x18,0x00,0x0a,0xae,0xd3,0x10,0x31,0x82,0x64]
518// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU
519v_mfma_scale_f32_32x32x64_f8f6f4 v[0:15], v[16:21], v[24:29], v[32:47], v48, v49 op_sel:[0,1,0] op_sel_hi:[0,1,0] cbsz:2 blgp:3
520
521//===----------------------------------------------------------------------===//
522// v_mfma_f32_16x16x128_f8f6f4 with appropriate register widths
523//===----------------------------------------------------------------------===//
524
525// bf8 x fp4
526// GFX950: v_mfma_f32_16x16x128_f8f6f4 v[0:3], v[4:11], v[12:15], v[20:23] cbsz:1 blgp:4 ; encoding: [0x00,0x01,0xad,0xd3,0x04,0x19,0x52,0x84]
527v_mfma_f32_16x16x128_f8f6f4 v[0:3], v[4:11], v[12:15], v[20:23] cbsz:1 blgp:4
528
529// GFX950: v_mfma_f32_16x16x128_f8f6f4 v[0:3], v[4:11], v[12:15], v[20:23] cbsz:1 blgp:4 ; encoding: [0x00,0x01,0xad,0xd3,0x04,0x19,0x52,0x84]
530v_mfma_f32_16x16x128_f8f6f4 v[0:3], v[4:11], v[12:15], v[20:23] cbsz:1 blgp:4
531
532// GFX950: v_mfma_f32_16x16x128_f8f6f4 a[0:3], a[4:11], a[12:15], a[20:23] cbsz:1 blgp:4 ; encoding: [0x00,0x81,0xad,0xd3,0x04,0x19,0x52,0x9c]
533v_mfma_f32_16x16x128_f8f6f4 a[0:3], a[4:11], a[12:15], a[20:23] cbsz:1 blgp:4
534
535
536// fp4 x bf8
537// GFX950: v_mfma_f32_16x16x128_f8f6f4 v[0:3], v[4:7], v[12:19], v[20:23] cbsz:4 blgp:1 ; encoding: [0x00,0x04,0xad,0xd3,0x04,0x19,0x52,0x24]
538v_mfma_f32_16x16x128_f8f6f4 v[0:3], v[4:7], v[12:19], v[20:23] cbsz:4 blgp:1
539
540// GFX950: v_mfma_f32_16x16x128_f8f6f4 a[0:3], a[4:7], a[12:19], a[20:23] cbsz:4 blgp:1 ; encoding: [0x00,0x84,0xad,0xd3,0x04,0x19,0x52,0x3c]
541v_mfma_f32_16x16x128_f8f6f4 a[0:3], a[4:7], a[12:19], a[20:23] cbsz:4 blgp:1
542
543
544// bf6 x bf8
545// GFX950: v_mfma_f32_16x16x128_f8f6f4 v[0:3], v[4:9], v[12:19], v[20:23] cbsz:3 blgp:1 ; encoding: [0x00,0x03,0xad,0xd3,0x04,0x19,0x52,0x24]
546v_mfma_f32_16x16x128_f8f6f4 v[0:3], v[4:9], v[12:19], v[20:23] cbsz:3 blgp:1
547
548// GFX950: v_mfma_f32_16x16x128_f8f6f4 a[0:3], a[4:9], a[12:19], a[20:23] cbsz:3 blgp:1 ; encoding: [0x00,0x83,0xad,0xd3,0x04,0x19,0x52,0x3c]
549v_mfma_f32_16x16x128_f8f6f4 a[0:3], a[4:9], a[12:19], a[20:23] cbsz:3 blgp:1
550
551
552// bf8 x bf6
553// GFX950: v_mfma_f32_16x16x128_f8f6f4 v[0:3], v[4:11], v[12:17], v[20:23] cbsz:1 blgp:3 ; encoding: [0x00,0x01,0xad,0xd3,0x04,0x19,0x52,0x64]
554v_mfma_f32_16x16x128_f8f6f4 v[0:3], v[4:11], v[12:17], v[20:23] cbsz:1 blgp:3
555
556// GFX950: v_mfma_f32_16x16x128_f8f6f4 a[0:3], a[4:11], a[12:17], a[20:23] cbsz:1 blgp:3 ; encoding: [0x00,0x81,0xad,0xd3,0x04,0x19,0x52,0x7c]
557v_mfma_f32_16x16x128_f8f6f4 a[0:3], a[4:11], a[12:17], a[20:23] cbsz:1 blgp:3
558
559
560// bf6 x bf6
561// GFX950: v_mfma_f32_16x16x128_f8f6f4 v[0:3], v[4:9], v[12:17], v[20:23] cbsz:3 blgp:3 ; encoding: [0x00,0x03,0xad,0xd3,0x04,0x19,0x52,0x64]
562v_mfma_f32_16x16x128_f8f6f4 v[0:3], v[4:9], v[12:17], v[20:23] cbsz:3 blgp:3
563
564// GFX950: v_mfma_f32_16x16x128_f8f6f4 a[0:3], a[4:9], a[12:17], a[20:23] cbsz:3 blgp:3 ; encoding: [0x00,0x83,0xad,0xd3,0x04,0x19,0x52,0x7c]
565v_mfma_f32_16x16x128_f8f6f4 a[0:3], a[4:9], a[12:17], a[20:23] cbsz:3 blgp:3
566
567// bf6 x fp6
568// GFX950: v_mfma_f32_16x16x128_f8f6f4 v[0:3], v[4:9], v[12:17], v[20:23] cbsz:3 blgp:2 ; encoding: [0x00,0x03,0xad,0xd3,0x04,0x19,0x52,0x44]
569v_mfma_f32_16x16x128_f8f6f4 v[0:3], v[4:9], v[12:17], v[20:23] cbsz:3 blgp:2
570
571// GFX950: v_mfma_f32_16x16x128_f8f6f4 a[0:3], a[4:9], a[12:17], a[20:23] cbsz:3 blgp:2 ; encoding: [0x00,0x83,0xad,0xd3,0x04,0x19,0x52,0x5c]
572v_mfma_f32_16x16x128_f8f6f4 a[0:3], a[4:9], a[12:17], a[20:23] cbsz:3 blgp:2
573
574
575// fp6 x bf6
576// GFX950: v_mfma_f32_16x16x128_f8f6f4 v[0:3], v[4:9], v[12:17], v[20:23] cbsz:2 blgp:3 ; encoding: [0x00,0x02,0xad,0xd3,0x04,0x19,0x52,0x64]
577v_mfma_f32_16x16x128_f8f6f4 v[0:3], v[4:9], v[12:17], v[20:23] cbsz:2 blgp:3
578
579// GFX950: v_mfma_f32_16x16x128_f8f6f4 a[0:3], a[4:9], a[12:17], a[20:23] cbsz:2 blgp:3 ; encoding: [0x00,0x82,0xad,0xd3,0x04,0x19,0x52,0x7c]
580v_mfma_f32_16x16x128_f8f6f4 a[0:3], a[4:9], a[12:17], a[20:23] cbsz:2 blgp:3
581
582
583// fp6 x fp4
584// GFX950: v_mfma_f32_16x16x128_f8f6f4 v[0:3], v[4:9], v[12:15], v[20:23] cbsz:2 blgp:4 ; encoding: [0x00,0x02,0xad,0xd3,0x04,0x19,0x52,0x84]
585v_mfma_f32_16x16x128_f8f6f4 v[0:3], v[4:9], v[12:15], v[20:23] cbsz:2 blgp:4
586
587// GFX950: v_mfma_f32_16x16x128_f8f6f4 a[0:3], a[4:9], a[12:15], a[20:23] cbsz:2 blgp:4 ; encoding: [0x00,0x82,0xad,0xd3,0x04,0x19,0x52,0x9c]
588v_mfma_f32_16x16x128_f8f6f4 a[0:3], a[4:9], a[12:15], a[20:23] cbsz:2 blgp:4
589
590
591// fp4 x fp6
592// GFX950: v_mfma_f32_16x16x128_f8f6f4 v[0:3], v[4:7], v[12:17], v[20:23] cbsz:4 blgp:2 ; encoding: [0x00,0x04,0xad,0xd3,0x04,0x19,0x52,0x44]
593v_mfma_f32_16x16x128_f8f6f4 v[0:3], v[4:7], v[12:17], v[20:23] cbsz:4 blgp:2
594
595// GFX950: v_mfma_f32_16x16x128_f8f6f4 a[0:3], a[4:7], a[12:17], a[20:23] cbsz:4 blgp:2 ; encoding: [0x00,0x84,0xad,0xd3,0x04,0x19,0x52,0x5c]
596v_mfma_f32_16x16x128_f8f6f4 a[0:3], a[4:7], a[12:17], a[20:23] cbsz:4 blgp:2
597
598
599// fp4 x fp4
600// GFX950: v_mfma_f32_16x16x128_f8f6f4 v[0:3], v[4:7], v[12:15], v[20:23] cbsz:4 blgp:4 ; encoding: [0x00,0x04,0xad,0xd3,0x04,0x19,0x52,0x84]
601v_mfma_f32_16x16x128_f8f6f4 v[0:3], v[4:7], v[12:15], v[20:23] cbsz:4 blgp:4
602
603// GFX950: v_mfma_f32_16x16x128_f8f6f4 a[0:3], a[4:7], a[12:15], a[20:23] cbsz:4 blgp:4 ; encoding: [0x00,0x84,0xad,0xd3,0x04,0x19,0x52,0x9c]
604v_mfma_f32_16x16x128_f8f6f4 a[0:3], a[4:7], a[12:15], a[20:23] cbsz:4 blgp:4
605
606//===----------------------------------------------------------------------===//
607// v_mfma_f32_32x32x64_f8f6f4 with appropriate register widths
608//===----------------------------------------------------------------------===//
609
610// bf8 x fp4
611// GFX950: v_mfma_f32_32x32x64_f8f6f4 v[0:15], v[4:11], v[12:15], v[16:31] cbsz:1 blgp:4 ; encoding: [0x00,0x01,0xae,0xd3,0x04,0x19,0x42,0x84]
612v_mfma_f32_32x32x64_f8f6f4 v[0:15], v[4:11], v[12:15], v[16:31] cbsz:1 blgp:4
613
614// GFX950: v_mfma_f32_32x32x64_f8f6f4 a[0:15], a[4:11], a[12:15], a[16:31] cbsz:1 blgp:4 ; encoding: [0x00,0x81,0xae,0xd3,0x04,0x19,0x42,0x9c]
615v_mfma_f32_32x32x64_f8f6f4 a[0:15], a[4:11], a[12:15], a[16:31] cbsz:1 blgp:4
616
617
618// fp4 x bf8
619// GFX950: v_mfma_f32_32x32x64_f8f6f4 v[0:15], v[4:7], v[8:15], v[16:31] cbsz:4 blgp:1 ; encoding: [0x00,0x04,0xae,0xd3,0x04,0x11,0x42,0x24]
620v_mfma_f32_32x32x64_f8f6f4 v[0:15], v[4:7], v[8:15], v[16:31] cbsz:4 blgp:1
621
622// GFX950: v_mfma_f32_32x32x64_f8f6f4 a[0:15], a[4:7], a[8:15], a[16:31] cbsz:4 blgp:1 ; encoding: [0x00,0x84,0xae,0xd3,0x04,0x11,0x42,0x3c]
623v_mfma_f32_32x32x64_f8f6f4 a[0:15], a[4:7], a[8:15], a[16:31] cbsz:4 blgp:1
624
625
626// bf6 x bf8
627// GFX950: v_mfma_f32_32x32x64_f8f6f4 v[0:15], v[4:9], v[12:19], v[16:31] cbsz:3 blgp:1 ; encoding: [0x00,0x03,0xae,0xd3,0x04,0x19,0x42,0x24]
628v_mfma_f32_32x32x64_f8f6f4 v[0:15], v[4:9], v[12:19], v[16:31] cbsz:3 blgp:1
629
630// GFX950: v_mfma_f32_32x32x64_f8f6f4 a[0:15], a[4:9], a[12:19], a[16:31] cbsz:3 blgp:1 ; encoding: [0x00,0x83,0xae,0xd3,0x04,0x19,0x42,0x3c]
631v_mfma_f32_32x32x64_f8f6f4 a[0:15], a[4:9], a[12:19], a[16:31] cbsz:3 blgp:1
632
633// bf8 x bf6
634// GFX950: v_mfma_f32_32x32x64_f8f6f4 v[0:15], v[4:11], v[12:17], v[16:31] cbsz:1 blgp:3 ; encoding: [0x00,0x01,0xae,0xd3,0x04,0x19,0x42,0x64]
635v_mfma_f32_32x32x64_f8f6f4 v[0:15], v[4:11], v[12:17], v[16:31] cbsz:1 blgp:3
636
637// GFX950: v_mfma_f32_32x32x64_f8f6f4 a[0:15], a[4:11], a[12:17], a[16:31] cbsz:1 blgp:3 ; encoding: [0x00,0x81,0xae,0xd3,0x04,0x19,0x42,0x7c]
638v_mfma_f32_32x32x64_f8f6f4 a[0:15], a[4:11], a[12:17], a[16:31] cbsz:1 blgp:3
639
640
641// bf6 x bf6
642// GFX950: v_mfma_f32_32x32x64_f8f6f4 v[0:15], v[4:9], v[12:17], v[16:31] cbsz:3 blgp:3 ; encoding: [0x00,0x03,0xae,0xd3,0x04,0x19,0x42,0x64]
643v_mfma_f32_32x32x64_f8f6f4 v[0:15], v[4:9], v[12:17], v[16:31] cbsz:3 blgp:3
644
645// GFX950: v_mfma_f32_32x32x64_f8f6f4 a[0:15], a[4:9], a[12:17], a[16:31] cbsz:3 blgp:3 ; encoding: [0x00,0x83,0xae,0xd3,0x04,0x19,0x42,0x7c]
646v_mfma_f32_32x32x64_f8f6f4 a[0:15], a[4:9], a[12:17], a[16:31] cbsz:3 blgp:3
647
648
649// bf6 x fp6
650// GFX950: v_mfma_f32_32x32x64_f8f6f4 v[0:15], v[4:9], v[12:17], v[16:31] cbsz:3 blgp:2 ; encoding: [0x00,0x03,0xae,0xd3,0x04,0x19,0x42,0x44]
651v_mfma_f32_32x32x64_f8f6f4 v[0:15], v[4:9], v[12:17], v[16:31] cbsz:3 blgp:2
652
653// GFX950: v_mfma_f32_32x32x64_f8f6f4 a[0:15], a[4:9], a[12:17], a[16:31] cbsz:3 blgp:2 ; encoding: [0x00,0x83,0xae,0xd3,0x04,0x19,0x42,0x5c]
654v_mfma_f32_32x32x64_f8f6f4 a[0:15], a[4:9], a[12:17], a[16:31] cbsz:3 blgp:2
655
656
657// fp6 x bf6
658// GFX950: v_mfma_f32_32x32x64_f8f6f4 v[0:15], v[4:9], v[12:17], v[16:31] cbsz:2 blgp:3 ; encoding: [0x00,0x02,0xae,0xd3,0x04,0x19,0x42,0x64]
659v_mfma_f32_32x32x64_f8f6f4 v[0:15], v[4:9], v[12:17], v[16:31] cbsz:2 blgp:3
660
661// GFX950: v_mfma_f32_32x32x64_f8f6f4 a[0:15], a[4:9], a[12:17], a[16:31] cbsz:2 blgp:3 ; encoding: [0x00,0x82,0xae,0xd3,0x04,0x19,0x42,0x7c]
662v_mfma_f32_32x32x64_f8f6f4 a[0:15], a[4:9], a[12:17], a[16:31] cbsz:2 blgp:3
663
664
665// fp6 x fp4
666// GFX950: v_mfma_f32_32x32x64_f8f6f4 v[0:15], v[4:9], v[12:15], v[16:31] cbsz:2 blgp:4 ; encoding: [0x00,0x02,0xae,0xd3,0x04,0x19,0x42,0x84]
667v_mfma_f32_32x32x64_f8f6f4 v[0:15], v[4:9], v[12:15], v[16:31] cbsz:2 blgp:4
668
669// GFX950: v_mfma_f32_32x32x64_f8f6f4 a[0:15], a[4:9], a[12:15], a[16:31] cbsz:2 blgp:4 ; encoding: [0x00,0x82,0xae,0xd3,0x04,0x19,0x42,0x9c]
670v_mfma_f32_32x32x64_f8f6f4 a[0:15], a[4:9], a[12:15], a[16:31] cbsz:2 blgp:4
671
672
673// fp4 x fp6
674// GFX950: v_mfma_f32_32x32x64_f8f6f4 v[0:15], v[4:7], v[12:17], v[16:31] cbsz:4 blgp:2 ; encoding: [0x00,0x04,0xae,0xd3,0x04,0x19,0x42,0x44]
675v_mfma_f32_32x32x64_f8f6f4 v[0:15], v[4:7], v[12:17], v[16:31] cbsz:4 blgp:2
676
677// GFX950: v_mfma_f32_32x32x64_f8f6f4 a[0:15], a[4:7], a[12:17], a[16:31] cbsz:4 blgp:2 ; encoding: [0x00,0x84,0xae,0xd3,0x04,0x19,0x42,0x5c]
678v_mfma_f32_32x32x64_f8f6f4 a[0:15], a[4:7], a[12:17], a[16:31] cbsz:4 blgp:2
679
680
681// fp4 x fp4
682// GFX950: v_mfma_f32_32x32x64_f8f6f4 v[0:15], v[4:7], v[12:15], v[16:31] cbsz:4 blgp:4 ; encoding: [0x00,0x04,0xae,0xd3,0x04,0x19,0x42,0x84]
683v_mfma_f32_32x32x64_f8f6f4 v[0:15], v[4:7], v[12:15], v[16:31] cbsz:4 blgp:4
684
685// GFX950: v_mfma_f32_32x32x64_f8f6f4 a[0:15], a[4:7], a[12:15], a[16:31] cbsz:4 blgp:4 ; encoding: [0x00,0x84,0xae,0xd3,0x04,0x19,0x42,0x9c]
686v_mfma_f32_32x32x64_f8f6f4 a[0:15], a[4:7], a[12:15], a[16:31] cbsz:4 blgp:4
687
688//===----------------------------------------------------------------------===//
689// v_mfma_scale_f32_16x16x128_f8f6f4 corrected widths
690//===----------------------------------------------------------------------===//
691
692// fp8 x fp8
693// GFX950: v_mfma_scale_f32_16x16x128_f8f6f4 v[0:3], v[4:11], v[12:19], v[20:23], v24, v25 op_sel_hi:[0,0,0] ; encoding: [0x00,0x00,0xac,0xd3,0x18,0x33,0x02,0x00,0x00,0x08,0xad,0xd3,0x04,0x19,0x52,0x04]
694v_mfma_scale_f32_16x16x128_f8f6f4 v[0:3], v[4:11], v[12:19], v[20:23], v24, v25 cbsz:0 blgp:0
695
696// GFX950: v_mfma_scale_f32_16x16x128_f8f6f4 a[0:3], a[4:11], a[12:19], a[20:23], v24, v25 op_sel_hi:[0,0,0] ; encoding: [0x00,0x00,0xac,0xd3,0x18,0x33,0x02,0x00,0x00,0x88,0xad,0xd3,0x04,0x19,0x52,0x1c]
697v_mfma_scale_f32_16x16x128_f8f6f4 a[0:3], a[4:11], a[12:19], a[20:23], v24, v25 cbsz:0 blgp:0
698
699
700// bf8 x fp8
701// GFX950: v_mfma_scale_f32_16x16x128_f8f6f4 v[0:3], v[4:11], v[12:19], v[20:23], v24, v25 op_sel_hi:[0,0,0] cbsz:1 ; encoding: [0x00,0x00,0xac,0xd3,0x18,0x33,0x02,0x00,0x00,0x09,0xad,0xd3,0x04,0x19,0x52,0x04]
702v_mfma_scale_f32_16x16x128_f8f6f4 v[0:3], v[4:11], v[12:19], v[20:23], v24, v25 cbsz:1 blgp:0
703
704// GFX950: v_mfma_scale_f32_16x16x128_f8f6f4 a[0:3], a[4:11], a[12:19], a[20:23], v24, v25 op_sel_hi:[0,0,0] cbsz:1 ; encoding: [0x00,0x00,0xac,0xd3,0x18,0x33,0x02,0x00,0x00,0x89,0xad,0xd3,0x04,0x19,0x52,0x1c]
705v_mfma_scale_f32_16x16x128_f8f6f4 a[0:3], a[4:11], a[12:19], a[20:23], v24, v25 cbsz:1 blgp:0
706
707// fp8 x bf8
708// GFX950: v_mfma_scale_f32_16x16x128_f8f6f4 v[0:3], v[4:11], v[12:19], v[20:23], v24, v25 op_sel_hi:[0,0,0] blgp:1 ; encoding: [0x00,0x00,0xac,0xd3,0x18,0x33,0x02,0x00,0x00,0x08,0xad,0xd3,0x04,0x19,0x52,0x24]
709v_mfma_scale_f32_16x16x128_f8f6f4 v[0:3], v[4:11], v[12:19], v[20:23], v24, v25 cbsz:0 blgp:1
710
711// GFX950: v_mfma_scale_f32_16x16x128_f8f6f4 a[0:3], a[4:11], a[12:19], a[20:23], v24, v25 op_sel_hi:[0,0,0] blgp:1 ; encoding: [0x00,0x00,0xac,0xd3,0x18,0x33,0x02,0x00,0x00,0x88,0xad,0xd3,0x04,0x19,0x52,0x3c]
712v_mfma_scale_f32_16x16x128_f8f6f4 a[0:3], a[4:11], a[12:19], a[20:23], v24, v25 cbsz:0 blgp:1
713
714
715// bf8 x fp4
716// GFX950: v_mfma_scale_f32_16x16x128_f8f6f4 v[0:3], v[4:11], v[12:15], v[20:23], v24, v25 op_sel_hi:[0,0,0] cbsz:1 blgp:4 ; encoding: [0x00,0x00,0xac,0xd3,0x18,0x33,0x02,0x00,0x00,0x09,0xad,0xd3,0x04,0x19,0x52,0x84]
717v_mfma_scale_f32_16x16x128_f8f6f4 v[0:3], v[4:11], v[12:15], v[20:23], v24, v25 cbsz:1 blgp:4
718
719// GFX950: v_mfma_scale_f32_16x16x128_f8f6f4 a[0:3], a[4:11], a[12:15], a[20:23], v24, v25 op_sel_hi:[0,0,0] cbsz:1 blgp:4 ; encoding: [0x00,0x00,0xac,0xd3,0x18,0x33,0x02,0x00,0x00,0x89,0xad,0xd3,0x04,0x19,0x52,0x9c]
720v_mfma_scale_f32_16x16x128_f8f6f4 a[0:3], a[4:11], a[12:15], a[20:23], v24, v25 cbsz:1 blgp:4
721
722// fp4 x bf8
723// GFX950: v_mfma_scale_f32_16x16x128_f8f6f4 v[0:3], v[4:7], v[12:19], v[20:23], v24, v25 op_sel_hi:[0,0,0] cbsz:4 blgp:1 ; encoding: [0x00,0x00,0xac,0xd3,0x18,0x33,0x02,0x00,0x00,0x0c,0xad,0xd3,0x04,0x19,0x52,0x24]
724v_mfma_scale_f32_16x16x128_f8f6f4 v[0:3], v[4:7], v[12:19], v[20:23], v24, v25 cbsz:4 blgp:1
725
726// GFX950: v_mfma_scale_f32_16x16x128_f8f6f4 a[0:3], a[4:7], a[12:19], a[20:23], v24, v25 op_sel_hi:[0,0,0] cbsz:4 blgp:1 ; encoding: [0x00,0x00,0xac,0xd3,0x18,0x33,0x02,0x00,0x00,0x8c,0xad,0xd3,0x04,0x19,0x52,0x3c]
727v_mfma_scale_f32_16x16x128_f8f6f4 a[0:3], a[4:7], a[12:19], a[20:23], v24, v25 cbsz:4 blgp:1
728
729
730// bf6 x bf8
731// GFX950: v_mfma_scale_f32_16x16x128_f8f6f4 v[0:3], v[4:9], v[12:19], v[20:23], v24, v25 op_sel_hi:[0,0,0] cbsz:3 blgp:1 ; encoding: [0x00,0x00,0xac,0xd3,0x18,0x33,0x02,0x00,0x00,0x0b,0xad,0xd3,0x04,0x19,0x52,0x24]
732v_mfma_scale_f32_16x16x128_f8f6f4 v[0:3], v[4:9], v[12:19], v[20:23], v24, v25 cbsz:3 blgp:1
733
734// GFX950: v_mfma_scale_f32_16x16x128_f8f6f4 a[0:3], a[4:9], a[12:19], a[20:23], v24, v25 op_sel_hi:[0,0,0] cbsz:3 blgp:1 ; encoding: [0x00,0x00,0xac,0xd3,0x18,0x33,0x02,0x00,0x00,0x8b,0xad,0xd3,0x04,0x19,0x52,0x3c]
735v_mfma_scale_f32_16x16x128_f8f6f4 a[0:3], a[4:9], a[12:19], a[20:23], v24, v25 cbsz:3 blgp:1
736
737
738// bf8 x bf6
739// GFX950: v_mfma_scale_f32_16x16x128_f8f6f4 v[0:3], v[4:11], v[12:17], v[20:23], v24, v25 op_sel_hi:[0,0,0] cbsz:1 blgp:3 ; encoding: [0x00,0x00,0xac,0xd3,0x18,0x33,0x02,0x00,0x00,0x09,0xad,0xd3,0x04,0x19,0x52,0x64]
740v_mfma_scale_f32_16x16x128_f8f6f4 v[0:3], v[4:11], v[12:17], v[20:23], v24, v25 cbsz:1 blgp:3
741
742// GFX950: v_mfma_scale_f32_16x16x128_f8f6f4 a[0:3], a[4:11], a[12:17], a[20:23], v24, v25 op_sel_hi:[0,0,0] cbsz:1 blgp:3 ; encoding: [0x00,0x00,0xac,0xd3,0x18,0x33,0x02,0x00,0x00,0x89,0xad,0xd3,0x04,0x19,0x52,0x7c]
743v_mfma_scale_f32_16x16x128_f8f6f4 a[0:3], a[4:11], a[12:17], a[20:23], v24, v25 cbsz:1 blgp:3
744
745
746// bf6 x bf6
747// GFX950: v_mfma_scale_f32_16x16x128_f8f6f4 v[0:3], v[4:9], v[12:17], v[20:23], v24, v25 op_sel_hi:[0,0,0] cbsz:3 blgp:3 ; encoding: [0x00,0x00,0xac,0xd3,0x18,0x33,0x02,0x00,0x00,0x0b,0xad,0xd3,0x04,0x19,0x52,0x64]
748v_mfma_scale_f32_16x16x128_f8f6f4 v[0:3], v[4:9], v[12:17], v[20:23], v24, v25 cbsz:3 blgp:3
749
750// GFX950: v_mfma_scale_f32_16x16x128_f8f6f4 a[0:3], a[4:9], a[12:17], a[20:23], v24, v25 op_sel_hi:[0,0,0] cbsz:3 blgp:3 ; encoding: [0x00,0x00,0xac,0xd3,0x18,0x33,0x02,0x00,0x00,0x8b,0xad,0xd3,0x04,0x19,0x52,0x7c]
751v_mfma_scale_f32_16x16x128_f8f6f4 a[0:3], a[4:9], a[12:17], a[20:23], v24, v25 cbsz:3 blgp:3
752
753
754// bf6 x fp6
755// GFX950: v_mfma_scale_f32_16x16x128_f8f6f4 v[0:3], v[4:9], v[12:17], v[20:23], v24, v25 op_sel_hi:[0,0,0] cbsz:3 blgp:2 ; encoding: [0x00,0x00,0xac,0xd3,0x18,0x33,0x02,0x00,0x00,0x0b,0xad,0xd3,0x04,0x19,0x52,0x44]
756v_mfma_scale_f32_16x16x128_f8f6f4 v[0:3], v[4:9], v[12:17], v[20:23], v24, v25 cbsz:3 blgp:2
757
758// GFX950: v_mfma_scale_f32_16x16x128_f8f6f4 a[0:3], a[4:9], a[12:17], a[20:23], v24, v25 op_sel_hi:[0,0,0] cbsz:3 blgp:2 ; encoding: [0x00,0x00,0xac,0xd3,0x18,0x33,0x02,0x00,0x00,0x8b,0xad,0xd3,0x04,0x19,0x52,0x5c]
759v_mfma_scale_f32_16x16x128_f8f6f4 a[0:3], a[4:9], a[12:17], a[20:23], v24, v25 cbsz:3 blgp:2
760
761
762// fp6 x bf6
763// GFX950: v_mfma_scale_f32_16x16x128_f8f6f4 v[0:3], v[4:9], v[12:17], v[20:23], v24, v25 op_sel_hi:[0,0,0] cbsz:2 blgp:3 ; encoding: [0x00,0x00,0xac,0xd3,0x18,0x33,0x02,0x00,0x00,0x0a,0xad,0xd3,0x04,0x19,0x52,0x64
764v_mfma_scale_f32_16x16x128_f8f6f4 v[0:3], v[4:9], v[12:17], v[20:23], v24, v25 cbsz:2 blgp:3
765
766// GFX950: v_mfma_scale_f32_16x16x128_f8f6f4 a[0:3], a[4:9], a[12:17], a[20:23], v24, v25 op_sel_hi:[0,0,0] cbsz:2 blgp:3 ; encoding: [0x00,0x00,0xac,0xd3,0x18,0x33,0x02,0x00,0x00,0x8a,0xad,0xd3,0x04,0x19,0x52,0x7c]
767v_mfma_scale_f32_16x16x128_f8f6f4 a[0:3], a[4:9], a[12:17], a[20:23], v24, v25 cbsz:2 blgp:3
768
769
770// fp6 x fp4
771// GFX950: v_mfma_scale_f32_16x16x128_f8f6f4 v[0:3], v[4:9], v[12:15], v[20:23], v24, v25 op_sel_hi:[0,0,0] cbsz:2 blgp:4 ; encoding: [0x00,0x00,0xac,0xd3,0x18,0x33,0x02,0x00,0x00,0x0a,0xad,0xd3,0x04,0x19,0x52,0x84]
772v_mfma_scale_f32_16x16x128_f8f6f4 v[0:3], v[4:9], v[12:15], v[20:23], v24, v25 cbsz:2 blgp:4
773
774// GFX950: v_mfma_scale_f32_16x16x128_f8f6f4 a[0:3], a[4:9], a[12:15], a[20:23], v24, v25 op_sel_hi:[0,0,0] cbsz:2 blgp:4 ; encoding: [0x00,0x00,0xac,0xd3,0x18,0x33,0x02,0x00,0x00,0x8a,0xad,0xd3,0x04,0x19,0x52,0x9c]
775v_mfma_scale_f32_16x16x128_f8f6f4 a[0:3], a[4:9], a[12:15], a[20:23], v24, v25 cbsz:2 blgp:4
776
777// fp4 x fp6
778// GFX950: v_mfma_scale_f32_16x16x128_f8f6f4 v[0:3], v[4:7], v[12:17], v[20:23], v24, v25 op_sel_hi:[0,0,0] cbsz:4 blgp:2 ; encoding: [0x00,0x00,0xac,0xd3,0x18,0x33,0x02,0x00,0x00,0x0c,0xad,0xd3,0x04,0x19,0x52,0x44]
779v_mfma_scale_f32_16x16x128_f8f6f4 v[0:3], v[4:7], v[12:17], v[20:23], v24, v25 cbsz:4 blgp:2
780
781// GFX950: v_mfma_scale_f32_16x16x128_f8f6f4 a[0:3], a[4:7], a[12:17], a[20:23], v24, v25 op_sel_hi:[0,0,0] cbsz:4 blgp:2 ; encoding: [0x00,0x00,0xac,0xd3,0x18,0x33,0x02,0x00,0x00,0x8c,0xad,0xd3,0x04,0x19,0x52,0x5c]
782v_mfma_scale_f32_16x16x128_f8f6f4 a[0:3], a[4:7], a[12:17], a[20:23], v24, v25 cbsz:4 blgp:2
783
784// fp4 x fp4
785// GFX950: v_mfma_scale_f32_16x16x128_f8f6f4 v[0:3], v[4:7], v[12:15], v[20:23], v24, v25 op_sel_hi:[0,0,0] cbsz:4 blgp:4 ; encoding: [0x00,0x00,0xac,0xd3,0x18,0x33,0x02,0x00,0x00,0x0c,0xad,0xd3,0x04,0x19,0x52,0x84]
786v_mfma_scale_f32_16x16x128_f8f6f4 v[0:3], v[4:7], v[12:15], v[20:23], v24, v25 cbsz:4 blgp:4
787
788// GFX950: v_mfma_scale_f32_16x16x128_f8f6f4 a[0:3], a[4:7], a[12:15], a[20:23], v24, v25 op_sel_hi:[0,0,0] cbsz:4 blgp:4 ; encoding: [0x00,0x00,0xac,0xd3,0x18,0x33,0x02,0x00,0x00,0x8c,0xad,0xd3,0x04,0x19,0x52,0x9c]
789v_mfma_scale_f32_16x16x128_f8f6f4 a[0:3], a[4:7], a[12:15], a[20:23], v24, v25 cbsz:4 blgp:4
790
791
792//===----------------------------------------------------------------------===//
793// v_mfma_scale_f32_32x32x64_f8f6f4 corrected widths
794//===----------------------------------------------------------------------===//
795
796// fp8 x fp8
797// GFX950: v_mfma_scale_f32_32x32x64_f8f6f4 v[0:15], v[16:23], v[24:31], v[32:47], v48, v49 op_sel_hi:[0,0,0] ; encoding: [0x00,0x00,0xac,0xd3,0x30,0x63,0x02,0x00,0x00,0x08,0xae,0xd3,0x10,0x31,0x82,0x04]
798v_mfma_scale_f32_32x32x64_f8f6f4 v[0:15], v[16:23], v[24:31], v[32:47], v48, v49 cbsz:0 blgp:0
799
800// GFX950: v_mfma_scale_f32_32x32x64_f8f6f4 a[0:15], a[16:23], a[24:31], a[32:47], v48, v49 op_sel_hi:[0,0,0] ; encoding: [0x00,0x00,0xac,0xd3,0x30,0x63,0x02,0x00,0x00,0x88,0xae,0xd3,0x10,0x31,0x82,0x1c]
801v_mfma_scale_f32_32x32x64_f8f6f4 a[0:15], a[16:23], a[24:31], a[32:47], v48, v49 cbsz:0 blgp:0
802
803// bf8 x fp8
804// GFX950: v_mfma_scale_f32_32x32x64_f8f6f4 v[0:15], v[16:23], v[24:31], v[32:47], v48, v49 op_sel_hi:[0,0,0] cbsz:1 ; encoding: [0x00,0x00,0xac,0xd3,0x30,0x63,0x02,0x00,0x00,0x09,0xae,0xd3,0x10,0x31,0x82,0x04]
805v_mfma_scale_f32_32x32x64_f8f6f4 v[0:15], v[16:23], v[24:31], v[32:47], v48, v49 cbsz:1 blgp:0
806
807// GFX950: v_mfma_scale_f32_32x32x64_f8f6f4 a[0:15], a[16:23], a[24:31], a[32:47], v48, v49 op_sel_hi:[0,0,0] cbsz:1 ; encoding: [0x00,0x00,0xac,0xd3,0x30,0x63,0x02,0x00,0x00,0x89,0xae,0xd3,0x10,0x31,0x82,0x1c]
808v_mfma_scale_f32_32x32x64_f8f6f4 a[0:15], a[16:23], a[24:31], a[32:47], v48, v49 cbsz:1 blgp:0
809
810
811// fp8 x bf8
812// GFX950: v_mfma_scale_f32_32x32x64_f8f6f4 v[0:15], v[16:23], v[24:31], v[32:47], v48, v49 op_sel_hi:[0,0,0] blgp:1 ; encoding: [0x00,0x00,0xac,0xd3,0x30,0x63,0x02,0x00,0x00,0x08,0xae,0xd3,0x10,0x31,0x82,0x24]
813v_mfma_scale_f32_32x32x64_f8f6f4 v[0:15], v[16:23], v[24:31], v[32:47], v48, v49 cbsz:0 blgp:1
814
815// GFX950: v_mfma_scale_f32_32x32x64_f8f6f4 a[0:15], a[16:23], a[24:31], a[32:47], v48, v49 op_sel_hi:[0,0,0] blgp:1 ; encoding: [0x00,0x00,0xac,0xd3,0x30,0x63,0x02,0x00,0x00,0x88,0xae,0xd3,0x10,0x31,0x82,0x3c]
816v_mfma_scale_f32_32x32x64_f8f6f4 a[0:15], a[16:23], a[24:31], a[32:47], v48, v49 cbsz:0 blgp:1
817
818// bf8 x fp4
819// GFX950: v_mfma_scale_f32_32x32x64_f8f6f4 v[0:15], v[16:23], v[24:27], v[32:47], v48, v49 op_sel_hi:[0,0,0] cbsz:1 blgp:4 ; encoding: [0x00,0x00,0xac,0xd3,0x30,0x63,0x02,0x00,0x00,0x09,0xae,0xd3,0x10,0x31,0x82,0x84]
820v_mfma_scale_f32_32x32x64_f8f6f4 v[0:15], v[16:23], v[24:27], v[32:47], v48, v49 cbsz:1 blgp:4
821
822// GFX950: v_mfma_scale_f32_32x32x64_f8f6f4 a[0:15], a[16:23], a[24:27], a[32:47], v48, v49 op_sel_hi:[0,0,0] cbsz:1 blgp:4 ; encoding: [0x00,0x00,0xac,0xd3,0x30,0x63,0x02,0x00,0x00,0x89,0xae,0xd3,0x10,0x31,0x82,0x9c]
823v_mfma_scale_f32_32x32x64_f8f6f4 a[0:15], a[16:23], a[24:27], a[32:47], v48, v49 cbsz:1 blgp:4
824
825
826// fp4 x bf8
827// GFX950: v_mfma_scale_f32_32x32x64_f8f6f4 v[0:15], v[16:19], v[24:31], v[32:47], v48, v49 op_sel_hi:[0,0,0] cbsz:4 blgp:1 ; encoding: [0x00,0x00,0xac,0xd3,0x30,0x63,0x02,0x00,0x00,0x0c,0xae,0xd3,0x10,0x31,0x82,0x24]
828v_mfma_scale_f32_32x32x64_f8f6f4 v[0:15], v[16:19], v[24:31], v[32:47], v48, v49 cbsz:4 blgp:1
829
830// GFX950: v_mfma_scale_f32_32x32x64_f8f6f4 a[0:15], a[16:19], a[24:31], a[32:47], v48, v49 op_sel_hi:[0,0,0] cbsz:4 blgp:1 ; encoding: [0x00,0x00,0xac,0xd3,0x30,0x63,0x02,0x00,0x00,0x8c,0xae,0xd3,0x10,0x31,0x82,0x3c]
831v_mfma_scale_f32_32x32x64_f8f6f4 a[0:15], a[16:19], a[24:31], a[32:47], v48, v49 cbsz:4 blgp:1
832
833
834// bf6 x bf8
835// GFX950: v_mfma_scale_f32_32x32x64_f8f6f4 v[0:15], v[16:21], v[24:31], v[32:47], v48, v49 op_sel_hi:[0,0,0] cbsz:3 blgp:1 ; encoding: [0x00,0x00,0xac,0xd3,0x30,0x63,0x02,0x00,0x00,0x0b,0xae,0xd3,0x10,0x31,0x82,0x24]
836v_mfma_scale_f32_32x32x64_f8f6f4 v[0:15], v[16:21], v[24:31], v[32:47], v48, v49 cbsz:3 blgp:1
837
838// GFX950: v_mfma_scale_f32_32x32x64_f8f6f4 a[0:15], a[16:21], a[24:31], a[32:47], v48, v49 op_sel_hi:[0,0,0] cbsz:3 blgp:1 ; encoding: [0x00,0x00,0xac,0xd3,0x30,0x63,0x02,0x00,0x00,0x8b,0xae,0xd3,0x10,0x31,0x82,0x3c]
839v_mfma_scale_f32_32x32x64_f8f6f4 a[0:15], a[16:21], a[24:31], a[32:47], v48, v49 cbsz:3 blgp:1
840
841
842// bf8 x bf6
843// GFX950: v_mfma_scale_f32_32x32x64_f8f6f4 v[0:15], v[16:23], v[24:29], v[32:47], v48, v49 op_sel_hi:[0,0,0] cbsz:1 blgp:3 ; encoding: [0x00,0x00,0xac,0xd3,0x30,0x63,0x02,0x00,0x00,0x09,0xae,0xd3,0x10,0x31,0x82,0x64]
844v_mfma_scale_f32_32x32x64_f8f6f4 v[0:15], v[16:23], v[24:29], v[32:47], v48, v49 cbsz:1 blgp:3
845
846// GFX950: v_mfma_scale_f32_32x32x64_f8f6f4 a[0:15], a[16:23], a[24:29], a[32:47], v48, v49 op_sel_hi:[0,0,0] cbsz:1 blgp:3 ; encoding: [0x00,0x00,0xac,0xd3,0x30,0x63,0x02,0x00,0x00,0x89,0xae,0xd3,0x10,0x31,0x82,0x7c]
847v_mfma_scale_f32_32x32x64_f8f6f4 a[0:15], a[16:23], a[24:29], a[32:47], v48, v49 cbsz:1 blgp:3
848
849// bf6 x bf6
850// GFX950: v_mfma_scale_f32_32x32x64_f8f6f4 v[0:15], v[16:21], v[24:29], v[32:47], v48, v49 op_sel_hi:[0,0,0] cbsz:3 blgp:3 ; encoding: [0x00,0x00,0xac,0xd3,0x30,0x63,0x02,0x00,0x00,0x0b,0xae,0xd3,0x10,0x31,0x82,0x64]
851v_mfma_scale_f32_32x32x64_f8f6f4 v[0:15], v[16:21], v[24:29], v[32:47], v48, v49 cbsz:3 blgp:3
852
853// GFX950: v_mfma_scale_f32_32x32x64_f8f6f4 a[0:15], a[16:21], a[24:29], a[32:47], v48, v49 op_sel_hi:[0,0,0] cbsz:3 blgp:3 ; encoding: [0x00,0x00,0xac,0xd3,0x30,0x63,0x02,0x00,0x00,0x8b,0xae,0xd3,0x10,0x31,0x82,0x7c]
854v_mfma_scale_f32_32x32x64_f8f6f4 a[0:15], a[16:21], a[24:29], a[32:47], v48, v49 cbsz:3 blgp:3
855
856// bf6 x fp6
857// GFX950: v_mfma_scale_f32_32x32x64_f8f6f4 v[0:15], v[16:21], v[24:29], v[32:47], v48, v49 op_sel_hi:[0,0,0] cbsz:3 blgp:2 ; encoding: [0x00,0x00,0xac,0xd3,0x30,0x63,0x02,0x00,0x00,0x0b,0xae,0xd3,0x10,0x31,0x82,0x44]
858v_mfma_scale_f32_32x32x64_f8f6f4 v[0:15], v[16:21], v[24:29], v[32:47], v48, v49 cbsz:3 blgp:2
859
860// GFX950: v_mfma_scale_f32_32x32x64_f8f6f4 a[0:15], a[16:21], a[24:29], a[32:47], v48, v49 op_sel_hi:[0,0,0] cbsz:3 blgp:2 ; encoding: [0x00,0x00,0xac,0xd3,0x30,0x63,0x02,0x00,0x00,0x8b,0xae,0xd3,0x10,0x31,0x82,0x5c]
861v_mfma_scale_f32_32x32x64_f8f6f4 a[0:15], a[16:21], a[24:29], a[32:47], v48, v49 cbsz:3 blgp:2
862
863
864// fp6 x bf6
865// GFX950: v_mfma_scale_f32_32x32x64_f8f6f4 v[0:15], v[16:21], v[24:29], v[32:47], v48, v49 op_sel_hi:[0,0,0] cbsz:2 blgp:3 ; encoding: [0x00,0x00,0xac,0xd3,0x30,0x63,0x02,0x00,0x00,0x0a,0xae,0xd3,0x10,0x31,0x82,0x64]
866v_mfma_scale_f32_32x32x64_f8f6f4 v[0:15], v[16:21], v[24:29], v[32:47], v48, v49 cbsz:2 blgp:3
867
868// GFX950: v_mfma_scale_f32_32x32x64_f8f6f4 a[0:15], a[16:21], a[24:29], a[32:47], v48, v49 op_sel_hi:[0,0,0] cbsz:2 blgp:3 ; encoding: [0x00,0x00,0xac,0xd3,0x30,0x63,0x02,0x00,0x00,0x8a,0xae,0xd3,0x10,0x31,0x82,0x7c]
869v_mfma_scale_f32_32x32x64_f8f6f4 a[0:15], a[16:21], a[24:29], a[32:47], v48, v49 cbsz:2 blgp:3
870
871
872// fp6 x fp4
873// GFX950: v_mfma_scale_f32_32x32x64_f8f6f4 v[0:15], v[16:21], v[24:27], v[32:47], v48, v49 op_sel_hi:[0,0,0] cbsz:2 blgp:4 ; encoding: [0x00,0x00,0xac,0xd3,0x30,0x63,0x02,0x00,0x00,0x0a,0xae,0xd3,0x10,0x31,0x82,0x84]
874v_mfma_scale_f32_32x32x64_f8f6f4 v[0:15], v[16:21], v[24:27], v[32:47], v48, v49 cbsz:2 blgp:4
875
876// GFX950: v_mfma_scale_f32_32x32x64_f8f6f4 a[0:15], a[16:21], a[24:27], a[32:47], v48, v49 op_sel_hi:[0,0,0] cbsz:2 blgp:4 ; encoding: [0x00,0x00,0xac,0xd3,0x30,0x63,0x02,0x00,0x00,0x8a,0xae,0xd3,0x10,0x31,0x82,0x9c]
877v_mfma_scale_f32_32x32x64_f8f6f4 a[0:15], a[16:21], a[24:27], a[32:47], v48, v49 cbsz:2 blgp:4
878
879
880// fp4 x fp6
881// GFX950: v_mfma_scale_f32_32x32x64_f8f6f4 v[0:15], v[16:19], v[24:29], v[32:47], v48, v49 op_sel_hi:[0,0,0] cbsz:4 blgp:2 ; encoding: [0x00,0x00,0xac,0xd3,0x30,0x63,0x02,0x00,0x00,0x0c,0xae,0xd3,0x10,0x31,0x82,0x44]
882v_mfma_scale_f32_32x32x64_f8f6f4 v[0:15], v[16:19], v[24:29], v[32:47], v48, v49 cbsz:4 blgp:2
883
884// GFX950: v_mfma_scale_f32_32x32x64_f8f6f4 a[0:15], a[16:19], a[24:29], a[32:47], v48, v49 op_sel_hi:[0,0,0] cbsz:4 blgp:2 ; encoding: [0x00,0x00,0xac,0xd3,0x30,0x63,0x02,0x00,0x00,0x8c,0xae,0xd3,0x10,0x31,0x82,0x5c]
885v_mfma_scale_f32_32x32x64_f8f6f4 a[0:15], a[16:19], a[24:29], a[32:47], v48, v49 cbsz:4 blgp:2
886
887
888// fp4 x fp4
889// GFX950: v_mfma_scale_f32_32x32x64_f8f6f4 v[0:15], v[16:19], v[24:27], v[32:47], v48, v49 op_sel_hi:[0,0,0] cbsz:4 blgp:4 ; encoding: [0x00,0x00,0xac,0xd3,0x30,0x63,0x02,0x00,0x00,0x0c,0xae,0xd3,0x10,0x31,0x82,0x84]
890v_mfma_scale_f32_32x32x64_f8f6f4 v[0:15], v[16:19], v[24:27], v[32:47], v48, v49 cbsz:4 blgp:4
891
892// GFX950: v_mfma_scale_f32_32x32x64_f8f6f4 a[0:15], a[16:19], a[24:27], a[32:47], v48, v49 op_sel_hi:[0,0,0] cbsz:4 blgp:4 ; encoding: [0x00,0x00,0xac,0xd3,0x30,0x63,0x02,0x00,0x00,0x8c,0xae,0xd3,0x10,0x31,0x82,0x9c]
893v_mfma_scale_f32_32x32x64_f8f6f4 a[0:15], a[16:19], a[24:27], a[32:47], v48, v49 cbsz:4 blgp:4
894
895//===----------------------------------------------------------------------===//
896// v_mfma_i32_16x16x64_i8
897//===----------------------------------------------------------------------===//
898
899// GFX950: v_mfma_i32_16x16x64_i8 v[0:3], v[0:3], v[0:3], v[0:3] ; encoding: [0x00,0x00,0xb6,0xd3,0x00,0x01,0x02,0x04]
900// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU
901v_mfma_i32_16x16x64_i8 v[0:3], v[0:3], v[0:3], v[0:3]
902
903// GFX950: v_mfma_i32_16x16x64_i8 v[0:3], v[0:3], v[0:3], v[0:3] ; encoding: [0x00,0x00,0xb6,0xd3,0x00,0x01,0x02,0x04]
904// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU
905v_mfma_i32_16x16x64i8 v[0:3], v[0:3], v[0:3], v[0:3]
906
907// GFX950: v_mfma_i32_16x16x64_i8 a[0:3], a[0:3], a[0:3], a[0:3] ; encoding: [0x00,0x80,0xb6,0xd3,0x00,0x01,0x02,0x1c]
908// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU
909v_mfma_i32_16x16x64_i8 a[0:3], a[0:3], a[0:3], a[0:3]
910
911// GFX950: v_mfma_i32_16x16x64_i8 a[0:3], a[0:3], a[0:3], a[0:3] ; encoding: [0x00,0x80,0xb6,0xd3,0x00,0x01,0x02,0x1c]
912// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU
913v_mfma_i32_16x16x64i8 a[0:3], a[0:3], a[0:3], a[0:3]
914
915// GFX950: v_mfma_i32_16x16x64_i8 v[0:3], a[0:3], v[0:3], 1.0 ; encoding: [0x00,0x00,0xb6,0xd3,0x00,0x01,0xca,0x0b]
916// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU
917v_mfma_i32_16x16x64_i8 v[0:3], a[0:3], v[0:3], 1.0
918
919// GFX950: v_mfma_i32_16x16x64_i8 a[0:3], v[0:3], a[0:3], 1.0 ; encoding: [0x00,0x80,0xb6,0xd3,0x00,0x01,0xca,0x13]
920// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU
921v_mfma_i32_16x16x64_i8 a[0:3], v[0:3], a[0:3], 1.0
922
923// GFX950: v_mfma_i32_16x16x64_i8 v[0:3], v[0:3], v[0:3], v[0:3] blgp:5 ; encoding: [0x00,0x00,0xb6,0xd3,0x00,0x01,0x02,0xa4]
924// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU
925v_mfma_i32_16x16x64_i8 v[0:3], v[0:3], v[0:3], v[0:3] blgp:5
926
927// GFX950: v_mfma_i32_16x16x64_i8 a[0:3], a[0:3], a[0:3], a[0:3] blgp:1 ; encoding: [0x00,0x80,0xb6,0xd3,0x00,0x01,0x02,0x3c]
928// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU
929v_mfma_i32_16x16x64_i8 a[0:3], a[0:3], a[0:3], a[0:3] blgp:1
930
931// GFX950: v_mfma_i32_16x16x64_i8 a[0:3], a[0:3], a[0:3], a[0:3] cbsz:3 ; encoding: [0x00,0x83,0xb6,0xd3,0x00,0x01,0x02,0x1c]
932// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU
933v_mfma_i32_16x16x64_i8 a[0:3], a[0:3], a[0:3], a[0:3] cbsz:3
934
935// GFX950:  v_mfma_i32_16x16x64_i8 a[0:3], a[0:3], a[0:3], a[0:3] abid:1 ; encoding: [0x00,0x88,0xb6,0xd3,0x00,0x01,0x02,0x1c]
936// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU
937v_mfma_i32_16x16x64_i8 a[0:3], a[0:3], a[0:3], a[0:3] abid:1
938
939// GFX950:  v_mfma_i32_16x16x64_i8 a[0:3], a[0:3], a[0:3], a[0:3] cbsz:3 abid:1 ; encoding: [0x00,0x8b,0xb6,0xd3,0x00,0x01,0x02,0x1c]
940// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU
941v_mfma_i32_16x16x64_i8 a[0:3], a[0:3], a[0:3], a[0:3] cbsz:3 abid:1
942
943// GFX950: v_mfma_i32_16x16x64_i8 a[0:3], v[0:3], v[0:3], a[4:7] ; encoding: [0x00,0x80,0xb6,0xd3,0x00,0x01,0x12,0x04]
944// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU
945v_mfma_i32_16x16x64_i8 a[0:3], v[0:3], v[0:3], a[4:7]
946
947// GFX950: v_mfma_i32_16x16x64_i8 v[0:3], a[0:3], a[0:3], v[4:7] ; encoding: [0x00,0x00,0xb6,0xd3,0x00,0x01,0x12,0x1c]
948// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU
949v_mfma_i32_16x16x64_i8 v[0:3], a[0:3], a[0:3], v[4:7]
950
951//===----------------------------------------------------------------------===//
952// v_mfma_i32_32x32x32_i8
953//===----------------------------------------------------------------------===//
954
955// GFX950: v_mfma_i32_32x32x32_i8 v[0:15], v[0:3], v[0:3], v[0:15] ; encoding: [0x00,0x00,0xb8,0xd3,0x00,0x01,0x02,0x04]
956// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU
957v_mfma_i32_32x32x32_i8 v[0:15], v[0:3], v[0:3], v[0:15]
958
959// GFX950: v_mfma_i32_32x32x32_i8 a[0:15], a[0:3], a[0:3], a[0:15] ; encoding: [0x00,0x80,0xb8,0xd3,0x00,0x01,0x02,0x1c]
960// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU
961v_mfma_i32_32x32x32_i8 a[0:15], a[0:3], a[0:3], a[0:15]
962
963// GFX950:  v_mfma_i32_32x32x32_i8 v[0:15], v[0:3], v[0:3], v[0:15] ; encoding: [0x00,0x00,0xb8,0xd3,0x00,0x01,0x02,0x04]
964// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU
965v_mfma_i32_32x32x32i8 v[0:15], v[0:3], v[0:3], v[0:15]
966
967// GFX950: v_mfma_i32_32x32x32_i8 a[0:15], a[0:3], a[0:3], a[0:15] ; encoding: [0x00,0x80,0xb8,0xd3,0x00,0x01,0x02,0x1c]
968// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU
969v_mfma_i32_32x32x32i8 a[0:15], a[0:3], a[0:3], a[0:15]
970
971// GFX950: v_mfma_i32_32x32x32_i8 v[0:15], v[0:3], v[0:3], 1.0 ; encoding: [0x00,0x00,0xb8,0xd3,0x00,0x01,0xca,0x03]
972// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU
973v_mfma_i32_32x32x32_i8 v[0:15], v[0:3], v[0:3], 1.0
974
975// GFX950: v_mfma_i32_32x32x32_i8 a[0:15], a[0:3], a[0:3], 1.0 ; encoding: [0x00,0x80,0xb8,0xd3,0x00,0x01,0xca,0x1b]
976// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU
977v_mfma_i32_32x32x32_i8 a[0:15], a[0:3], a[0:3], 1.0
978
979// GFX950:  v_mfma_i32_32x32x32_i8 v[0:15], v[0:3], v[0:3], v[0:15] blgp:5 ; encoding: [0x00,0x00,0xb8,0xd3,0x00,0x01,0x02,0xa4]
980// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU
981v_mfma_i32_32x32x32_i8 v[0:15], v[0:3], v[0:3], v[0:15] blgp:5
982
983// GFX950: v_mfma_i32_32x32x32_i8 a[0:15], a[0:3], a[0:3], a[0:15] blgp:2 ; encoding: [0x00,0x80,0xb8,0xd3,0x00,0x01,0x02,0x5c]
984// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU
985v_mfma_i32_32x32x32_i8 a[0:15], a[0:3], a[0:3], a[0:15] blgp:2
986
987// GFX950: v_mfma_i32_32x32x32_i8 v[0:15], v[0:3], v[0:3], v[0:15] cbsz:3 ; encoding: [0x00,0x03,0xb8,0xd3,0x00,0x01,0x02,0x04]
988// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU
989v_mfma_i32_32x32x32_i8 v[0:15], v[0:3], v[0:3], v[0:15] cbsz:3
990
991// GFX950: v_mfma_i32_32x32x32_i8 v[0:15], v[0:3], v[0:3], v[0:15] abid:1 ; encoding: [0x00,0x08,0xb8,0xd3,0x00,0x01,0x02,0x04]
992// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU
993v_mfma_i32_32x32x32_i8 v[0:15], v[0:3], v[0:3], v[0:15] abid:1
994
995// GFX950: v_mfma_i32_32x32x32_i8 a[0:15], a[0:3], a[0:3], a[0:15] cbsz:3 abid:1 ; encoding: [0x00,0x8b,0xb8,0xd3,0x00,0x01,0x02,0x1c]
996// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU
997v_mfma_i32_32x32x32_i8 a[0:15], a[0:3], a[0:3], a[0:15] cbsz:3 abid:1
998
999//===----------------------------------------------------------------------===//
1000// v_mfma_f32_16x16x32_bf16
1001//===----------------------------------------------------------------------===//
1002
1003// GFX950: v_mfma_f32_16x16x32_bf16 v[0:3], v[0:3], v[0:3], v[0:3] ; encoding: [0x00,0x00,0xb5,0xd3,0x00,0x01,0x02,0x04]
1004// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1005v_mfma_f32_16x16x32_bf16 v[0:3], v[0:3], v[0:3], v[0:3]
1006
1007// GFX950: v_mfma_f32_16x16x32_bf16 v[0:3], v[0:3], v[0:3], v[0:3] ; encoding: [0x00,0x00,0xb5,0xd3,0x00,0x01,0x02,0x04]
1008// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1009v_mfma_f32_16x16x32bf16 v[0:3], v[0:3], v[0:3], v[0:3]
1010
1011// GFX950: v_mfma_f32_16x16x32_bf16 a[0:3], a[0:3], a[0:3], a[0:3] ; encoding: [0x00,0x80,0xb5,0xd3,0x00,0x01,0x02,0x1c]
1012// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1013v_mfma_f32_16x16x32_bf16 a[0:3], a[0:3], a[0:3], a[0:3]
1014
1015// GFX950: v_mfma_f32_16x16x32_bf16 a[0:3], a[0:3], a[0:3], a[0:3] ; encoding: [0x00,0x80,0xb5,0xd3,0x00,0x01,0x02,0x1c]
1016// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1017v_mfma_f32_16x16x32bf16 a[0:3], a[0:3], a[0:3], a[0:3]
1018
1019// GFX950: v_mfma_f32_16x16x32_bf16 v[0:3], a[0:3], v[0:3], 1.0 ; encoding: [0x00,0x00,0xb5,0xd3,0x00,0x01,0xca,0x0b]
1020// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1021v_mfma_f32_16x16x32_bf16 v[0:3], a[0:3], v[0:3], 1.0
1022
1023// GFX950: v_mfma_f32_16x16x32_bf16 a[0:3], v[0:3], a[0:3], 1.0 ; encoding: [0x00,0x80,0xb5,0xd3,0x00,0x01,0xca,0x13]
1024// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1025v_mfma_f32_16x16x32_bf16 a[0:3], v[0:3], a[0:3], 1.0
1026
1027// GFX950: v_mfma_f32_16x16x32_bf16 v[0:3], v[0:3], v[0:3], v[0:3] blgp:5 ; encoding: [0x00,0x00,0xb5,0xd3,0x00,0x01,0x02,0xa4]
1028// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1029v_mfma_f32_16x16x32_bf16 v[0:3], v[0:3], v[0:3], v[0:3] blgp:5
1030
1031// GFX950: v_mfma_f32_16x16x32_bf16 a[0:3], a[0:3], a[0:3], a[0:3] blgp:1 ; encoding: [0x00,0x80,0xb5,0xd3,0x00,0x01,0x02,0x3c]
1032// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1033v_mfma_f32_16x16x32_bf16 a[0:3], a[0:3], a[0:3], a[0:3] blgp:1
1034
1035// GFX950: v_mfma_f32_16x16x32_bf16 a[0:3], a[0:3], a[0:3], a[0:3] cbsz:3 ; encoding: [0x00,0x83,0xb5,0xd3,0x00,0x01,0x02,0x1c]
1036// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1037v_mfma_f32_16x16x32_bf16 a[0:3], a[0:3], a[0:3], a[0:3] cbsz:3
1038
1039// GFX950:  v_mfma_f32_16x16x32_bf16 a[0:3], a[0:3], a[0:3], a[0:3] abid:1 ; encoding: [0x00,0x88,0xb5,0xd3,0x00,0x01,0x02,0x1c]
1040// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1041v_mfma_f32_16x16x32_bf16 a[0:3], a[0:3], a[0:3], a[0:3] abid:1
1042
1043// GFX950:  v_mfma_f32_16x16x32_bf16 a[0:3], a[0:3], a[0:3], a[0:3] cbsz:3 abid:1 ; encoding: [0x00,0x8b,0xb5,0xd3,0x00,0x01,0x02,0x1c]
1044// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1045v_mfma_f32_16x16x32_bf16 a[0:3], a[0:3], a[0:3], a[0:3] cbsz:3 abid:1
1046
1047// GFX950: v_mfma_f32_16x16x32_bf16 a[0:3], v[0:3], v[0:3], a[4:7] ; encoding: [0x00,0x80,0xb5,0xd3,0x00,0x01,0x12,0x04]
1048// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1049v_mfma_f32_16x16x32_bf16 a[0:3], v[0:3], v[0:3], a[4:7]
1050
1051// GFX950: v_mfma_f32_16x16x32_bf16 v[0:3], a[0:3], a[0:3], v[4:7] ; encoding: [0x00,0x00,0xb5,0xd3,0x00,0x01,0x12,0x1c]
1052// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1053v_mfma_f32_16x16x32_bf16 v[0:3], a[0:3], a[0:3], v[4:7]
1054
1055
1056//===----------------------------------------------------------------------===//
1057// SMFMAC opcodes.
1058//===----------------------------------------------------------------------===//
1059
1060//===----------------------------------------------------------------------===//
1061// v_smfmac_f32_16x16x64_f16
1062//===----------------------------------------------------------------------===//
1063
1064// GFX950: v_smfmac_f32_16x16x64_f16 v[10:13], a[2:5], v[4:11], v3 cbsz:3 abid:1 ; encoding: [0x0a,0x0b,0xda,0xd3,0x02,0x09,0x0e,0x0c]
1065// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1066v_smfmac_f32_16x16x64_f16 v[10:13], a[2:5], v[4:11], v3 cbsz:3 abid:1
1067
1068// GFX950: v_smfmac_f32_16x16x64_f16 v[10:13], a[2:5], v[4:11], v3 cbsz:3 abid:1 ; encoding: [0x0a,0x0b,0xda,0xd3,0x02,0x09,0x0e,0x0c]
1069// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1070v_smfmac_f32_16x16x64f16 v[10:13], a[2:5], v[4:11], v3 cbsz:3 abid:1
1071
1072// GFX950: v_smfmac_f32_16x16x64_f16 a[10:13], v[2:5], a[4:11], v1 ; encoding: [0x0a,0x80,0xda,0xd3,0x02,0x09,0x06,0x14]
1073// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1074v_smfmac_f32_16x16x64_f16 a[10:13], v[2:5], a[4:11], v1
1075
1076// GFX950: v_smfmac_f32_16x16x64_f16 v[10:13], a[2:5], v[4:11], v2 cbsz:3 abid:1 ; encoding: [0x0a,0x0b,0xda,0xd3,0x02,0x09,0x0a,0x0c]
1077// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1078v_smfmac_f32_16x16x64_f16 v[10:13], a[2:5], v[4:11], v2 cbsz:3 abid:1
1079
1080// GFX950: v_smfmac_f32_16x16x64_f16 a[10:13], v[2:5], a[4:11], v3 ; encoding: [0x0a,0x80,0xda,0xd3,0x02,0x09,0x0e,0x14]
1081// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1082v_smfmac_f32_16x16x64_f16 a[10:13], v[2:5], a[4:11], v3
1083
1084// GFX950: v_smfmac_f32_16x16x64_f16 v[10:13], v[2:5], v[6:13], v2 cbsz:3 abid:1 ; encoding: [0x0a,0x0b,0xda,0xd3,0x02,0x0d,0x0a,0x04]
1085// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1086v_smfmac_f32_16x16x64_f16 v[10:13], v[2:5], v[6:13], v2 cbsz:3 abid:1
1087
1088// GFX950: v_smfmac_f32_16x16x64_f16 a[10:13], a[2:5], a[6:13], v2 cbsz:3 abid:1 ; encoding: [0x0a,0x8b,0xda,0xd3,0x02,0x0d,0x0a,0x1c]
1089// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1090v_smfmac_f32_16x16x64_f16 a[10:13], a[2:5], a[6:13], v2 cbsz:3 abid:1
1091
1092// GFX950: v_smfmac_f32_16x16x64_f16 a[10:13], a[2:5], a[6:13], v3 cbsz:1 abid:3 ; encoding: [0x0a,0x99,0xda,0xd3,0x02,0x0d,0x0e,0x1c]
1093// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1094v_smfmac_f32_16x16x64_f16 a[10:13], a[2:5], a[6:13], v3 cbsz:1 abid:3
1095
1096//===----------------------------------------------------------------------===//
1097// v_smfmac_f32_32x32x32_f16
1098//===----------------------------------------------------------------------===//
1099
1100// GFX950: v_smfmac_f32_32x32x32_f16 v[10:25], a[2:5], v[4:11], v3 cbsz:3 abid:1 ; encoding: [0x0a,0x0b,0xdb,0xd3,0x02,0x09,0x0e,0x0c]
1101// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1102v_smfmac_f32_32x32x32_f16 v[10:25], a[2:5], v[4:11], v3 cbsz:3 abid:1
1103
1104// GFX950: v_smfmac_f32_32x32x32_f16 v[10:25], a[2:5], v[4:11], v3 cbsz:3 abid:1 ; encoding: [0x0a,0x0b,0xdb,0xd3,0x02,0x09,0x0e,0x0c]
1105// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1106v_smfmac_f32_32x32x32f16 v[10:25], a[2:5], v[4:11], v3 cbsz:3 abid:1
1107
1108// GFX950: v_smfmac_f32_32x32x32_f16 a[10:25], v[2:5], a[4:11], v1 ; encoding: [0x0a,0x80,0xdb,0xd3,0x02,0x09,0x06,0x14]
1109// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1110v_smfmac_f32_32x32x32_f16 a[10:25], v[2:5], a[4:11], v1
1111
1112// GFX950: v_smfmac_f32_32x32x32_f16 v[10:25], a[2:5], v[4:11], v2 cbsz:3 abid:1 ; encoding: [0x0a,0x0b,0xdb,0xd3,0x02,0x09,0x0a,0x0c]
1113// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1114v_smfmac_f32_32x32x32_f16 v[10:25], a[2:5], v[4:11], v2 cbsz:3 abid:1
1115
1116// GFX950: v_smfmac_f32_32x32x32_f16 a[10:25], v[2:5], a[4:11], v3 ; encoding: [0x0a,0x80,0xdb,0xd3,0x02,0x09,0x0e,0x14]
1117// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1118v_smfmac_f32_32x32x32_f16 a[10:25], v[2:5], a[4:11], v3
1119
1120// GFX950: v_smfmac_f32_32x32x32_f16 v[10:25], v[2:5], v[6:13], v2 cbsz:3 abid:1 ; encoding: [0x0a,0x0b,0xdb,0xd3,0x02,0x0d,0x0a,0x04]
1121// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1122v_smfmac_f32_32x32x32_f16 v[10:25], v[2:5], v[6:13], v2 cbsz:3 abid:1
1123
1124// GFX950: v_smfmac_f32_32x32x32_f16 a[10:25], a[2:5], a[6:13], v2 cbsz:3 abid:1 ; encoding: [0x0a,0x8b,0xdb,0xd3,0x02,0x0d,0x0a,0x1c]
1125// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1126v_smfmac_f32_32x32x32_f16 a[10:25], a[2:5], a[6:13], v2 cbsz:3 abid:1
1127
1128// GFX950: v_smfmac_f32_32x32x32_f16 a[10:25], a[2:5], a[6:13], v3 cbsz:1 abid:3 ; encoding: [0x0a,0x99,0xdb,0xd3,0x02,0x0d,0x0e,0x1c]
1129// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1130v_smfmac_f32_32x32x32_f16 a[10:25], a[2:5], a[6:13], v3 cbsz:1 abid:3
1131
1132//===----------------------------------------------------------------------===//
1133// v_smfmac_f32_16x16x64_bf16
1134//===----------------------------------------------------------------------===//
1135
1136// GFX950: v_smfmac_f32_16x16x64_bf16 v[10:13], a[2:5], v[4:11], v3 cbsz:3 abid:1 ; encoding: [0x0a,0x0b,0xb9,0xd3,0x02,0x09,0x0e,0x0c]
1137// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1138v_smfmac_f32_16x16x64_bf16 v[10:13], a[2:5], v[4:11], v3 cbsz:3 abid:1
1139
1140// GFX950: v_smfmac_f32_16x16x64_bf16 v[10:13], a[2:5], v[4:11], v3 cbsz:3 abid:1 ; encoding: [0x0a,0x0b,0xb9,0xd3,0x02,0x09,0x0e,0x0c]
1141// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1142v_smfmac_f32_16x16x64bf16 v[10:13], a[2:5], v[4:11], v3 cbsz:3 abid:1
1143
1144// GFX950: v_smfmac_f32_16x16x64_bf16 a[10:13], v[2:5], a[4:11], v1 ; encoding: [0x0a,0x80,0xb9,0xd3,0x02,0x09,0x06,0x14]
1145// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1146v_smfmac_f32_16x16x64_bf16 a[10:13], v[2:5], a[4:11], v1
1147
1148// GFX950: v_smfmac_f32_16x16x64_bf16 v[10:13], a[2:5], v[4:11], v2 cbsz:3 abid:1 ; encoding: [0x0a,0x0b,0xb9,0xd3,0x02,0x09,0x0a,0x0c]
1149// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1150v_smfmac_f32_16x16x64_bf16 v[10:13], a[2:5], v[4:11], v2 cbsz:3 abid:1
1151
1152// GFX950: v_smfmac_f32_16x16x64_bf16 a[10:13], v[2:5], a[4:11], v3 ; encoding: [0x0a,0x80,0xb9,0xd3,0x02,0x09,0x0e,0x14]
1153// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1154v_smfmac_f32_16x16x64_bf16 a[10:13], v[2:5], a[4:11], v3
1155
1156// GFX950: v_smfmac_f32_16x16x64_bf16 v[10:13], v[2:5], v[6:13], v2 cbsz:3 abid:1 ; encoding: [0x0a,0x0b,0xb9,0xd3,0x02,0x0d,0x0a,0x04]
1157// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1158v_smfmac_f32_16x16x64_bf16 v[10:13], v[2:5], v[6:13], v2 cbsz:3 abid:1
1159
1160// GFX950: v_smfmac_f32_16x16x64_bf16 a[10:13], a[2:5], a[6:13], v2 cbsz:3 abid:1 ; encoding: [0x0a,0x8b,0xb9,0xd3,0x02,0x0d,0x0a,0x1c]
1161// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1162v_smfmac_f32_16x16x64_bf16 a[10:13], a[2:5], a[6:13], v2 cbsz:3 abid:1
1163
1164// GFX950: v_smfmac_f32_16x16x64_bf16 a[10:13], a[2:5], a[6:13], v3 cbsz:1 abid:3 ; encoding: [0x0a,0x99,0xb9,0xd3,0x02,0x0d,0x0e,0x1c]
1165// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1166v_smfmac_f32_16x16x64_bf16 a[10:13], a[2:5], a[6:13], v3 cbsz:1 abid:3
1167
1168//===----------------------------------------------------------------------===//
1169// v_smfmac_f32_32x32x32_bf16
1170//===----------------------------------------------------------------------===//
1171
1172// GFX950: v_smfmac_f32_32x32x32_bf16 v[10:25], a[2:5], v[4:11], v3 cbsz:3 abid:1 ; encoding: [0x0a,0x0b,0xc6,0xd3,0x02,0x09,0x0e,0x0c]
1173// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1174v_smfmac_f32_32x32x32_bf16 v[10:25], a[2:5], v[4:11], v3 cbsz:3 abid:1
1175
1176// GFX950: v_smfmac_f32_32x32x32_bf16 v[10:25], a[2:5], v[4:11], v3 cbsz:3 abid:1 ; encoding: [0x0a,0x0b,0xc6,0xd3,0x02,0x09,0x0e,0x0c]
1177// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1178v_smfmac_f32_32x32x32bf16 v[10:25], a[2:5], v[4:11], v3 cbsz:3 abid:1
1179
1180// GFX950: v_smfmac_f32_32x32x32_bf16 a[10:25], v[2:5], a[4:11], v1 ; encoding: [0x0a,0x80,0xc6,0xd3,0x02,0x09,0x06,0x14]
1181// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1182v_smfmac_f32_32x32x32_bf16 a[10:25], v[2:5], a[4:11], v1
1183
1184// GFX950: v_smfmac_f32_32x32x32_bf16 v[10:25], a[2:5], v[4:11], v2 cbsz:3 abid:1 ; encoding: [0x0a,0x0b,0xc6,0xd3,0x02,0x09,0x0a,0x0c]
1185// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1186v_smfmac_f32_32x32x32_bf16 v[10:25], a[2:5], v[4:11], v2 cbsz:3 abid:1
1187
1188// GFX950: v_smfmac_f32_32x32x32_bf16 a[10:25], v[2:5], a[4:11], v3 ; encoding: [0x0a,0x80,0xc6,0xd3,0x02,0x09,0x0e,0x14]
1189// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1190v_smfmac_f32_32x32x32_bf16 a[10:25], v[2:5], a[4:11], v3
1191
1192// GFX950: v_smfmac_f32_32x32x32_bf16 v[10:25], v[2:5], v[6:13], v2 cbsz:3 abid:1 ; encoding: [0x0a,0x0b,0xc6,0xd3,0x02,0x0d,0x0a,0x04]
1193// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1194v_smfmac_f32_32x32x32_bf16 v[10:25], v[2:5], v[6:13], v2 cbsz:3 abid:1
1195
1196// GFX950: v_smfmac_f32_32x32x32_bf16 a[10:25], a[2:5], a[6:13], v2 cbsz:3 abid:1 ; encoding: [0x0a,0x8b,0xc6,0xd3,0x02,0x0d,0x0a,0x1c]
1197// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1198v_smfmac_f32_32x32x32_bf16 a[10:25], a[2:5], a[6:13], v2 cbsz:3 abid:1
1199
1200// GFX950: v_smfmac_f32_32x32x32_bf16 a[10:25], a[2:5], a[6:13], v3 cbsz:1 abid:3 ; encoding: [0x0a,0x99,0xc6,0xd3,0x02,0x0d,0x0e,0x1c]
1201// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1202v_smfmac_f32_32x32x32_bf16 a[10:25], a[2:5], a[6:13], v3 cbsz:1 abid:3
1203
1204//===----------------------------------------------------------------------===//
1205// v_smfmac_i32_16x16x128_i8
1206//===----------------------------------------------------------------------===//
1207
1208// GFX950: v_smfmac_i32_16x16x128_i8 v[10:13], a[2:5], v[4:11], v3 cbsz:3 abid:1 ; encoding: [0x0a,0x0b,0xba,0xd3,0x02,0x09,0x0e,0x0c]
1209// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1210v_smfmac_i32_16x16x128_i8 v[10:13], a[2:5], v[4:11], v3 cbsz:3 abid:1
1211
1212// GFX950: v_smfmac_i32_16x16x128_i8 v[10:13], a[2:5], v[4:11], v3 cbsz:3 abid:1 ; encoding: [0x0a,0x0b,0xba,0xd3,0x02,0x09,0x0e,0x0c]
1213// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1214v_smfmac_i32_16x16x128i8 v[10:13], a[2:5], v[4:11], v3 cbsz:3 abid:1
1215
1216// GFX950: v_smfmac_i32_16x16x128_i8 a[10:13], v[2:5], a[4:11], v1 ; encoding: [0x0a,0x80,0xba,0xd3,0x02,0x09,0x06,0x14]
1217// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1218v_smfmac_i32_16x16x128_i8 a[10:13], v[2:5], a[4:11], v1
1219
1220// GFX950: v_smfmac_i32_16x16x128_i8 v[10:13], a[2:5], v[4:11], v2 cbsz:3 abid:1 ; encoding: [0x0a,0x0b,0xba,0xd3,0x02,0x09,0x0a,0x0c]
1221// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1222v_smfmac_i32_16x16x128_i8 v[10:13], a[2:5], v[4:11], v2 cbsz:3 abid:1
1223
1224// GFX950: v_smfmac_i32_16x16x128_i8 a[10:13], v[2:5], a[4:11], v3 ; encoding: [0x0a,0x80,0xba,0xd3,0x02,0x09,0x0e,0x14]
1225// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1226v_smfmac_i32_16x16x128_i8 a[10:13], v[2:5], a[4:11], v3
1227
1228// GFX950: v_smfmac_i32_16x16x128_i8 v[10:13], v[2:5], v[6:13], v2 cbsz:3 abid:1 ; encoding: [0x0a,0x0b,0xba,0xd3,0x02,0x0d,0x0a,0x04]
1229// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1230v_smfmac_i32_16x16x128_i8 v[10:13], v[2:5], v[6:13], v2 cbsz:3 abid:1
1231
1232// GFX950: v_smfmac_i32_16x16x128_i8 a[10:13], a[2:5], a[6:13], v2 cbsz:3 abid:1 ; encoding: [0x0a,0x8b,0xba,0xd3,0x02,0x0d,0x0a,0x1c]
1233// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1234v_smfmac_i32_16x16x128_i8 a[10:13], a[2:5], a[6:13], v2 cbsz:3 abid:1
1235
1236// GFX950: v_smfmac_i32_16x16x128_i8 a[10:13], a[2:5], a[6:13], v3 cbsz:1 abid:3 ; encoding: [0x0a,0x99,0xba,0xd3,0x02,0x0d,0x0e,0x1c]
1237// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1238v_smfmac_i32_16x16x128_i8 a[10:13], a[2:5], a[6:13], v3 cbsz:1 abid:3
1239
1240//===----------------------------------------------------------------------===//
1241// v_smfmac_i32_32x32x64_i8
1242//===----------------------------------------------------------------------===//
1243
1244// GFX950: v_smfmac_i32_32x32x64_i8 v[10:25], a[2:5], v[4:11], v3 cbsz:3 abid:1 ; encoding: [0x0a,0x0b,0xc7,0xd3,0x02,0x09,0x0e,0x0c]
1245// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1246v_smfmac_i32_32x32x64_i8 v[10:25], a[2:5], v[4:11], v3 cbsz:3 abid:1
1247
1248// GFX950: v_smfmac_i32_32x32x64_i8 v[10:25], a[2:5], v[4:11], v3 cbsz:3 abid:1 ; encoding: [0x0a,0x0b,0xc7,0xd3,0x02,0x09,0x0e,0x0c]
1249// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1250v_smfmac_i32_32x32x64i8 v[10:25], a[2:5], v[4:11], v3 cbsz:3 abid:1
1251
1252// GFX950: v_smfmac_i32_32x32x64_i8 a[10:25], v[2:5], a[4:11], v1 ; encoding: [0x0a,0x80,0xc7,0xd3,0x02,0x09,0x06,0x14]
1253// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1254v_smfmac_i32_32x32x64_i8 a[10:25], v[2:5], a[4:11], v1
1255
1256// GFX950: v_smfmac_i32_32x32x64_i8 v[10:25], a[2:5], v[4:11], v2 cbsz:3 abid:1 ; encoding: [0x0a,0x0b,0xc7,0xd3,0x02,0x09,0x0a,0x0c]
1257// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1258v_smfmac_i32_32x32x64_i8 v[10:25], a[2:5], v[4:11], v2 cbsz:3 abid:1
1259
1260// GFX950: v_smfmac_i32_32x32x64_i8 a[10:25], v[2:5], a[4:11], v3 ; encoding: [0x0a,0x80,0xc7,0xd3,0x02,0x09,0x0e,0x14]
1261// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1262v_smfmac_i32_32x32x64_i8 a[10:25], v[2:5], a[4:11], v3
1263
1264// GFX950: v_smfmac_i32_32x32x64_i8 v[10:25], v[2:5], v[6:13], v2 cbsz:3 abid:1 ; encoding: [0x0a,0x0b,0xc7,0xd3,0x02,0x0d,0x0a,0x04]
1265// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1266v_smfmac_i32_32x32x64_i8 v[10:25], v[2:5], v[6:13], v2 cbsz:3 abid:1
1267
1268// GFX950: v_smfmac_i32_32x32x64_i8 a[10:25], a[2:5], a[6:13], v2 cbsz:3 abid:1 ; encoding: [0x0a,0x8b,0xc7,0xd3,0x02,0x0d,0x0a,0x1c]
1269// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1270v_smfmac_i32_32x32x64_i8 a[10:25], a[2:5], a[6:13], v2 cbsz:3 abid:1
1271
1272// GFX950: v_smfmac_i32_32x32x64_i8 a[10:25], a[2:5], a[6:13], v3 cbsz:1 abid:3 ; encoding: [0x0a,0x99,0xc7,0xd3,0x02,0x0d,0x0e,0x1c]
1273// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1274v_smfmac_i32_32x32x64_i8 a[10:25], a[2:5], a[6:13], v3 cbsz:1 abid:3
1275
1276//===----------------------------------------------------------------------===//
1277// v_smfmac_f32_16x16x128_bf8_bf8
1278//===----------------------------------------------------------------------===//
1279
1280// GFX950: v_smfmac_f32_16x16x128_bf8_bf8 v[10:13], a[2:5], v[4:11], v3 cbsz:3 abid:1 ; encoding: [0x0a,0x0b,0xbb,0xd3,0x02,0x09,0x0e,0x0c]
1281// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1282v_smfmac_f32_16x16x128_bf8_bf8 v[10:13], a[2:5], v[4:11], v3 cbsz:3 abid:1
1283
1284// GFX950: v_smfmac_f32_16x16x128_bf8_bf8 v[10:13], a[2:5], v[4:11], v3 cbsz:3 abid:1 ; encoding: [0x0a,0x0b,0xbb,0xd3,0x02,0x09,0x0e,0x0c]
1285// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1286v_smfmac_f32_16x16x128bf8bf8 v[10:13], a[2:5], v[4:11], v3 cbsz:3 abid:1
1287
1288// GFX950: v_smfmac_f32_16x16x128_bf8_bf8 a[10:13], v[2:5], a[4:11], v1 ; encoding: [0x0a,0x80,0xbb,0xd3,0x02,0x09,0x06,0x14]
1289// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1290v_smfmac_f32_16x16x128_bf8_bf8 a[10:13], v[2:5], a[4:11], v1
1291
1292// GFX950: v_smfmac_f32_16x16x128_bf8_bf8 v[10:13], a[2:5], v[4:11], v2 cbsz:3 abid:1 ; encoding: [0x0a,0x0b,0xbb,0xd3,0x02,0x09,0x0a,0x0c]
1293// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1294v_smfmac_f32_16x16x128_bf8_bf8 v[10:13], a[2:5], v[4:11], v2 cbsz:3 abid:1
1295
1296// GFX950: v_smfmac_f32_16x16x128_bf8_bf8 a[10:13], v[2:5], a[4:11], v3 ; encoding: [0x0a,0x80,0xbb,0xd3,0x02,0x09,0x0e,0x14]
1297// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1298v_smfmac_f32_16x16x128_bf8_bf8 a[10:13], v[2:5], a[4:11], v3
1299
1300// GFX950: v_smfmac_f32_16x16x128_bf8_bf8 v[10:13], v[2:5], v[6:13], v2 cbsz:3 abid:1 ; encoding: [0x0a,0x0b,0xbb,0xd3,0x02,0x0d,0x0a,0x04]
1301// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1302v_smfmac_f32_16x16x128_bf8_bf8 v[10:13], v[2:5], v[6:13], v2 cbsz:3 abid:1
1303
1304// GFX950: v_smfmac_f32_16x16x128_bf8_bf8 a[10:13], a[2:5], a[6:13], v2 cbsz:3 abid:1 ; encoding: [0x0a,0x8b,0xbb,0xd3,0x02,0x0d,0x0a,0x1c]
1305// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1306v_smfmac_f32_16x16x128_bf8_bf8 a[10:13], a[2:5], a[6:13], v2 cbsz:3 abid:1
1307
1308// GFX950: v_smfmac_f32_16x16x128_bf8_bf8 a[10:13], a[2:5], a[6:13], v3 cbsz:1 abid:3 ; encoding: [0x0a,0x99,0xbb,0xd3,0x02,0x0d,0x0e,0x1c]
1309// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1310v_smfmac_f32_16x16x128_bf8_bf8 a[10:13], a[2:5], a[6:13], v3 cbsz:1 abid:3
1311
1312//===----------------------------------------------------------------------===//
1313// v_smfmac_f32_16x16x128_bf8_fp8
1314//===----------------------------------------------------------------------===//
1315
1316// GFX950: v_smfmac_f32_16x16x128_bf8_fp8 v[10:13], a[2:5], v[4:11], v3 cbsz:3 abid:1 ; encoding: [0x0a,0x0b,0xbc,0xd3,0x02,0x09,0x0e,0x0c]
1317// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1318v_smfmac_f32_16x16x128_bf8_fp8 v[10:13], a[2:5], v[4:11], v3 cbsz:3 abid:1
1319
1320// GFX950: v_smfmac_f32_16x16x128_bf8_fp8 v[10:13], a[2:5], v[4:11], v3 cbsz:3 abid:1 ; encoding: [0x0a,0x0b,0xbc,0xd3,0x02,0x09,0x0e,0x0c]
1321// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1322v_smfmac_f32_16x16x128bf8fp8 v[10:13], a[2:5], v[4:11], v3 cbsz:3 abid:1
1323
1324// GFX950: v_smfmac_f32_16x16x128_bf8_fp8 a[10:13], v[2:5], a[4:11], v1 ; encoding: [0x0a,0x80,0xbc,0xd3,0x02,0x09,0x06,0x14]
1325// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1326v_smfmac_f32_16x16x128_bf8_fp8 a[10:13], v[2:5], a[4:11], v1
1327
1328// GFX950: v_smfmac_f32_16x16x128_bf8_fp8 v[10:13], a[2:5], v[4:11], v2 cbsz:3 abid:1 ; encoding: [0x0a,0x0b,0xbc,0xd3,0x02,0x09,0x0a,0x0c]
1329// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1330v_smfmac_f32_16x16x128_bf8_fp8 v[10:13], a[2:5], v[4:11], v2 cbsz:3 abid:1
1331
1332// GFX950: v_smfmac_f32_16x16x128_bf8_fp8 a[10:13], v[2:5], a[4:11], v3 ; encoding: [0x0a,0x80,0xbc,0xd3,0x02,0x09,0x0e,0x14]
1333// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1334v_smfmac_f32_16x16x128_bf8_fp8 a[10:13], v[2:5], a[4:11], v3
1335
1336// GFX950: v_smfmac_f32_16x16x128_bf8_fp8 v[10:13], v[2:5], v[6:13], v2 cbsz:3 abid:1 ; encoding: [0x0a,0x0b,0xbc,0xd3,0x02,0x0d,0x0a,0x04]
1337// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1338v_smfmac_f32_16x16x128_bf8_fp8 v[10:13], v[2:5], v[6:13], v2 cbsz:3 abid:1
1339
1340// GFX950: v_smfmac_f32_16x16x128_bf8_fp8 a[10:13], a[2:5], a[6:13], v2 cbsz:3 abid:1 ; encoding: [0x0a,0x8b,0xbc,0xd3,0x02,0x0d,0x0a,0x1c]
1341// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1342v_smfmac_f32_16x16x128_bf8_fp8 a[10:13], a[2:5], a[6:13], v2 cbsz:3 abid:1
1343
1344// GFX950: v_smfmac_f32_16x16x128_bf8_fp8 a[10:13], a[2:5], a[6:13], v3 cbsz:1 abid:3 ; encoding: [0x0a,0x99,0xbc,0xd3,0x02,0x0d,0x0e,0x1c]
1345// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1346v_smfmac_f32_16x16x128_bf8_fp8 a[10:13], a[2:5], a[6:13], v3 cbsz:1 abid:3
1347
1348//===----------------------------------------------------------------------===//
1349// v_smfmac_f32_16x16x128_fp8_bf8
1350//===----------------------------------------------------------------------===//
1351
1352// GFX950: v_smfmac_f32_16x16x128_fp8_bf8 v[10:13], a[2:5], v[4:11], v3 cbsz:3 abid:1 ; encoding: [0x0a,0x0b,0xbd,0xd3,0x02,0x09,0x0e,0x0c]
1353// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1354v_smfmac_f32_16x16x128_fp8_bf8 v[10:13], a[2:5], v[4:11], v3 cbsz:3 abid:1
1355
1356// GFX950: v_smfmac_f32_16x16x128_fp8_bf8 v[10:13], a[2:5], v[4:11], v3 cbsz:3 abid:1 ; encoding: [0x0a,0x0b,0xbd,0xd3,0x02,0x09,0x0e,0x0c]
1357// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1358v_smfmac_f32_16x16x128fp8bf8 v[10:13], a[2:5], v[4:11], v3 cbsz:3 abid:1
1359
1360// GFX950: v_smfmac_f32_16x16x128_fp8_bf8 a[10:13], v[2:5], a[4:11], v1 ; encoding: [0x0a,0x80,0xbd,0xd3,0x02,0x09,0x06,0x14]
1361// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1362v_smfmac_f32_16x16x128_fp8_bf8 a[10:13], v[2:5], a[4:11], v1
1363
1364// GFX950: v_smfmac_f32_16x16x128_fp8_bf8 v[10:13], a[2:5], v[4:11], v2 cbsz:3 abid:1 ; encoding: [0x0a,0x0b,0xbd,0xd3,0x02,0x09,0x0a,0x0c]
1365// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1366v_smfmac_f32_16x16x128_fp8_bf8 v[10:13], a[2:5], v[4:11], v2 cbsz:3 abid:1
1367
1368// GFX950: v_smfmac_f32_16x16x128_fp8_bf8 a[10:13], v[2:5], a[4:11], v3 ; encoding: [0x0a,0x80,0xbd,0xd3,0x02,0x09,0x0e,0x14]
1369// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1370v_smfmac_f32_16x16x128_fp8_bf8 a[10:13], v[2:5], a[4:11], v3
1371
1372// GFX950: v_smfmac_f32_16x16x128_fp8_bf8 v[10:13], v[2:5], v[6:13], v2 cbsz:3 abid:1 ; encoding: [0x0a,0x0b,0xbd,0xd3,0x02,0x0d,0x0a,0x04]
1373// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1374v_smfmac_f32_16x16x128_fp8_bf8 v[10:13], v[2:5], v[6:13], v2 cbsz:3 abid:1
1375
1376// GFX950: v_smfmac_f32_16x16x128_fp8_bf8 a[10:13], a[2:5], a[6:13], v2 cbsz:3 abid:1 ; encoding: [0x0a,0x8b,0xbd,0xd3,0x02,0x0d,0x0a,0x1c]
1377// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1378v_smfmac_f32_16x16x128_fp8_bf8 a[10:13], a[2:5], a[6:13], v2 cbsz:3 abid:1
1379
1380// GFX950: v_smfmac_f32_16x16x128_fp8_bf8 a[10:13], a[2:5], a[6:13], v3 cbsz:1 abid:3 ; encoding: [0x0a,0x99,0xbd,0xd3,0x02,0x0d,0x0e,0x1c]
1381// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1382v_smfmac_f32_16x16x128_fp8_bf8 a[10:13], a[2:5], a[6:13], v3 cbsz:1 abid:3
1383
1384//===----------------------------------------------------------------------===//
1385// v_smfmac_f32_16x16x128_fp8_fp8
1386//===----------------------------------------------------------------------===//
1387
1388// GFX950: v_smfmac_f32_16x16x128_fp8_fp8 v[10:13], a[2:5], v[4:11], v3 cbsz:3 abid:1 ; encoding: [0x0a,0x0b,0xc3,0xd3,0x02,0x09,0x0e,0x0c]
1389// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1390v_smfmac_f32_16x16x128_fp8_fp8 v[10:13], a[2:5], v[4:11], v3 cbsz:3 abid:1
1391
1392// GFX950: v_smfmac_f32_16x16x128_fp8_fp8 v[10:13], a[2:5], v[4:11], v3 cbsz:3 abid:1 ; encoding: [0x0a,0x0b,0xc3,0xd3,0x02,0x09,0x0e,0x0c]
1393// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1394v_smfmac_f32_16x16x128fp8fp8 v[10:13], a[2:5], v[4:11], v3 cbsz:3 abid:1
1395
1396// GFX950: v_smfmac_f32_16x16x128_fp8_fp8 a[10:13], v[2:5], a[4:11], v1 ; encoding: [0x0a,0x80,0xc3,0xd3,0x02,0x09,0x06,0x14]
1397// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1398v_smfmac_f32_16x16x128_fp8_fp8 a[10:13], v[2:5], a[4:11], v1
1399
1400// GFX950: v_smfmac_f32_16x16x128_fp8_fp8 v[10:13], a[2:5], v[4:11], v2 cbsz:3 abid:1 ; encoding: [0x0a,0x0b,0xc3,0xd3,0x02,0x09,0x0a,0x0c]
1401// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1402v_smfmac_f32_16x16x128_fp8_fp8 v[10:13], a[2:5], v[4:11], v2 cbsz:3 abid:1
1403
1404// GFX950: v_smfmac_f32_16x16x128_fp8_fp8 a[10:13], v[2:5], a[4:11], v3 ; encoding: [0x0a,0x80,0xc3,0xd3,0x02,0x09,0x0e,0x14]
1405// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1406v_smfmac_f32_16x16x128_fp8_fp8 a[10:13], v[2:5], a[4:11], v3
1407
1408// GFX950: v_smfmac_f32_16x16x128_fp8_fp8 v[10:13], v[2:5], v[6:13], v2 cbsz:3 abid:1 ; encoding: [0x0a,0x0b,0xc3,0xd3,0x02,0x0d,0x0a,0x04]
1409// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1410v_smfmac_f32_16x16x128_fp8_fp8 v[10:13], v[2:5], v[6:13], v2 cbsz:3 abid:1
1411
1412// GFX950: v_smfmac_f32_16x16x128_fp8_fp8 a[10:13], a[2:5], a[6:13], v2 cbsz:3 abid:1 ; encoding: [0x0a,0x8b,0xc3,0xd3,0x02,0x0d,0x0a,0x1c]
1413// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1414v_smfmac_f32_16x16x128_fp8_fp8 a[10:13], a[2:5], a[6:13], v2 cbsz:3 abid:1
1415
1416// GFX950: v_smfmac_f32_16x16x128_fp8_fp8 a[10:13], a[2:5], a[6:13], v3 cbsz:1 abid:3 ; encoding: [0x0a,0x99,0xc3,0xd3,0x02,0x0d,0x0e,0x1c]
1417// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1418v_smfmac_f32_16x16x128_fp8_fp8 a[10:13], a[2:5], a[6:13], v3 cbsz:1 abid:3
1419
1420//===----------------------------------------------------------------------===//
1421// v_smfmac_f32_32x32x64_bf8_bf8
1422//===----------------------------------------------------------------------===//
1423
1424// GFX950: v_smfmac_f32_32x32x64_bf8_bf8 v[10:25], a[2:5], v[4:11], v3 cbsz:3 abid:1 ; encoding: [0x0a,0x0b,0xcb,0xd3,0x02,0x09,0x0e,0x0c]
1425// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1426v_smfmac_f32_32x32x64_bf8_bf8 v[10:25], a[2:5], v[4:11], v3 cbsz:3 abid:1
1427
1428// GFX950: v_smfmac_f32_32x32x64_bf8_bf8 v[10:25], a[2:5], v[4:11], v3 cbsz:3 abid:1 ; encoding: [0x0a,0x0b,0xcb,0xd3,0x02,0x09,0x0e,0x0c]
1429// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1430v_smfmac_f32_32x32x64bf8bf8 v[10:25], a[2:5], v[4:11], v3 cbsz:3 abid:1
1431
1432// GFX950: v_smfmac_f32_32x32x64_bf8_bf8 a[10:25], v[2:5], a[4:11], v1 ; encoding: [0x0a,0x80,0xcb,0xd3,0x02,0x09,0x06,0x14]
1433// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1434v_smfmac_f32_32x32x64_bf8_bf8 a[10:25], v[2:5], a[4:11], v1
1435
1436// GFX950: v_smfmac_f32_32x32x64_bf8_bf8 v[10:25], a[2:5], v[4:11], v2 cbsz:3 abid:1 ; encoding: [0x0a,0x0b,0xcb,0xd3,0x02,0x09,0x0a,0x0c]
1437// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1438v_smfmac_f32_32x32x64_bf8_bf8 v[10:25], a[2:5], v[4:11], v2 cbsz:3 abid:1
1439
1440// GFX950: v_smfmac_f32_32x32x64_bf8_bf8 a[10:25], v[2:5], a[4:11], v3 ; encoding: [0x0a,0x80,0xcb,0xd3,0x02,0x09,0x0e,0x14]
1441// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1442v_smfmac_f32_32x32x64_bf8_bf8 a[10:25], v[2:5], a[4:11], v3
1443
1444// GFX950: v_smfmac_f32_32x32x64_bf8_bf8 v[10:25], v[2:5], v[6:13], v2 cbsz:3 abid:1 ; encoding: [0x0a,0x0b,0xcb,0xd3,0x02,0x0d,0x0a,0x04]
1445// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1446v_smfmac_f32_32x32x64_bf8_bf8 v[10:25], v[2:5], v[6:13], v2 cbsz:3 abid:1
1447
1448// GFX950: v_smfmac_f32_32x32x64_bf8_bf8 a[10:25], a[2:5], a[6:13], v2 cbsz:3 abid:1 ; encoding: [0x0a,0x8b,0xcb,0xd3,0x02,0x0d,0x0a,0x1c]
1449// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1450v_smfmac_f32_32x32x64_bf8_bf8 a[10:25], a[2:5], a[6:13], v2 cbsz:3 abid:1
1451
1452// GFX950: v_smfmac_f32_32x32x64_bf8_bf8 a[10:25], a[2:5], a[6:13], v3 cbsz:1 abid:3 ; encoding: [0x0a,0x99,0xcb,0xd3,0x02,0x0d,0x0e,0x1c]
1453// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1454v_smfmac_f32_32x32x64_bf8_bf8 a[10:25], a[2:5], a[6:13], v3 cbsz:1 abid:3
1455
1456//===----------------------------------------------------------------------===//
1457// v_smfmac_f32_32x32x64_bf8_fp8
1458//===----------------------------------------------------------------------===//
1459
1460// GFX950: v_smfmac_f32_32x32x64_bf8_fp8 v[10:25], a[2:5], v[4:11], v3 cbsz:3 abid:1 ; encoding: [0x0a,0x0b,0xce,0xd3,0x02,0x09,0x0e,0x0c]
1461// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1462v_smfmac_f32_32x32x64_bf8_fp8 v[10:25], a[2:5], v[4:11], v3 cbsz:3 abid:1
1463
1464// GFX950: v_smfmac_f32_32x32x64_bf8_fp8 v[10:25], a[2:5], v[4:11], v3 cbsz:3 abid:1 ; encoding: [0x0a,0x0b,0xce,0xd3,0x02,0x09,0x0e,0x0c]
1465// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1466v_smfmac_f32_32x32x64bf8fp8 v[10:25], a[2:5], v[4:11], v3 cbsz:3 abid:1
1467
1468// GFX950: v_smfmac_f32_32x32x64_bf8_fp8 a[10:25], v[2:5], a[4:11], v1 ; encoding: [0x0a,0x80,0xce,0xd3,0x02,0x09,0x06,0x14]
1469// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1470v_smfmac_f32_32x32x64_bf8_fp8 a[10:25], v[2:5], a[4:11], v1
1471
1472// GFX950: v_smfmac_f32_32x32x64_bf8_fp8 v[10:25], a[2:5], v[4:11], v2 cbsz:3 abid:1 ; encoding: [0x0a,0x0b,0xce,0xd3,0x02,0x09,0x0a,0x0c]
1473// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1474v_smfmac_f32_32x32x64_bf8_fp8 v[10:25], a[2:5], v[4:11], v2 cbsz:3 abid:1
1475
1476// GFX950: v_smfmac_f32_32x32x64_bf8_fp8 a[10:25], v[2:5], a[4:11], v3 ; encoding: [0x0a,0x80,0xce,0xd3,0x02,0x09,0x0e,0x14]
1477// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1478v_smfmac_f32_32x32x64_bf8_fp8 a[10:25], v[2:5], a[4:11], v3
1479
1480// GFX950: v_smfmac_f32_32x32x64_bf8_fp8 v[10:25], v[2:5], v[6:13], v2 cbsz:3 abid:1 ; encoding: [0x0a,0x0b,0xce,0xd3,0x02,0x0d,0x0a,0x04]
1481// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1482v_smfmac_f32_32x32x64_bf8_fp8 v[10:25], v[2:5], v[6:13], v2 cbsz:3 abid:1
1483
1484// GFX950: v_smfmac_f32_32x32x64_bf8_fp8 a[10:25], a[2:5], a[6:13], v2 cbsz:3 abid:1 ; encoding: [0x0a,0x8b,0xce,0xd3,0x02,0x0d,0x0a,0x1c]
1485// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1486v_smfmac_f32_32x32x64_bf8_fp8 a[10:25], a[2:5], a[6:13], v2 cbsz:3 abid:1
1487
1488// GFX950: v_smfmac_f32_32x32x64_bf8_fp8 a[10:25], a[2:5], a[6:13], v3 cbsz:1 abid:3 ; encoding: [0x0a,0x99,0xce,0xd3,0x02,0x0d,0x0e,0x1c]
1489// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1490v_smfmac_f32_32x32x64_bf8_fp8 a[10:25], a[2:5], a[6:13], v3 cbsz:1 abid:3
1491
1492//===----------------------------------------------------------------------===//
1493// v_smfmac_f32_32x32x64_fp8_bf8
1494//===----------------------------------------------------------------------===//
1495
1496// GFX950: v_smfmac_f32_32x32x64_fp8_bf8 v[10:25], a[2:5], v[4:11], v3 cbsz:3 abid:1 ; encoding: [0x0a,0x0b,0xcf,0xd3,0x02,0x09,0x0e,0x0c]
1497// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1498v_smfmac_f32_32x32x64_fp8_bf8 v[10:25], a[2:5], v[4:11], v3 cbsz:3 abid:1
1499
1500// GFX950: v_smfmac_f32_32x32x64_fp8_bf8 v[10:25], a[2:5], v[4:11], v3 cbsz:3 abid:1 ; encoding: [0x0a,0x0b,0xcf,0xd3,0x02,0x09,0x0e,0x0c]
1501// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1502v_smfmac_f32_32x32x64fp8bf8 v[10:25], a[2:5], v[4:11], v3 cbsz:3 abid:1
1503
1504// GFX950: v_smfmac_f32_32x32x64_fp8_bf8 a[10:25], v[2:5], a[4:11], v1 ; encoding: [0x0a,0x80,0xcf,0xd3,0x02,0x09,0x06,0x14]
1505// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1506v_smfmac_f32_32x32x64_fp8_bf8 a[10:25], v[2:5], a[4:11], v1
1507
1508// GFX950: v_smfmac_f32_32x32x64_fp8_bf8 v[10:25], a[2:5], v[4:11], v2 cbsz:3 abid:1 ; encoding: [0x0a,0x0b,0xcf,0xd3,0x02,0x09,0x0a,0x0c]
1509// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1510v_smfmac_f32_32x32x64_fp8_bf8 v[10:25], a[2:5], v[4:11], v2 cbsz:3 abid:1
1511
1512// GFX950: v_smfmac_f32_32x32x64_fp8_bf8 a[10:25], v[2:5], a[4:11], v3 ; encoding: [0x0a,0x80,0xcf,0xd3,0x02,0x09,0x0e,0x14]
1513// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1514v_smfmac_f32_32x32x64_fp8_bf8 a[10:25], v[2:5], a[4:11], v3
1515
1516// GFX950: v_smfmac_f32_32x32x64_fp8_bf8 v[10:25], v[2:5], v[6:13], v2 cbsz:3 abid:1 ; encoding: [0x0a,0x0b,0xcf,0xd3,0x02,0x0d,0x0a,0x04]
1517// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1518v_smfmac_f32_32x32x64_fp8_bf8 v[10:25], v[2:5], v[6:13], v2 cbsz:3 abid:1
1519
1520// GFX950: v_smfmac_f32_32x32x64_fp8_bf8 a[10:25], a[2:5], a[6:13], v2 cbsz:3 abid:1 ; encoding: [0x0a,0x8b,0xcf,0xd3,0x02,0x0d,0x0a,0x1c]
1521// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1522v_smfmac_f32_32x32x64_fp8_bf8 a[10:25], a[2:5], a[6:13], v2 cbsz:3 abid:1
1523
1524// GFX950: v_smfmac_f32_32x32x64_fp8_bf8 a[10:25], a[2:5], a[6:13], v3 cbsz:1 abid:3 ; encoding: [0x0a,0x99,0xcf,0xd3,0x02,0x0d,0x0e,0x1c]
1525// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1526v_smfmac_f32_32x32x64_fp8_bf8 a[10:25], a[2:5], a[6:13], v3 cbsz:1 abid:3
1527
1528//===----------------------------------------------------------------------===//
1529// v_smfmac_f32_32x32x64_fp8_fp8
1530//===----------------------------------------------------------------------===//
1531
1532// GFX950: v_smfmac_f32_32x32x64_fp8_fp8 v[10:25], a[2:5], v[4:11], v3 cbsz:3 abid:1 ; encoding: [0x0a,0x0b,0xd3,0xd3,0x02,0x09,0x0e,0x0c]
1533// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1534v_smfmac_f32_32x32x64_fp8_fp8 v[10:25], a[2:5], v[4:11], v3 cbsz:3 abid:1
1535
1536// GFX950: v_smfmac_f32_32x32x64_fp8_fp8 v[10:25], a[2:5], v[4:11], v3 cbsz:3 abid:1 ; encoding: [0x0a,0x0b,0xd3,0xd3,0x02,0x09,0x0e,0x0c]
1537// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1538v_smfmac_f32_32x32x64fp8fp8 v[10:25], a[2:5], v[4:11], v3 cbsz:3 abid:1
1539
1540// GFX950: v_smfmac_f32_32x32x64_fp8_fp8 a[10:25], v[2:5], a[4:11], v1 ; encoding: [0x0a,0x80,0xd3,0xd3,0x02,0x09,0x06,0x14]
1541// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1542v_smfmac_f32_32x32x64_fp8_fp8 a[10:25], v[2:5], a[4:11], v1
1543
1544// GFX950: v_smfmac_f32_32x32x64_fp8_fp8 v[10:25], a[2:5], v[4:11], v2 cbsz:3 abid:1 ; encoding: [0x0a,0x0b,0xd3,0xd3,0x02,0x09,0x0a,0x0c]
1545// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1546v_smfmac_f32_32x32x64_fp8_fp8 v[10:25], a[2:5], v[4:11], v2 cbsz:3 abid:1
1547
1548// GFX950: v_smfmac_f32_32x32x64_fp8_fp8 a[10:25], v[2:5], a[4:11], v3 ; encoding: [0x0a,0x80,0xd3,0xd3,0x02,0x09,0x0e,0x14]
1549// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1550v_smfmac_f32_32x32x64_fp8_fp8 a[10:25], v[2:5], a[4:11], v3
1551
1552// GFX950: v_smfmac_f32_32x32x64_fp8_fp8 v[10:25], v[2:5], v[6:13], v2 cbsz:3 abid:1 ; encoding: [0x0a,0x0b,0xd3,0xd3,0x02,0x0d,0x0a,0x04]
1553// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1554v_smfmac_f32_32x32x64_fp8_fp8 v[10:25], v[2:5], v[6:13], v2 cbsz:3 abid:1
1555
1556// GFX950: v_smfmac_f32_32x32x64_fp8_fp8 a[10:25], a[2:5], a[6:13], v2 cbsz:3 abid:1 ; encoding: [0x0a,0x8b,0xd3,0xd3,0x02,0x0d,0x0a,0x1c]
1557// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1558v_smfmac_f32_32x32x64_fp8_fp8 a[10:25], a[2:5], a[6:13], v2 cbsz:3 abid:1
1559
1560// GFX950: v_smfmac_f32_32x32x64_fp8_fp8 a[10:25], a[2:5], a[6:13], v3 cbsz:1 abid:3 ; encoding: [0x0a,0x99,0xd3,0xd3,0x02,0x0d,0x0e,0x1c]
1561// ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1562v_smfmac_f32_32x32x64_fp8_fp8 a[10:25], a[2:5], a[6:13], v3 cbsz:1 abid:3
1563