Skip to content

Commit f405e04

Browse files
arsenmarsenm
authored
AMDGPU: Test a few more cases for assembler errors for misaligned gfx90a vgprs (#156998)
1 parent 9a0ed70 commit f405e04

File tree

1 file changed

+81
-0
lines changed

1 file changed

+81
-0
lines changed

‎llvm/test/MC/AMDGPU/misaligned-vgpr-tuples-err.s‎

Lines changed: 81 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -3,15 +3,96 @@
33
v_add_f64 v[1:2], v[1:2], v[1:2]
44
// GFX90A: :[[@LINE-1]]:{{[0-9]+}}: error: invalid register class: vgpr tuples must be 64 bit aligned
55

6+
global_load_dwordx2 v[1:2], v[0:1], off
7+
// GFX90A: :[[@LINE-1]]:{{[0-9]+}}: error: invalid register class: vgpr tuples must be 64 bit aligned
8+
69
global_load_dwordx3 v[1:3], v[0:1], off
710
// GFX90A: :[[@LINE-1]]:{{[0-9]+}}: error: invalid register class: vgpr tuples must be 64 bit aligned
811

912
global_load_dwordx4 v[1:4], v[0:1], off
1013
// GFX90A: :[[@LINE-1]]:{{[0-9]+}}: error: invalid register class: vgpr tuples must be 64 bit aligned
1114

15+
global_load_dwordx2 a[1:2], v[0:1], off
16+
// GFX90A: :[[@LINE-1]]:{{[0-9]+}}: error: invalid register class: vgpr tuples must be 64 bit aligned
17+
18+
global_load_dwordx3 a[1:3], v[0:1], off
19+
// GFX90A: :[[@LINE-1]]:{{[0-9]+}}: error: invalid register class: vgpr tuples must be 64 bit aligned
20+
21+
global_load_dwordx4 a[1:4], v[0:1], off
22+
// GFX90A: :[[@LINE-1]]:{{[0-9]+}}: error: invalid register class: vgpr tuples must be 64 bit aligned
23+
24+
25+
image_load v[1:2], v2, s[0:7] dmask:0x3 unorm
26+
// GFX90A: :[[@LINE-1]]:{{[0-9]+}}: error: invalid register class: vgpr tuples must be 64 bit aligned
27+
28+
image_load v[1:3], v2, s[0:7] dmask:0x7 unorm
29+
// GFX90A: :[[@LINE-1]]:{{[0-9]+}}: error: invalid register class: vgpr tuples must be 64 bit aligned
30+
1231
image_load v[1:4], v2, s[0:7] dmask:0xf unorm
1332
// GFX90A: :[[@LINE-1]]:{{[0-9]+}}: error: invalid register class: vgpr tuples must be 64 bit aligned
1433

34+
image_load a[1:2], v2, s[0:7] dmask:0x3 unorm
35+
// GFX90A: :[[@LINE-1]]:{{[0-9]+}}: error: invalid register class: vgpr tuples must be 64 bit aligned
36+
37+
image_load a[1:3], v2, s[0:7] dmask:0x7 unorm
38+
// GFX90A: :[[@LINE-1]]:{{[0-9]+}}: error: invalid register class: vgpr tuples must be 64 bit aligned
39+
40+
image_load a[1:4], v2, s[0:7] dmask:0xf unorm
41+
// GFX90A: :[[@LINE-1]]:{{[0-9]+}}: error: invalid register class: vgpr tuples must be 64 bit aligned
42+
43+
44+
image_store v[193:194], v[238:241], s[28:35] dmask:0x3 unorm
45+
// GFX90A: :[[@LINE-1]]:{{[0-9]+}}: error: invalid register class: vgpr tuples must be 64 bit aligned
46+
47+
image_store v[193:195], v[238:241], s[28:35] dmask:0x7 unorm
48+
// GFX90A: :[[@LINE-1]]:{{[0-9]+}}: error: invalid register class: vgpr tuples must be 64 bit aligned
49+
50+
image_store v[193:196], v[238:241], s[28:35] dmask:0xf unorm
51+
// GFX90A: :[[@LINE-1]]:{{[0-9]+}}: error: invalid register class: vgpr tuples must be 64 bit aligned
52+
53+
image_store a[193:194], v[238:241], s[28:35] dmask:0x3 unorm
54+
// GFX90A: :[[@LINE-1]]:{{[0-9]+}}: error: invalid register class: vgpr tuples must be 64 bit aligned
55+
56+
image_store a[193:195], v[238:241], s[28:35] dmask:0x7 unorm
57+
// GFX90A: :[[@LINE-1]]:{{[0-9]+}}: error: invalid register class: vgpr tuples must be 64 bit aligned
58+
59+
image_store a[193:196], v[238:241], s[28:35] dmask:0xf unorm
60+
// GFX90A: :[[@LINE-1]]:{{[0-9]+}}: error: invalid register class: vgpr tuples must be 64 bit aligned
61+
62+
63+
image_atomic_swap v4, v[193:196], s[28:35] dmask:0x1 unorm glc
64+
// GFX90A: :[[@LINE-1]]:{{[0-9]+}}: error: invalid register class: vgpr tuples must be 64 bit aligned
65+
66+
image_atomic_swap v[5:6], v1, s[8:15] dmask:0x3 unorm
67+
// GFX90A: :[[@LINE-1]]:{{[0-9]+}}: error: invalid register class: vgpr tuples must be 64 bit aligned
68+
69+
70+
image_atomic_cmpswap v[5:6], v[192:195], s[28:35] dmask:0x3 unorm glc
71+
// GFX90A: :[[@LINE-1]]:{{[0-9]+}}: error: invalid register class: vgpr tuples must be 64 bit aligned
72+
73+
image_atomic_cmpswap v[4:5], v[193:196], s[28:35] dmask:0x3 unorm glc
74+
// GFX90A: :[[@LINE-1]]:{{[0-9]+}}: error: invalid register class: vgpr tuples must be 64 bit aligned
75+
76+
image_atomic_cmpswap v[5:8], v[192:195], s[28:35] dmask:0xf unorm glc
77+
// GFX90A: :[[@LINE-1]]:{{[0-9]+}}: error: invalid register class: vgpr tuples must be 64 bit aligned
78+
79+
image_atomic_cmpswap v[4:7], v[193:196], s[28:35] dmask:0xf unorm glc
80+
// GFX90A: :[[@LINE-1]]:{{[0-9]+}}: error: invalid register class: vgpr tuples must be 64 bit aligned
81+
82+
83+
image_atomic_cmpswap a[5:6], v[192:195], s[28:35] dmask:0x3 unorm glc
84+
// GFX90A: :[[@LINE-1]]:{{[0-9]+}}: error: invalid register class: vgpr tuples must be 64 bit aligned
85+
86+
image_atomic_cmpswap a[4:5], v[193:196], s[28:35] dmask:0x3 unorm glc
87+
// GFX90A: :[[@LINE-1]]:{{[0-9]+}}: error: invalid register class: vgpr tuples must be 64 bit aligned
88+
89+
image_atomic_cmpswap a[5:8], v[192:195], s[28:35] dmask:0xf unorm glc
90+
// GFX90A: :[[@LINE-1]]:{{[0-9]+}}: error: invalid register class: vgpr tuples must be 64 bit aligned
91+
92+
image_atomic_cmpswap a[4:7], v[193:196], s[28:35] dmask:0xf unorm glc
93+
// GFX90A: :[[@LINE-1]]:{{[0-9]+}}: error: invalid register class: vgpr tuples must be 64 bit aligned
94+
95+
1596
v_mfma_f32_32x32x8f16 a[0:15], a[1:2], v[0:1], a[0:15]
1697
// GFX90A: :[[@LINE-1]]:{{[0-9]+}}: error: invalid register class: vgpr tuples must be 64 bit aligned
1798

0 commit comments

Comments
 (0)