[llvm] [AMDGPU] Add AsmParser tests for misaligned VGPR tuples in custom-predicate operands (PR #221986)

[Valery Pykhtin via llvm-commits]

https://github.com/vpykhtin created https://github.com/llvm/llvm-project/pull/221986

MFMA VGPR sources, DP-ALU DPP, v_pk_mov_b32 and VOPD3 f64 take VGPR tuple
operands validated by custom AsmParser predicates. Odd-aligned tuples are
currently only caught late by validateVGPRAlign, so the error points at the
whole instruction (column 1). A follow-up moves the diagnostic to the
offending operand.

>From 439676da3dbfae367537a089cfcab302a891d941 Mon Sep 17 00:00:00 2001
From: Valery Pykhtin <valery.pykhtin at amd.com>
Date: Sat, 29 Aug 2026 08:56:47 +0000
Subject: [PATCH] [AMDGPU] Add AsmParser tests for misaligned VGPR tuples in
 custom-predicate operands

MFMA VGPR sources, DP-ALU DPP, v_pk_mov_b32 and VOPD3 f64 take VGPR tuple
operands validated by custom AsmParser predicates. Odd-aligned tuples are
currently only caught late by validateVGPRAlign, so the error points at the
whole instruction (column 1). A follow-up moves the diagnostic to the
offending operand.
---
 llvm/test/MC/AMDGPU/gfx1250_asm_vopd_errs.s   | 20 +++++++++++++++++++
 .../MC/AMDGPU/misaligned-vgpr-tuples-err.s    | 18 +++++++++++++++++
 2 files changed, 38 insertions(+)

diff --git a/llvm/test/MC/AMDGPU/gfx1250_asm_vopd_errs.s b/llvm/test/MC/AMDGPU/gfx1250_asm_vopd_errs.s
index 8d6b6ad93d942..3014207ef2b92 100644
--- a/llvm/test/MC/AMDGPU/gfx1250_asm_vopd_errs.s
+++ b/llvm/test/MC/AMDGPU/gfx1250_asm_vopd_errs.s
@@ -324,3 +324,23 @@ v_dual_cndmask_b32 v28, -v15, v15, s46 :: v_dual_cndmask_b32 v29, sext(v13), -v1
 // GFX12: :[[@LINE-1]]:{{[0-9]+}}: error: not a valid operand
 // GFX12-NEXT:{{^}}v_dual_cndmask_b32 v28, -v15, v15, s46 :: v_dual_cndmask_b32 v29, sext(v13), -v13, s46
 // GFX12-NEXT:{{^}}                                                                  ^
+
+//===----------------------------------------------------------------------===//
+// VOPD3 f64 sources must be 64-bit aligned (FeatureRequiresAlignedVGPRs).
+//===----------------------------------------------------------------------===//
+
+v_dual_fma_f64 v[252:253], v[7:8], v[4:5], v[10:11] :: v_dual_add_f32 v8, v1, v3
+// GFX12: :[[@LINE-1]]:1: error: invalid register class: vgpr tuples must be 64 bit aligned
+// GFX12-NEXT:{{^}}v_dual_fma_f64 v[252:253], v[7:8], v[4:5], v[10:11] :: v_dual_add_f32 v8, v1, v3
+// GFX12-NEXT:{{^}}^
+
+v_dual_fma_f64 v[252:253], v[6:7], v[5:6], v[10:11] :: v_dual_add_f32 v8, v1, v3
+// GFX12: :[[@LINE-1]]:1: error: invalid register class: vgpr tuples must be 64 bit aligned
+// GFX12-NEXT:{{^}}v_dual_fma_f64 v[252:253], v[6:7], v[5:6], v[10:11] :: v_dual_add_f32 v8, v1, v3
+// GFX12-NEXT:{{^}}^
+
+v_dual_fma_f64 v[252:253], v[6:7], v[4:5], v[11:12] :: v_dual_add_f32 v8, v1, v3
+// GFX12: :[[@LINE-1]]:1: error: invalid register class: vgpr tuples must be 64 bit aligned
+// GFX12-NEXT:{{^}}v_dual_fma_f64 v[252:253], v[6:7], v[4:5], v[11:12] :: v_dual_add_f32 v8, v1, v3
+// GFX12-NEXT:{{^}}^
+
diff --git a/llvm/test/MC/AMDGPU/misaligned-vgpr-tuples-err.s b/llvm/test/MC/AMDGPU/misaligned-vgpr-tuples-err.s
index 4e2a41819ea54..4646517f9f4db 100644
--- a/llvm/test/MC/AMDGPU/misaligned-vgpr-tuples-err.s
+++ b/llvm/test/MC/AMDGPU/misaligned-vgpr-tuples-err.s
@@ -105,3 +105,21 @@ v_mfma_f32_32x32x8f16 a[0:15], a[1:2], v[0:1], a[0:15]
 
 v_mfma_i32_4x4x4i8 a[1:4], a0, v1, 2
 // GFX90A: :[[@LINE-1]]:20: error: invalid operand for instruction
+
+v_mfma_f32_4x4x1f32 v[0:3], v0, v1, v[5:8]
+// GFX90A: :[[@LINE-1]]:1: error: invalid register class: vgpr tuples must be 64 bit aligned
+
+v_mfma_f32_16x16x1f32 v[0:15], v0, v1, v[17:32]
+// GFX90A: :[[@LINE-1]]:1: error: invalid register class: vgpr tuples must be 64 bit aligned
+
+v_mfma_f32_32x32x1f32 v[0:31], v0, v1, v[33:64]
+// GFX90A: :[[@LINE-1]]:1: error: invalid register class: vgpr tuples must be 64 bit aligned
+
+v_ceil_f64_dpp v[0:1], v[3:4] row_newbcast:1 row_mask:0xf bank_mask:0xf
+// GFX90A: :[[@LINE-1]]:1: error: invalid register class: vgpr tuples must be 64 bit aligned
+
+v_pk_mov_b32 v[0:1], v[3:4], v[4:5]
+// GFX90A: :[[@LINE-1]]:1: error: invalid register class: vgpr tuples must be 64 bit aligned
+
+v_pk_mov_b32 v[0:1], v[2:3], v[5:6]
+// GFX90A: :[[@LINE-1]]:1: error: invalid register class: vgpr tuples must be 64 bit aligned