[llvm-branch-commits] [llvm] [AMDGPU] Route no-modifier reg-or-inline AsmParser operands through HwMode predicate (PR #221989)
Valery Pykhtin via llvm-branch-commits
llvm-branch-commits at lists.llvm.org
Tue Sep 8 06:04:24 PDT 2026
https://github.com/vpykhtin created https://github.com/llvm/llvm-project/pull/221989
Convert the reg-or-inline operands with no modifiers (MFMA VGPR/AGPR
sources, VCSrc, v_pk_mov_b32, VOP scalar f64) from the fixed-class
isRegOrInlineNoMods to the HwMode-aware isRegOrInlineNoModsByHwMode, so an
odd-aligned tuple is rejected at the offending operand column instead of by
the validateVGPRAlign catch-all.
Co-Authored-By: Claude <noreply at anthropic.com>
>From a943ffe28b415f2cb4f73c5812da4093f35f7d2b Mon Sep 17 00:00:00 2001
From: Valery Pykhtin <valery.pykhtin at amd.com>
Date: Mon, 31 Aug 2026 18:48:14 +0000
Subject: [PATCH] [AMDGPU] Route no-modifier reg-or-inline AsmParser operands
through HwMode predicate
Convert the reg-or-inline operands with no modifiers (MFMA VGPR/AGPR
sources, VCSrc, v_pk_mov_b32, VOP scalar f64) from the fixed-class
isRegOrInlineNoMods to the HwMode-aware isRegOrInlineNoModsByHwMode, so an
odd-aligned tuple is rejected at the offending operand column instead of by
the validateVGPRAlign catch-all.
Co-Authored-By: Claude <noreply at anthropic.com>
---
.../AMDGPU/AsmParser/AMDGPUAsmParser.cpp | 99 ++++++++++---------
.../MC/AMDGPU/misaligned-vgpr-tuples-err.s | 14 +--
2 files changed, 60 insertions(+), 53 deletions(-)
diff --git a/llvm/lib/Target/AMDGPU/AsmParser/AMDGPUAsmParser.cpp b/llvm/lib/Target/AMDGPU/AsmParser/AMDGPUAsmParser.cpp
index 19b08b4b77a89..db024282e4f9d 100644
--- a/llvm/lib/Target/AMDGPU/AsmParser/AMDGPUAsmParser.cpp
+++ b/llvm/lib/Target/AMDGPU/AsmParser/AMDGPUAsmParser.cpp
@@ -449,6 +449,10 @@ class AMDGPUOperand : public MCParsedAsmOperand {
bool isRegOrInlineNoMods(unsigned RCID, MVT type) const {
return isRegOrInline(RCID, type) && !hasModifiers();
}
+
+ bool isRegOrInlineNoModsByHwMode(unsigned RCByHwModeIdx, MVT type) const {
+ return isRegOrInlineByHwMode(RCByHwModeIdx, type) && !hasModifiers();
+ }
bool isSCSrcB16() const {
return isRegOrInlineNoMods(AMDGPU::SReg_32RegClassID, MVT::i16);
@@ -553,7 +557,7 @@ class AMDGPUOperand : public MCParsedAsmOperand {
}
bool isVCSrc_b64() const {
- return isRegOrInlineNoMods(AMDGPU::VS_64RegClassID, MVT::i64);
+ return isRegOrInlineNoModsByHwMode(AMDGPU::VS_64_AlignTarget, MVT::i64);
}
bool isVCSrcT_b16() const {
@@ -579,7 +583,7 @@ class AMDGPUOperand : public MCParsedAsmOperand {
}
bool isVCSrc_f64() const {
- return isRegOrInlineNoMods(AMDGPU::VS_64RegClassID, MVT::f64);
+ return isRegOrInlineNoModsByHwMode(AMDGPU::VS_64_AlignTarget, MVT::f64);
}
bool isVCSrcTBF16() const {
@@ -664,7 +668,10 @@ class AMDGPUOperand : public MCParsedAsmOperand {
return isVCSrc_f32() || isLiteralImm(MVT::f32) || isExpr();
}
- bool isVSrc_f64() const { return isVCSrc_f64() || isLiteralImm(MVT::f64); }
+ bool isVSrc_f64() const {
+ return isRegOrInlineNoModsByHwMode(AMDGPU::VS_64_AlignTarget, MVT::f64) ||
+ isLiteralImm(MVT::f64);
+ }
bool isVSrcT_bf16() const {
return isVCSrcTBF16() || isLiteralImm(MVT::bf16);
@@ -723,91 +730,91 @@ class AMDGPUOperand : public MCParsedAsmOperand {
bool isVISrcV2F16() const { return isVISrcF16() || isVISrcB32(); }
bool isVISrc_64_bf16() const {
- return isRegOrInlineNoMods(AMDGPU::VReg_64RegClassID, MVT::bf16);
+ return isRegOrInlineNoModsByHwMode(AMDGPU::VReg_64_AlignTarget, MVT::bf16);
}
bool isVISrc_64_f16() const {
- return isRegOrInlineNoMods(AMDGPU::VReg_64RegClassID, MVT::f16);
+ return isRegOrInlineNoModsByHwMode(AMDGPU::VReg_64_AlignTarget, MVT::f16);
}
bool isVISrc_64_b32() const {
- return isRegOrInlineNoMods(AMDGPU::VReg_64RegClassID, MVT::i32);
+ return isRegOrInlineNoModsByHwMode(AMDGPU::VReg_64_AlignTarget, MVT::i32);
}
bool isVISrc_64B64() const {
- return isRegOrInlineNoMods(AMDGPU::VReg_64RegClassID, MVT::i64);
+ return isRegOrInlineNoModsByHwMode(AMDGPU::VReg_64_AlignTarget, MVT::i64);
}
bool isVISrc_64_f64() const {
- return isRegOrInlineNoMods(AMDGPU::VReg_64RegClassID, MVT::f64);
+ return isRegOrInlineNoModsByHwMode(AMDGPU::VReg_64_AlignTarget, MVT::f64);
}
bool isVISrc_64V2FP32() const {
- return isRegOrInlineNoMods(AMDGPU::VReg_64RegClassID, MVT::f32);
+ return isRegOrInlineNoModsByHwMode(AMDGPU::VReg_64_AlignTarget, MVT::f32);
}
bool isVISrc_64V2INT32() const {
- return isRegOrInlineNoMods(AMDGPU::VReg_64RegClassID, MVT::i32);
+ return isRegOrInlineNoModsByHwMode(AMDGPU::VReg_64_AlignTarget, MVT::i32);
}
bool isVISrc_256_b32() const {
- return isRegOrInlineNoMods(AMDGPU::VReg_256RegClassID, MVT::i32);
+ return isRegOrInlineNoModsByHwMode(AMDGPU::VReg_256_AlignTarget, MVT::i32);
}
bool isVISrc_256_f32() const {
- return isRegOrInlineNoMods(AMDGPU::VReg_256RegClassID, MVT::f32);
+ return isRegOrInlineNoModsByHwMode(AMDGPU::VReg_256_AlignTarget, MVT::f32);
}
bool isVISrc_256B64() const {
- return isRegOrInlineNoMods(AMDGPU::VReg_256RegClassID, MVT::i64);
+ return isRegOrInlineNoModsByHwMode(AMDGPU::VReg_256_AlignTarget, MVT::i64);
}
bool isVISrc_256_f64() const {
- return isRegOrInlineNoMods(AMDGPU::VReg_256RegClassID, MVT::f64);
+ return isRegOrInlineNoModsByHwMode(AMDGPU::VReg_256_AlignTarget, MVT::f64);
}
bool isVISrc_512_f64() const {
- return isRegOrInlineNoMods(AMDGPU::VReg_512RegClassID, MVT::f64);
+ return isRegOrInlineNoModsByHwMode(AMDGPU::VReg_512_AlignTarget, MVT::f64);
}
bool isVISrc_128B16() const {
- return isRegOrInlineNoMods(AMDGPU::VReg_128RegClassID, MVT::i16);
+ return isRegOrInlineNoModsByHwMode(AMDGPU::VReg_128_AlignTarget, MVT::i16);
}
bool isVISrc_128V2B16() const { return isVISrc_128B16(); }
bool isVISrc_128_b32() const {
- return isRegOrInlineNoMods(AMDGPU::VReg_128RegClassID, MVT::i32);
+ return isRegOrInlineNoModsByHwMode(AMDGPU::VReg_128_AlignTarget, MVT::i32);
}
bool isVISrc_128_f32() const {
- return isRegOrInlineNoMods(AMDGPU::VReg_128RegClassID, MVT::f32);
+ return isRegOrInlineNoModsByHwMode(AMDGPU::VReg_128_AlignTarget, MVT::f32);
}
bool isVISrc_256V2FP32() const {
- return isRegOrInlineNoMods(AMDGPU::VReg_256RegClassID, MVT::f32);
+ return isRegOrInlineNoModsByHwMode(AMDGPU::VReg_256_AlignTarget, MVT::f32);
}
bool isVISrc_256V2INT32() const {
- return isRegOrInlineNoMods(AMDGPU::VReg_256RegClassID, MVT::i32);
+ return isRegOrInlineNoModsByHwMode(AMDGPU::VReg_256_AlignTarget, MVT::i32);
}
bool isVISrc_512_b32() const {
- return isRegOrInlineNoMods(AMDGPU::VReg_512RegClassID, MVT::i32);
+ return isRegOrInlineNoModsByHwMode(AMDGPU::VReg_512_AlignTarget, MVT::i32);
}
bool isVISrc_512B16() const {
- return isRegOrInlineNoMods(AMDGPU::VReg_512RegClassID, MVT::i16);
+ return isRegOrInlineNoModsByHwMode(AMDGPU::VReg_512_AlignTarget, MVT::i16);
}
bool isVISrc_512V2B16() const { return isVISrc_512B16(); }
bool isVISrc_512_f32() const {
- return isRegOrInlineNoMods(AMDGPU::VReg_512RegClassID, MVT::f32);
+ return isRegOrInlineNoModsByHwMode(AMDGPU::VReg_512_AlignTarget, MVT::f32);
}
bool isVISrc_512F16() const {
- return isRegOrInlineNoMods(AMDGPU::VReg_512RegClassID, MVT::f16);
+ return isRegOrInlineNoModsByHwMode(AMDGPU::VReg_512_AlignTarget, MVT::f16);
}
bool isVISrc_512V2F16() const {
@@ -815,21 +822,21 @@ class AMDGPUOperand : public MCParsedAsmOperand {
}
bool isVISrc_1024_b32() const {
- return isRegOrInlineNoMods(AMDGPU::VReg_1024RegClassID, MVT::i32);
+ return isRegOrInlineNoModsByHwMode(AMDGPU::VReg_1024_AlignTarget, MVT::i32);
}
bool isVISrc_1024B16() const {
- return isRegOrInlineNoMods(AMDGPU::VReg_1024RegClassID, MVT::i16);
+ return isRegOrInlineNoModsByHwMode(AMDGPU::VReg_1024_AlignTarget, MVT::i16);
}
bool isVISrc_1024V2B16() const { return isVISrc_1024B16(); }
bool isVISrc_1024_f32() const {
- return isRegOrInlineNoMods(AMDGPU::VReg_1024RegClassID, MVT::f32);
+ return isRegOrInlineNoModsByHwMode(AMDGPU::VReg_1024_AlignTarget, MVT::f32);
}
bool isVISrc_1024F16() const {
- return isRegOrInlineNoMods(AMDGPU::VReg_1024RegClassID, MVT::f16);
+ return isRegOrInlineNoModsByHwMode(AMDGPU::VReg_1024_AlignTarget, MVT::f16);
}
bool isVISrc_1024V2F16() const {
@@ -857,29 +864,29 @@ class AMDGPUOperand : public MCParsedAsmOperand {
bool isAISrcV2F16() const { return isAISrcF16() || isAISrcB32(); }
bool isAISrc_64B64() const {
- return isRegOrInlineNoMods(AMDGPU::AReg_64RegClassID, MVT::i64);
+ return isRegOrInlineNoModsByHwMode(AMDGPU::AReg_64_AlignTarget, MVT::i64);
}
bool isAISrc_64_f64() const {
- return isRegOrInlineNoMods(AMDGPU::AReg_64RegClassID, MVT::f64);
+ return isRegOrInlineNoModsByHwMode(AMDGPU::AReg_64_AlignTarget, MVT::f64);
}
bool isAISrc_128_b32() const {
- return isRegOrInlineNoMods(AMDGPU::AReg_128RegClassID, MVT::i32);
+ return isRegOrInlineNoModsByHwMode(AMDGPU::AReg_128_AlignTarget, MVT::i32);
}
bool isAISrc_128B16() const {
- return isRegOrInlineNoMods(AMDGPU::AReg_128RegClassID, MVT::i16);
+ return isRegOrInlineNoModsByHwMode(AMDGPU::AReg_128_AlignTarget, MVT::i16);
}
bool isAISrc_128V2B16() const { return isAISrc_128B16(); }
bool isAISrc_128_f32() const {
- return isRegOrInlineNoMods(AMDGPU::AReg_128RegClassID, MVT::f32);
+ return isRegOrInlineNoModsByHwMode(AMDGPU::AReg_128_AlignTarget, MVT::f32);
}
bool isAISrc_128F16() const {
- return isRegOrInlineNoMods(AMDGPU::AReg_128RegClassID, MVT::f16);
+ return isRegOrInlineNoModsByHwMode(AMDGPU::AReg_128_AlignTarget, MVT::f16);
}
bool isAISrc_128V2F16() const {
@@ -887,11 +894,11 @@ class AMDGPUOperand : public MCParsedAsmOperand {
}
bool isVISrc_128_bf16() const {
- return isRegOrInlineNoMods(AMDGPU::VReg_128RegClassID, MVT::bf16);
+ return isRegOrInlineNoModsByHwMode(AMDGPU::VReg_128_AlignTarget, MVT::bf16);
}
bool isVISrc_128_f16() const {
- return isRegOrInlineNoMods(AMDGPU::VReg_128RegClassID, MVT::f16);
+ return isRegOrInlineNoModsByHwMode(AMDGPU::VReg_128_AlignTarget, MVT::f16);
}
bool isVISrc_128V2F16() const {
@@ -899,29 +906,29 @@ class AMDGPUOperand : public MCParsedAsmOperand {
}
bool isAISrc_256B64() const {
- return isRegOrInlineNoMods(AMDGPU::AReg_256RegClassID, MVT::i64);
+ return isRegOrInlineNoModsByHwMode(AMDGPU::AReg_256_AlignTarget, MVT::i64);
}
bool isAISrc_256_f64() const {
- return isRegOrInlineNoMods(AMDGPU::AReg_256RegClassID, MVT::f64);
+ return isRegOrInlineNoModsByHwMode(AMDGPU::AReg_256_AlignTarget, MVT::f64);
}
bool isAISrc_512_b32() const {
- return isRegOrInlineNoMods(AMDGPU::AReg_512RegClassID, MVT::i32);
+ return isRegOrInlineNoModsByHwMode(AMDGPU::AReg_512_AlignTarget, MVT::i32);
}
bool isAISrc_512B16() const {
- return isRegOrInlineNoMods(AMDGPU::AReg_512RegClassID, MVT::i16);
+ return isRegOrInlineNoModsByHwMode(AMDGPU::AReg_512_AlignTarget, MVT::i16);
}
bool isAISrc_512V2B16() const { return isAISrc_512B16(); }
bool isAISrc_512_f32() const {
- return isRegOrInlineNoMods(AMDGPU::AReg_512RegClassID, MVT::f32);
+ return isRegOrInlineNoModsByHwMode(AMDGPU::AReg_512_AlignTarget, MVT::f32);
}
bool isAISrc_512F16() const {
- return isRegOrInlineNoMods(AMDGPU::AReg_512RegClassID, MVT::f16);
+ return isRegOrInlineNoModsByHwMode(AMDGPU::AReg_512_AlignTarget, MVT::f16);
}
bool isAISrc_512V2F16() const {
@@ -929,21 +936,21 @@ class AMDGPUOperand : public MCParsedAsmOperand {
}
bool isAISrc_1024_b32() const {
- return isRegOrInlineNoMods(AMDGPU::AReg_1024RegClassID, MVT::i32);
+ return isRegOrInlineNoModsByHwMode(AMDGPU::AReg_1024_AlignTarget, MVT::i32);
}
bool isAISrc_1024B16() const {
- return isRegOrInlineNoMods(AMDGPU::AReg_1024RegClassID, MVT::i16);
+ return isRegOrInlineNoModsByHwMode(AMDGPU::AReg_1024_AlignTarget, MVT::i16);
}
bool isAISrc_1024V2B16() const { return isAISrc_1024B16(); }
bool isAISrc_1024_f32() const {
- return isRegOrInlineNoMods(AMDGPU::AReg_1024RegClassID, MVT::f32);
+ return isRegOrInlineNoModsByHwMode(AMDGPU::AReg_1024_AlignTarget, MVT::f32);
}
bool isAISrc_1024F16() const {
- return isRegOrInlineNoMods(AMDGPU::AReg_1024RegClassID, MVT::f16);
+ return isRegOrInlineNoModsByHwMode(AMDGPU::AReg_1024_AlignTarget, MVT::f16);
}
bool isAISrc_1024V2F16() const {
diff --git a/llvm/test/MC/AMDGPU/misaligned-vgpr-tuples-err.s b/llvm/test/MC/AMDGPU/misaligned-vgpr-tuples-err.s
index ac9a1c7c65ba8..07b196c4b3542 100644
--- a/llvm/test/MC/AMDGPU/misaligned-vgpr-tuples-err.s
+++ b/llvm/test/MC/AMDGPU/misaligned-vgpr-tuples-err.s
@@ -94,10 +94,10 @@ v_add_f64 v[1:2], v[1:2], v[1:2]
// GFX90A: :[[@LINE-1]]:11: error: invalid operand for instruction
v_mfma_f32_16x16x1f32 a[0:15], a0, v1, a[17:32]
-// GFX90A: :[[@LINE-1]]:1: error: invalid register class: vgpr tuples must be 64 bit aligned
+// GFX90A: :[[@LINE-1]]:40: error: invalid register class: vgpr tuples must be 64 bit aligned
v_mfma_f32_32x32x1f32 a[0:31], v0, v1, a[33:64]
-// GFX90A: :[[@LINE-1]]:1: error: invalid register class: vgpr tuples must be 64 bit aligned
+// GFX90A: :[[@LINE-1]]:40: 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]
@@ -107,19 +107,19 @@ 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
+// GFX90A: :[[@LINE-1]]:37: 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
+// GFX90A: :[[@LINE-1]]:40: 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
+// GFX90A: :[[@LINE-1]]:40: 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]]:24: 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
+// GFX90A: :[[@LINE-1]]:22: 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
+// GFX90A: :[[@LINE-1]]:30: error: invalid register class: vgpr tuples must be 64 bit aligned
More information about the llvm-branch-commits
mailing list