Skip to content

AMDGPU: Test a few more cases for assembler errors for misaligned gfx90a vgprs #156998

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Merged
merged 1 commit into from
Sep 5, 2025
Merged

AMDGPU: Test a few more cases for assembler errors for misaligned gfx90a vgprs #156998

Changes from all commits
Commits
Show all changes
1 commit
Select commit
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
81 changes: 81 additions & 0 deletions llvm/test/MC/AMDGPU/misaligned-vgpr-tuples-err.s
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -3,15 +3,96 @@
v_add_f64 v[1:2], v[1:2], v[1:2]
// GFX90A: :[[@LINE-1]]:{{[0-9]+}}: error: invalid register class: vgpr tuples must be 64 bit aligned

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

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

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

global_load_dwordx2 a[1:2], v[0:1], off
// GFX90A: :[[@LINE-1]]:{{[0-9]+}}: error: invalid register class: vgpr tuples must be 64 bit aligned

global_load_dwordx3 a[1:3], v[0:1], off
// GFX90A: :[[@LINE-1]]:{{[0-9]+}}: error: invalid register class: vgpr tuples must be 64 bit aligned

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


image_load v[1:2], v2, s[0:7] dmask:0x3 unorm
// GFX90A: :[[@LINE-1]]:{{[0-9]+}}: error: invalid register class: vgpr tuples must be 64 bit aligned

image_load v[1:3], v2, s[0:7] dmask:0x7 unorm
// GFX90A: :[[@LINE-1]]:{{[0-9]+}}: error: invalid register class: vgpr tuples must be 64 bit aligned

image_load v[1:4], v2, s[0:7] dmask:0xf unorm
// GFX90A: :[[@LINE-1]]:{{[0-9]+}}: error: invalid register class: vgpr tuples must be 64 bit aligned

image_load a[1:2], v2, s[0:7] dmask:0x3 unorm
// GFX90A: :[[@LINE-1]]:{{[0-9]+}}: error: invalid register class: vgpr tuples must be 64 bit aligned

image_load a[1:3], v2, s[0:7] dmask:0x7 unorm
// GFX90A: :[[@LINE-1]]:{{[0-9]+}}: error: invalid register class: vgpr tuples must be 64 bit aligned

image_load a[1:4], v2, s[0:7] dmask:0xf unorm
// GFX90A: :[[@LINE-1]]:{{[0-9]+}}: error: invalid register class: vgpr tuples must be 64 bit aligned


image_store v[193:194], v[238:241], s[28:35] dmask:0x3 unorm
// GFX90A: :[[@LINE-1]]:{{[0-9]+}}: error: invalid register class: vgpr tuples must be 64 bit aligned

image_store v[193:195], v[238:241], s[28:35] dmask:0x7 unorm
// GFX90A: :[[@LINE-1]]:{{[0-9]+}}: error: invalid register class: vgpr tuples must be 64 bit aligned

image_store v[193:196], v[238:241], s[28:35] dmask:0xf unorm
// GFX90A: :[[@LINE-1]]:{{[0-9]+}}: error: invalid register class: vgpr tuples must be 64 bit aligned

image_store a[193:194], v[238:241], s[28:35] dmask:0x3 unorm
// GFX90A: :[[@LINE-1]]:{{[0-9]+}}: error: invalid register class: vgpr tuples must be 64 bit aligned

image_store a[193:195], v[238:241], s[28:35] dmask:0x7 unorm
// GFX90A: :[[@LINE-1]]:{{[0-9]+}}: error: invalid register class: vgpr tuples must be 64 bit aligned

image_store a[193:196], v[238:241], s[28:35] dmask:0xf unorm
// GFX90A: :[[@LINE-1]]:{{[0-9]+}}: error: invalid register class: vgpr tuples must be 64 bit aligned


image_atomic_swap v4, v[193:196], s[28:35] dmask:0x1 unorm glc
// GFX90A: :[[@LINE-1]]:{{[0-9]+}}: error: invalid register class: vgpr tuples must be 64 bit aligned

image_atomic_swap v[5:6], v1, s[8:15] dmask:0x3 unorm
// GFX90A: :[[@LINE-1]]:{{[0-9]+}}: error: invalid register class: vgpr tuples must be 64 bit aligned


image_atomic_cmpswap v[5:6], v[192:195], s[28:35] dmask:0x3 unorm glc
// GFX90A: :[[@LINE-1]]:{{[0-9]+}}: error: invalid register class: vgpr tuples must be 64 bit aligned

image_atomic_cmpswap v[4:5], v[193:196], s[28:35] dmask:0x3 unorm glc
// GFX90A: :[[@LINE-1]]:{{[0-9]+}}: error: invalid register class: vgpr tuples must be 64 bit aligned

image_atomic_cmpswap v[5:8], v[192:195], s[28:35] dmask:0xf unorm glc
// GFX90A: :[[@LINE-1]]:{{[0-9]+}}: error: invalid register class: vgpr tuples must be 64 bit aligned

image_atomic_cmpswap v[4:7], v[193:196], s[28:35] dmask:0xf unorm glc
// GFX90A: :[[@LINE-1]]:{{[0-9]+}}: error: invalid register class: vgpr tuples must be 64 bit aligned


image_atomic_cmpswap a[5:6], v[192:195], s[28:35] dmask:0x3 unorm glc
// GFX90A: :[[@LINE-1]]:{{[0-9]+}}: error: invalid register class: vgpr tuples must be 64 bit aligned

image_atomic_cmpswap a[4:5], v[193:196], s[28:35] dmask:0x3 unorm glc
// GFX90A: :[[@LINE-1]]:{{[0-9]+}}: error: invalid register class: vgpr tuples must be 64 bit aligned

image_atomic_cmpswap a[5:8], v[192:195], s[28:35] dmask:0xf unorm glc
// GFX90A: :[[@LINE-1]]:{{[0-9]+}}: error: invalid register class: vgpr tuples must be 64 bit aligned

image_atomic_cmpswap a[4:7], v[193:196], s[28:35] dmask:0xf unorm glc
// GFX90A: :[[@LINE-1]]:{{[0-9]+}}: error: invalid register class: vgpr tuples must be 64 bit aligned


v_mfma_f32_32x32x8f16 a[0:15], a[1:2], v[0:1], a[0:15]
// GFX90A: :[[@LINE-1]]:{{[0-9]+}}: error: invalid register class: vgpr tuples must be 64 bit aligned

Expand Down