[llvm] AMDGPU: Test a few more cases for assembler errors for misaligned gfx90a vgprs (PR #156998)

[Matt Arsenault via llvm-commits]

https://github.com/arsenm updated https://github.com/llvm/llvm-project/pull/156998

>From e870b8ea2358be4fc056223b7cef46516a9989c9 Mon Sep 17 00:00:00 2001
From: Matt Arsenault <Matthew.Arsenault at amd.com>
Date: Fri, 5 Sep 2025 09:52:52 +0900
Subject: [PATCH] AMDGPU: Test a few more cases for assembler errors for
 misaligned gfx90a vgprs

---
 .../MC/AMDGPU/misaligned-vgpr-tuples-err.s    | 81 +++++++++++++++++++
 1 file changed, 81 insertions(+)

diff --git a/llvm/test/MC/AMDGPU/misaligned-vgpr-tuples-err.s b/llvm/test/MC/AMDGPU/misaligned-vgpr-tuples-err.s
index d76dc8c9fff63..7834fb5372ec7 100644
--- a/llvm/test/MC/AMDGPU/misaligned-vgpr-tuples-err.s
+++ b/llvm/test/MC/AMDGPU/misaligned-vgpr-tuples-err.s
@@ -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