Skip to content

Commit c310306

Browse files
rampitecrampitec
authored
[AMDGPU] Add more gfx1250 MC tests. NFC. (#152388)
These are already working, but left downstream.
1 parent 87404ea commit c310306

14 files changed

+11761
-0
lines changed

‎llvm/test/MC/AMDGPU/gfx1250_asm_features.s‎

Lines changed: 32 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,32 @@
1+
// RUN: llvm-mc -triple=amdgcn -mcpu=gfx1250 -show-encoding %s | FileCheck --check-prefixes=GFX1250 %s
2+
3+
//
4+
// Elements of CPol operand can be given in any order
5+
//
6+
7+
s_load_b32 s4, s[2:3], 10 th:TH_LOAD_NT scope:SCOPE_SE nv
8+
// GFX1250: encoding: [0x01,0x01,0xb0,0xf4,0x0a,0x00,0x00,0xf8]
9+
10+
s_load_b32 s4, s[2:3], 10 scope:SCOPE_SE nv th:TH_LOAD_NT
11+
// GFX1250: encoding: [0x01,0x01,0xb0,0xf4,0x0a,0x00,0x00,0xf8]
12+
13+
s_load_b32 s4, s[2:3], 10 nv scope:SCOPE_SE th:TH_LOAD_NT
14+
// GFX1250: encoding: [0x01,0x01,0xb0,0xf4,0x0a,0x00,0x00,0xf8]
15+
16+
buffer_load_b32 v5, v1, s[8:11], s3 offen offset:4095 th:TH_LOAD_NT scope:SCOPE_SE nv
17+
// GFX1250: encoding: [0x83,0x00,0x05,0xc4,0x05,0x10,0x94,0x40,0x01,0xff,0x0f,0x00]
18+
19+
buffer_load_b32 v5, v1, s[8:11], s3 offen offset:4095 scope:SCOPE_SE nv th:TH_LOAD_NT
20+
// GFX1250: encoding: [0x83,0x00,0x05,0xc4,0x05,0x10,0x94,0x40,0x01,0xff,0x0f,0x00]
21+
22+
buffer_load_b32 v5, v1, s[8:11], s3 offen offset:4095 nv scope:SCOPE_SE th:TH_LOAD_NT
23+
// GFX1250: encoding: [0x83,0x00,0x05,0xc4,0x05,0x10,0x94,0x40,0x01,0xff,0x0f,0x00]
24+
25+
global_load_b32 v0, v[2:3], off th:TH_LOAD_NT scope:SCOPE_SE nv
26+
// GFX1250: encoding: [0xfc,0x00,0x05,0xee,0x00,0x00,0x14,0x00,0x02,0x00,0x00,0x00]
27+
28+
global_load_b32 v0, v[2:3], off scope:SCOPE_SE nv th:TH_LOAD_NT
29+
// GFX1250: encoding: [0xfc,0x00,0x05,0xee,0x00,0x00,0x14,0x00,0x02,0x00,0x00,0x00]
30+
31+
global_load_b32 v0, v[2:3], off nv scope:SCOPE_SE th:TH_LOAD_NT
32+
// GFX1250: encoding: [0xfc,0x00,0x05,0xee,0x00,0x00,0x14,0x00,0x02,0x00,0x00,0x00]

‎llvm/test/MC/AMDGPU/gfx1250_asm_unsupported.s‎

Lines changed: 14 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -97,6 +97,20 @@ v_interp_p10_rtz_f16_f32 v0, v1, v2, v3
9797
v_interp_p2_rtz_f16_f32 v0, v1, v2, v3
9898
// GFX1250-ERR: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
9999

100+
;; *xf32
101+
102+
v_mfma_f32_16x16x8_xf32 a[0:3], v[2:3], v[4:5], a[2:5]
103+
// GFX1250-ERR: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
104+
105+
v_mfma_f32_16x16x8xf32 a[0:3], v[2:3], v[4:5], a[2:5]
106+
// GFX1250-ERR: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
107+
108+
v_mfma_f32_32x32x4_xf32 a[0:15], v[2:3], v[4:5], a[18:33]
109+
// GFX1250-ERR: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
110+
111+
v_mfma_f32_32x32x4xf32 a[0:15], v[2:3], v[4:5], a[18:33]
112+
// GFX1250-ERR: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
113+
100114
;; Export, S_WAIT_EXPCNT and S_WAIT_EVENT
101115

102116
export mrt0 off, off, off, off

‎llvm/test/MC/AMDGPU/gfx1250_asm_vbuffer_mubuf.s‎

Lines changed: 2304 additions & 0 deletions
Large diffs are not rendered by default.

‎llvm/test/MC/AMDGPU/gfx1250_asm_vop3_err.s‎

Lines changed: 165 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -5,6 +5,76 @@ v_lshl_add_u64 v[2:3], v[4:5], v7, v[8:9] dpp8:[7,6,5,4,3,2,1,0]
55
// GFX125X-ERR-NEXT:{{^}}v_lshl_add_u64 v[2:3], v[4:5], v7, v[8:9] dpp8:[7,6,5,4,3,2,1,0]
66
// GFX125X-ERR-NEXT:{{^}} ^
77

8+
v_fma_f64 v[4:5], v[2:3], v[6:7], v[8:9] dpp8:[7,6,5,4,3,2,1,0]
9+
// GFX125X-ERR: :[[@LINE-1]]:{{[0-9]+}}: error: not a valid operand.
10+
// GFX125X-ERR-NEXT:{{^}}v_fma_f64 v[4:5], v[2:3], v[6:7], v[8:9] dpp8:[7,6,5,4,3,2,1,0]
11+
// GFX125X-ERR-NEXT:{{^}} ^
12+
13+
v_div_fixup_f64 v[4:5], v[2:3], v[6:7], v[8:9] dpp8:[7,6,5,4,3,2,1,0]
14+
// GFX125X-ERR: :[[@LINE-1]]:{{[0-9]+}}: error: not a valid operand.
15+
// GFX125X-ERR-NEXT:{{^}}v_div_fixup_f64 v[4:5], v[2:3], v[6:7], v[8:9] dpp8:[7,6,5,4,3,2,1,0]
16+
// GFX125X-ERR-NEXT:{{^}} ^
17+
18+
v_div_fmas_f64 v[4:5], v[2:3], v[6:7], v[8:9] dpp8:[7,6,5,4,3,2,1,0]
19+
// GFX125X-ERR: :[[@LINE-1]]:{{[0-9]+}}: error: not a valid operand.
20+
// GFX125X-ERR-NEXT:{{^}}v_div_fmas_f64 v[4:5], v[2:3], v[6:7], v[8:9] dpp8:[7,6,5,4,3,2,1,0]
21+
// GFX125X-ERR-NEXT:{{^}} ^
22+
23+
v_div_scale_f64 v[4:5], s2, v[2:3], v[6:7], v[8:9] dpp8:[7,6,5,4,3,2,1,0]
24+
// GFX125X-ERR: :[[@LINE-1]]:{{[0-9]+}}: error: not a valid operand.
25+
// GFX125X-ERR-NEXT:{{^}}v_div_scale_f64 v[4:5], s2, v[2:3], v[6:7], v[8:9] dpp8:[7,6,5,4,3,2,1,0]
26+
// GFX125X-ERR-NEXT:{{^}} ^
27+
28+
v_mad_co_u64_u32 v[4:5], s2, v2, v6, v[8:9] dpp8:[7,6,5,4,3,2,1,0]
29+
// GFX125X-ERR: :[[@LINE-1]]:{{[0-9]+}}: error: not a valid operand.
30+
// GFX125X-ERR-NEXT:{{^}}v_mad_co_u64_u32 v[4:5], s2, v2, v6, v[8:9] dpp8:[7,6,5,4,3,2,1,0]
31+
// GFX125X-ERR-NEXT:{{^}} ^
32+
33+
v_mad_co_i64_i32 v[4:5], s2, v2, v6, v[8:9] dpp8:[7,6,5,4,3,2,1,0]
34+
// GFX125X-ERR: :[[@LINE-1]]:{{[0-9]+}}: error: not a valid operand.
35+
// GFX125X-ERR-NEXT:{{^}}v_mad_co_i64_i32 v[4:5], s2, v2, v6, v[8:9] dpp8:[7,6,5,4,3,2,1,0]
36+
// GFX125X-ERR-NEXT:{{^}} ^
37+
38+
v_minimum_f64 v[4:5], v[2:3], v[6:7] dpp8:[7,6,5,4,3,2,1,0]
39+
// GFX125X-ERR: :[[@LINE-1]]:{{[0-9]+}}: error: not a valid operand.
40+
// GFX125X-ERR-NEXT:{{^}}v_minimum_f64 v[4:5], v[2:3], v[6:7] dpp8:[7,6,5,4,3,2,1,0]
41+
// GFX125X-ERR-NEXT:{{^}} ^
42+
43+
v_maximum_f64 v[4:5], v[2:3], v[6:7] dpp8:[7,6,5,4,3,2,1,0]
44+
// GFX125X-ERR: :[[@LINE-1]]:{{[0-9]+}}: error: not a valid operand.
45+
// GFX125X-ERR-NEXT:{{^}}v_maximum_f64 v[4:5], v[2:3], v[6:7] dpp8:[7,6,5,4,3,2,1,0]
46+
// GFX125X-ERR-NEXT:{{^}} ^
47+
48+
v_ldexp_f64 v[4:5], v[2:3], v6 dpp8:[7,6,5,4,3,2,1,0]
49+
// GFX125X-ERR: :[[@LINE-1]]:{{[0-9]+}}: error: not a valid operand.
50+
// GFX125X-ERR-NEXT:{{^}}v_ldexp_f64 v[4:5], v[2:3], v6 dpp8:[7,6,5,4,3,2,1,0]
51+
// GFX125X-ERR-NEXT:{{^}} ^
52+
53+
v_mul_lo_u32 v4, v2, v6 dpp8:[7,6,5,4,3,2,1,0]
54+
// GFX125X-ERR: :[[@LINE-1]]:{{[0-9]+}}: error: not a valid operand.
55+
// GFX125X-ERR-NEXT:{{^}}v_mul_lo_u32 v4, v2, v6 dpp8:[7,6,5,4,3,2,1,0]
56+
// GFX125X-ERR-NEXT:{{^}} ^
57+
58+
v_mul_hi_u32 v4, v2, v6 dpp8:[7,6,5,4,3,2,1,0]
59+
// GFX125X-ERR: :[[@LINE-1]]:{{[0-9]+}}: error: not a valid operand.
60+
// GFX125X-ERR-NEXT:{{^}}v_mul_hi_u32 v4, v2, v6 dpp8:[7,6,5,4,3,2,1,0]
61+
// GFX125X-ERR-NEXT:{{^}} ^
62+
63+
v_mul_hi_i32 v4, v2, v6 dpp8:[7,6,5,4,3,2,1,0]
64+
// GFX125X-ERR: :[[@LINE-1]]:{{[0-9]+}}: error: not a valid operand.
65+
// GFX125X-ERR-NEXT:{{^}}v_mul_hi_i32 v4, v2, v6 dpp8:[7,6,5,4,3,2,1,0]
66+
// GFX125X-ERR-NEXT:{{^}} ^
67+
68+
v_lshrrev_b64 v[4:5], v2, v[6:7] dpp8:[7,6,5,4,3,2,1,0]
69+
// GFX125X-ERR: :[[@LINE-1]]:{{[0-9]+}}: error: not a valid operand.
70+
// GFX125X-ERR-NEXT:{{^}}v_lshrrev_b64 v[4:5], v2, v[6:7] dpp8:[7,6,5,4,3,2,1,0]
71+
// GFX125X-ERR-NEXT:{{^}} ^
72+
73+
v_ashrrev_i64 v[4:5], v2, v[6:7] dpp8:[7,6,5,4,3,2,1,0]
74+
// GFX125X-ERR: :[[@LINE-1]]:{{[0-9]+}}: error: not a valid operand.
75+
// GFX125X-ERR-NEXT:{{^}}v_ashrrev_i64 v[4:5], v2, v[6:7] dpp8:[7,6,5,4,3,2,1,0]
76+
// GFX125X-ERR-NEXT:{{^}} ^
77+
878
v_mad_u32 v2, v4, v7, v8 dpp8:[7,6,5,4,3,2,1,0]
979
// GFX125X-ERR: :[[@LINE-1]]:{{[0-9]+}}: error: not a valid operand.
1080
// GFX125X-ERR-NEXT:{{^}}v_mad_u32 v2, v4, v7, v8 dpp8:[7,6,5,4,3,2,1,0]
@@ -42,9 +112,94 @@ v_mad_nc_i64_i32 v[4:5], v2, v5, v[6:7] dpp8:[7,6,5,4,3,2,1,0]
42112

43113
v_lshl_add_u64 v[2:3], v[4:5], v7, v[8:9] quad_perm:[3,2,1,0]
44114
// GFX1250-ERR: :[[@LINE-1]]:{{[0-9]+}}: error: not a valid operand.
115+
// GFX1251-ERR: :[[@LINE-2]]:{{[0-9]+}}: error: DP ALU dpp only supports row_share
45116
// GFX125X-ERR-NEXT:{{^}}v_lshl_add_u64 v[2:3], v[4:5], v7, v[8:9] quad_perm:[3,2,1,0]
46117
// GFX125X-ERR-NEXT:{{^}} ^
47118

119+
v_fma_f64 v[4:5], v[2:3], v[6:7], v[8:9] quad_perm:[3,2,1,0]
120+
// GFX1250-ERR: :[[@LINE-1]]:{{[0-9]+}}: error: not a valid operand.
121+
// GFX1251-ERR: :[[@LINE-2]]:{{[0-9]+}}: error: DP ALU dpp only supports row_share
122+
// GFX125X-ERR-NEXT:{{^}}v_fma_f64 v[4:5], v[2:3], v[6:7], v[8:9] quad_perm:[3,2,1,0]
123+
// GFX125X-ERR-NEXT:{{^}} ^
124+
125+
v_div_fixup_f64 v[4:5], v[2:3], v[6:7], v[8:9] quad_perm:[3,2,1,0]
126+
// GFX1250-ERR: :[[@LINE-1]]:{{[0-9]+}}: error: not a valid operand.
127+
// GFX1251-ERR: :[[@LINE-2]]:{{[0-9]+}}: error: DP ALU dpp only supports row_share
128+
// GFX125X-ERR-NEXT:{{^}}v_div_fixup_f64 v[4:5], v[2:3], v[6:7], v[8:9] quad_perm:[3,2,1,0]
129+
// GFX125X-ERR-NEXT:{{^}} ^
130+
131+
v_div_fmas_f64 v[4:5], v[2:3], v[6:7], v[8:9] quad_perm:[3,2,1,0]
132+
// GFX1250-ERR: :[[@LINE-1]]:{{[0-9]+}}: error: not a valid operand.
133+
// GFX1251-ERR: :[[@LINE-2]]:{{[0-9]+}}: error: DP ALU dpp only supports row_share
134+
// GFX125X-ERR-NEXT:{{^}}v_div_fmas_f64 v[4:5], v[2:3], v[6:7], v[8:9] quad_perm:[3,2,1,0]
135+
// GFX125X-ERR-NEXT:{{^}} ^
136+
137+
v_div_scale_f64 v[4:5], s2, v[2:3], v[6:7], v[8:9] quad_perm:[3,2,1,0]
138+
// GFX1250-ERR: :[[@LINE-1]]:{{[0-9]+}}: error: not a valid operand.
139+
// GFX1251-ERR: :[[@LINE-2]]:{{[0-9]+}}: error: DP ALU dpp only supports row_share
140+
// GFX125X-ERR-NEXT:{{^}}v_div_scale_f64 v[4:5], s2, v[2:3], v[6:7], v[8:9] quad_perm:[3,2,1,0]
141+
// GFX125X-ERR-NEXT:{{^}} ^
142+
143+
v_mad_co_u64_u32 v[4:5], s2, v2, v6, v[8:9] quad_perm:[3,2,1,0]
144+
// GFX1250-ERR: :[[@LINE-1]]:{{[0-9]+}}: error: not a valid operand.
145+
// GFX1251-ERR: :[[@LINE-2]]:{{[0-9]+}}: error: DP ALU dpp only supports row_share
146+
// GFX125X-ERR-NEXT:{{^}}v_mad_co_u64_u32 v[4:5], s2, v2, v6, v[8:9] quad_perm:[3,2,1,0]
147+
// GFX125X-ERR-NEXT:{{^}} ^
148+
149+
v_mad_co_i64_i32 v[4:5], s2, v2, v6, v[8:9] quad_perm:[3,2,1,0]
150+
// GFX1250-ERR: :[[@LINE-1]]:{{[0-9]+}}: error: not a valid operand.
151+
// GFX1251-ERR: :[[@LINE-2]]:{{[0-9]+}}: error: DP ALU dpp only supports row_share
152+
// GFX125X-ERR-NEXT:{{^}}v_mad_co_i64_i32 v[4:5], s2, v2, v6, v[8:9] quad_perm:[3,2,1,0]
153+
// GFX125X-ERR-NEXT:{{^}} ^
154+
155+
v_minimum_f64 v[4:5], v[2:3], v[6:7] quad_perm:[3,2,1,0]
156+
// GFX1250-ERR: :[[@LINE-1]]:{{[0-9]+}}: error: not a valid operand.
157+
// GFX1251-ERR: :[[@LINE-2]]:{{[0-9]+}}: error: DP ALU dpp only supports row_share
158+
// GFX125X-ERR-NEXT:{{^}}v_minimum_f64 v[4:5], v[2:3], v[6:7] quad_perm:[3,2,1,0]
159+
// GFX125X-ERR-NEXT:{{^}} ^
160+
161+
v_maximum_f64 v[4:5], v[2:3], v[6:7] quad_perm:[3,2,1,0]
162+
// GFX1250-ERR: :[[@LINE-1]]:{{[0-9]+}}: error: not a valid operand.
163+
// GFX1251-ERR: :[[@LINE-2]]:{{[0-9]+}}: error: DP ALU dpp only supports row_share
164+
// GFX125X-ERR-NEXT:{{^}}v_maximum_f64 v[4:5], v[2:3], v[6:7] quad_perm:[3,2,1,0]
165+
// GFX125X-ERR-NEXT:{{^}} ^
166+
167+
v_ldexp_f64 v[4:5], v[2:3], v6 quad_perm:[3,2,1,0]
168+
// GFX1250-ERR: :[[@LINE-1]]:{{[0-9]+}}: error: not a valid operand.
169+
// GFX1251-ERR: :[[@LINE-2]]:{{[0-9]+}}: error: DP ALU dpp only supports row_share
170+
// GFX125X-ERR-NEXT:{{^}}v_ldexp_f64 v[4:5], v[2:3], v6 quad_perm:[3,2,1,0]
171+
// GFX125X-ERR-NEXT:{{^}} ^
172+
173+
v_mul_lo_u32 v4, v2, v6 quad_perm:[3,2,1,0]
174+
// GFX1250-ERR: :[[@LINE-1]]:{{[0-9]+}}: error: not a valid operand.
175+
// GFX1251-ERR: :[[@LINE-2]]:{{[0-9]+}}: error: DP ALU dpp only supports row_share
176+
// GFX125X-ERR-NEXT:{{^}}v_mul_lo_u32 v4, v2, v6 quad_perm:[3,2,1,0]
177+
// GFX125X-ERR-NEXT:{{^}} ^
178+
179+
v_mul_hi_u32 v4, v2, v6 quad_perm:[3,2,1,0]
180+
// GFX1250-ERR: :[[@LINE-1]]:{{[0-9]+}}: error: not a valid operand.
181+
// GFX1251-ERR: :[[@LINE-2]]:{{[0-9]+}}: error: DP ALU dpp only supports row_share
182+
// GFX125X-ERR-NEXT:{{^}}v_mul_hi_u32 v4, v2, v6 quad_perm:[3,2,1,0]
183+
// GFX125X-ERR-NEXT:{{^}} ^
184+
185+
v_mul_hi_i32 v4, v2, v6 quad_perm:[3,2,1,0]
186+
// GFX1250-ERR: :[[@LINE-1]]:{{[0-9]+}}: error: not a valid operand.
187+
// GFX1251-ERR: :[[@LINE-2]]:{{[0-9]+}}: error: DP ALU dpp only supports row_share
188+
// GFX125X-ERR-NEXT:{{^}}v_mul_hi_i32 v4, v2, v6 quad_perm:[3,2,1,0]
189+
// GFX125X-ERR-NEXT:{{^}} ^
190+
191+
v_lshrrev_b64 v[4:5], v2, v[6:7] quad_perm:[3,2,1,0]
192+
// GFX1250-ERR: :[[@LINE-1]]:{{[0-9]+}}: error: not a valid operand.
193+
// GFX1251-ERR: :[[@LINE-2]]:{{[0-9]+}}: error: DP ALU dpp only supports row_share
194+
// GFX125X-ERR-NEXT:{{^}}v_lshrrev_b64 v[4:5], v2, v[6:7] quad_perm:[3,2,1,0]
195+
// GFX125X-ERR-NEXT:{{^}} ^
196+
197+
v_ashrrev_i64 v[4:5], v2, v[6:7] quad_perm:[3,2,1,0]
198+
// GFX1250-ERR: :[[@LINE-1]]:{{[0-9]+}}: error: not a valid operand.
199+
// GFX1251-ERR: :[[@LINE-2]]:{{[0-9]+}}: error: DP ALU dpp only supports row_share
200+
// GFX125X-ERR-NEXT:{{^}}v_ashrrev_i64 v[4:5], v2, v[6:7] quad_perm:[3,2,1,0]
201+
// GFX125X-ERR-NEXT:{{^}} ^
202+
48203
v_mad_u32 v2, v4, v7, v8 quad_perm:[3,2,1,0]
49204
// GFX1250-ERR: :[[@LINE-1]]:{{[0-9]+}}: error: not a valid operand.
50205
// GFX1251-ERR: :[[@LINE-2]]:{{[0-9]+}}: error: DP ALU dpp only supports row_share
@@ -87,6 +242,11 @@ v_mad_nc_i64_i32 v[4:5], v2, v5, v[6:7] quad_perm:[3,2,1,0]
87242
// GFX125X-ERR-NEXT:{{^}}v_mad_nc_i64_i32 v[4:5], v2, v5, v[6:7] quad_perm:[3,2,1,0]
88243
// GFX125X-ERR-NEXT:{{^}} ^
89244

245+
v_trig_preop_f64 v[4:5], v[8:9], v2 row_share:1
246+
// GFX125X-ERR: :[[@LINE-1]]:{{[0-9]+}}: error: not a valid operand.
247+
// GFX125X-ERR-NEXT:{{^}}v_trig_preop_f64 v[4:5], v[8:9], v2 row_share:1
248+
// GFX125X-ERR-NEXT:{{^}} ^
249+
90250
v_ashr_pk_i8_i32 v1, v2, v3, v4 clamp
91251
// GFX125X-ERR: :[[@LINE-1]]:{{[0-9]+}}: error: invalid operand for instruction
92252
// GFX125X-ERR-NEXT:{{^}}v_ashr_pk_i8_i32 v1, v2, v3, v4 clamp
@@ -161,3 +321,8 @@ v_cvt_scale_pk8_f32_fp4 v[10:17], s20, v8
161321
// GFX125X-ERR: :[[@LINE-1]]:{{[0-9]+}}: error: invalid operand for instruction
162322
// GFX125X-ERR-NEXT:{{^}}v_cvt_scale_pk8_f32_fp4 v[10:17], s20, v8
163323
// GFX125X-ERR-NEXT:{{^}} ^
324+
325+
v_cvt_scale_pk16_bf16_bf6 v[10:17], s[20:22], 0xcf00
326+
// GFX125X-ERR: :[[@LINE-1]]:{{[0-9]+}}: error: invalid operand for instruction
327+
// GFX125X-ERR-NEXT:{{^}}v_cvt_scale_pk16_bf16_bf6 v[10:17], s[20:22], 0xcf00
328+
// GFX125X-ERR-NEXT:{{^}} ^

‎llvm/test/MC/AMDGPU/gfx1250_asm_vop3_from_vop2_err.s‎

Lines changed: 13 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,13 @@
1+
// RUN: not llvm-mc -triple=amdgcn -mcpu=gfx1250 -show-encoding %s 2>&1 | FileCheck --check-prefix=GFX1250-ERR --implicit-check-not=error: --strict-whitespace %s
2+
3+
v_fmaak_f32_e64_dpp v4, v2, v6, 3 row_share:1
4+
// GFX1250-ERR: :[[@LINE-1]]:{{[0-9]+}}: error: e64_dpp variant of this instruction is not supported
5+
6+
v_fmamk_f32_e64_dpp v4, v2, 3, v6 row_share:1
7+
// GFX1250-ERR: :[[@LINE-1]]:{{[0-9]+}}: error: e64_dpp variant of this instruction is not supported
8+
9+
v_fmaak_f16_e64_dpp v4, v2, v6, 3 row_share:1
10+
// GFX1250-ERR: :[[@LINE-1]]:{{[0-9]+}}: error: e64_dpp variant of this instruction is not supported
11+
12+
v_fmamk_f16_e64_dpp v4, v2, 3, v6 row_share:1
13+
// GFX1250-ERR: :[[@LINE-1]]:{{[0-9]+}}: error: e64_dpp variant of this instruction is not supported

0 commit comments

Comments
 (0)