[llvm-branch-commits] [llvm] [AMDGPU] Make custom AsmParser matchers alignment aware (PR #221988)
Valery Pykhtin via llvm-branch-commits
llvm-branch-commits at lists.llvm.org
Tue Sep 8 06:02:52 PDT 2026
https://github.com/vpykhtin created https://github.com/llvm/llvm-project/pull/221988
Currently custom matchers accept registers belonging to unaligned classes,
leaving the alignment check and its diagnostic to validateVGPRAlign. The problem
is that validateVGPRAlign does not account for the operand register class, which
may have a different alignment requirement on mixed-alignment targets.
The fix is to make a custom matcher resolve the operand's _AlignTarget class via
HwMode and use that resolved class to accept a register. However, doing so alone
would reject a misaligned register with a generic "invalid operand for
instruction" diagnostic, so extra infrastructure conveys the alignment diagnostic
instead: on a class miss, diagnoseRegAlign re-checks the register against the same
class with alignment relaxed (getUnalignedEquivalentRC, new in AMDGPUBaseInfo);
if it fits, the only problem is alignment, so it records
OperandMatchError::VGPRAlignMismatch on the operand. matchAndEmitInstruction
selects that reason and prints the alignment error, now at the offending operand
column instead of column 1.
This commit converts the input-mods and DP-ALU DPP matchers. The no-modifier
reg-or-inline operands come next. validateVGPRAlign stays as a catch-all until
every operand is converted.
Co-Authored-By: Claude <noreply at anthropic.com>
>From 0efcaa1b8aedc3993f49a5a7cfefa203045e43f6 Mon Sep 17 00:00:00 2001
From: Valery Pykhtin <valery.pykhtin at amd.com>
Date: Sat, 29 Aug 2026 10:42:45 +0000
Subject: [PATCH] [AMDGPU] Make custom AsmParser matchers alignment aware
Currently custom matchers accept registers belonging to unaligned classes,
leaving the alignment check and its diagnostic to validateVGPRAlign. The problem
is that validateVGPRAlign does not account for the operand register class, which
may have a different alignment requirement on mixed-alignment targets.
The fix is to make a custom matcher resolve the operand's _AlignTarget class via
HwMode and use that resolved class to accept a register. However, doing so alone
would reject a misaligned register with a generic "invalid operand for
instruction" diagnostic, so extra infrastructure conveys the alignment diagnostic
instead: on a class miss, diagnoseRegAlign re-checks the register against the same
class with alignment relaxed (getUnalignedEquivalentRC, new in AMDGPUBaseInfo);
if it fits, the only problem is alignment, so it records
OperandMatchError::VGPRAlignMismatch on the operand. matchAndEmitInstruction
selects that reason and prints the alignment error, now at the offending operand
column instead of column 1.
This commit converts the input-mods and DP-ALU DPP matchers. The no-modifier
reg-or-inline operands come next. validateVGPRAlign stays as a catch-all until
every operand is converted.
Co-Authored-By: Claude <noreply at anthropic.com>
---
.../AMDGPU/AsmParser/AMDGPUAsmParser.cpp | 107 ++++++++++++++----
.../Target/AMDGPU/Utils/AMDGPUBaseInfo.cpp | 87 ++++++++++++++
llvm/lib/Target/AMDGPU/Utils/AMDGPUBaseInfo.h | 5 +
llvm/test/MC/AMDGPU/gfx1250_asm_vopd_errs.s | 12 +-
.../MC/AMDGPU/misaligned-vgpr-tuples-err.s | 2 +-
5 files changed, 187 insertions(+), 26 deletions(-)
diff --git a/llvm/lib/Target/AMDGPU/AsmParser/AMDGPUAsmParser.cpp b/llvm/lib/Target/AMDGPU/AsmParser/AMDGPUAsmParser.cpp
index f0719a4dabf6b..19b08b4b77a89 100644
--- a/llvm/lib/Target/AMDGPU/AsmParser/AMDGPUAsmParser.cpp
+++ b/llvm/lib/Target/AMDGPU/AsmParser/AMDGPUAsmParser.cpp
@@ -66,13 +66,24 @@ enum RegisterKind {
// Operand
//===----------------------------------------------------------------------===//
+// Why an operand predicate rejected an operand; higher values are more specific
+// and win the tiebreak in matchAndEmitInstruction. Keep None first.
+enum class OperandMatchError {
+ None,
+ VGPRAlignMismatch,
+};
+
class AMDGPUOperand : public MCParsedAsmOperand {
enum KindTy { Token, Immediate, Register, Expression } Kind;
SMLoc StartLoc, EndLoc;
const AMDGPUAsmParser *AsmParser;
+ mutable OperandMatchError MatchError = OperandMatchError::None;
+
public:
+ OperandMatchError getMatchError() const { return MatchError; }
+
AMDGPUOperand(KindTy Kind_, const AMDGPUAsmParser *AsmParser_)
: Kind(Kind_), AsmParser(AsmParser_) {}
@@ -244,10 +255,18 @@ class AMDGPUOperand : public MCParsedAsmOperand {
return isRegClass(RCID) || isInlinableImm(type);
}
+ bool isRegOrInlineByHwMode(unsigned RCByHwModeIdx, MVT type) const {
+ return isRegClassByHwMode(RCByHwModeIdx) || isInlinableImm(type);
+ }
+
bool isRegOrImmWithInputMods(unsigned RCID, MVT type) const {
return isRegOrInline(RCID, type) || isLiteralImm(type);
}
+ bool isRegOrImmWithInputModsByHwMode(unsigned RCByHwModeIdx, MVT type) const {
+ return isRegOrInlineByHwMode(RCByHwModeIdx, type) || isLiteralImm(type);
+ }
+
bool isRegOrImmWithInt16InputMods() const {
return isRegOrImmWithInputMods(AMDGPU::VS_32RegClassID, MVT::i16);
}
@@ -275,7 +294,7 @@ class AMDGPUOperand : public MCParsedAsmOperand {
}
bool isRegOrImmWithInt64InputMods() const {
- return isRegOrImmWithInputMods(AMDGPU::VS_64RegClassID, MVT::i64);
+ return isRegOrImmWithInputModsByHwMode(AMDGPU::VS_64_AlignTarget, MVT::i64);
}
bool isRegOrImmWithFP16InputMods() const {
@@ -292,7 +311,7 @@ class AMDGPUOperand : public MCParsedAsmOperand {
}
bool isRegOrImmWithFP64InputMods() const {
- return isRegOrImmWithInputMods(AMDGPU::VS_64RegClassID, MVT::f64);
+ return isRegOrImmWithInputModsByHwMode(AMDGPU::VS_64_AlignTarget, MVT::f64);
}
template <bool IsFake16> bool isRegOrInlineImmWithFP16InputMods() const {
@@ -305,7 +324,7 @@ class AMDGPUOperand : public MCParsedAsmOperand {
}
bool isRegOrInlineImmWithFP64InputMods() const {
- return isRegOrInline(AMDGPU::VS_64RegClassID, MVT::f64);
+ return isRegOrInlineByHwMode(AMDGPU::VS_64_AlignTarget, MVT::f64);
}
bool isVRegWithInputMods(unsigned RCID) const { return isRegClass(RCID); }
@@ -315,7 +334,7 @@ class AMDGPUOperand : public MCParsedAsmOperand {
}
bool isVRegWithFP64InputMods() const {
- return isVRegWithInputMods(AMDGPU::VReg_64RegClassID);
+ return isRegClassByHwMode(AMDGPU::VReg_64_AlignTarget);
}
bool isPackedFP16InputMods() const {
@@ -417,6 +436,14 @@ class AMDGPUOperand : public MCParsedAsmOperand {
bool isRegClass(unsigned RCID) const;
+ // Check the register against the HwMode-resolved operand class; on failure
+ // also record the alignment diagnostic via diagnoseRegAlign.
+ bool isRegClassByHwMode(unsigned RCByHwModeIdx) const;
+
+ // Record an alignment diagnostic if the register failed operand class RCID
+ // only for being odd-aligned; always returns false.
+ bool diagnoseRegAlign(int16_t RCID) const;
+
bool isInlineValue() const;
bool isRegOrInlineNoMods(unsigned RCID, MVT type) const {
@@ -1613,6 +1640,12 @@ class AMDGPUAsmParser : public MCTargetAsmParser {
const MCInstrInfo *getMII() const { return &MII; }
+ // Resolve a RegClassByHwModeUses index to a register class id for the active
+ // HwMode; -1 if the mode has no entry.
+ int16_t getRegClassByHwMode(unsigned RCByHwModeIdx) const {
+ return MII.getRegClassByHwModeTable(HwMode)[RCByHwModeIdx];
+ }
+
// FIXME: This should not be used. Instead, should use queries derived from
// getAvailableFeatures().
const FeatureBitset &getFeatureBits() const {
@@ -2255,6 +2288,23 @@ bool AMDGPUOperand::isLiteralImm(MVT type) const {
return canLosslesslyConvertToFPType(FPLiteral, ExpectedType);
}
+bool AMDGPUOperand::isRegClassByHwMode(unsigned RCByHwModeIdx) const {
+ if (!isRegKind())
+ return false;
+ int16_t RCID = AsmParser->getRegClassByHwMode(RCByHwModeIdx);
+ // On a class miss diagnoseRegAlign records a misalignment (a no-op on
+ // subtargets without aligned VGPRs); it always returns false.
+ return RCID >= 0 && (isRegClass(RCID) || diagnoseRegAlign(RCID));
+}
+
+bool AMDGPUOperand::diagnoseRegAlign(int16_t RCID) const {
+ const MCRegisterInfo *MRI = AsmParser->getMRI();
+ int UnalignedRCID = AMDGPU::getUnalignedEquivalentRC(RCID);
+ if (UnalignedRCID >= 0 && MRI->getRegClass(UnalignedRCID).contains(getReg()))
+ MatchError = OperandMatchError::VGPRAlignMismatch;
+ return false;
+}
+
bool AMDGPUOperand::isRegClass(unsigned RCID) const {
return isRegKind() &&
AsmParser->getMRI()->getRegClass(RCID).contains(getReg());
@@ -2263,8 +2313,8 @@ bool AMDGPUOperand::isRegClass(unsigned RCID) const {
bool AMDGPUOperand::isVRegWithInputMods() const {
return isRegClass(AMDGPU::VGPR_32RegClassID) ||
// GFX90A allows DPP on 64-bit operands.
- (isRegClass(AMDGPU::VReg_64RegClassID) &&
- AsmParser->getFeatureBits()[AMDGPU::FeatureDPALU_DPP]);
+ (AsmParser->getFeatureBits()[AMDGPU::FeatureDPALU_DPP] &&
+ isRegClassByHwMode(AMDGPU::VReg_64_AlignTarget));
}
template <bool IsFake16>
@@ -5888,25 +5938,32 @@ bool AMDGPUAsmParser::matchAndEmitInstruction(SMLoc IDLoc, unsigned &Opcode,
MCInst Inst;
Inst.setLoc(IDLoc);
unsigned Result = Match_Success;
-
- // Order match statuses from least to most specific and keep the most
- // specific one:
- // Match_MnemonicFail < Match_InvalidOperand < Match_MissingFeature
- auto atLeastAsSpecific = [](unsigned New, unsigned Cur) {
- auto rank = [](unsigned M) {
- return M == Match_MnemonicFail ? 1
- : M == Match_InvalidOperand ? 2
- : M == Match_MissingFeature ? 3
- : 0; // Match_Success sentinel
+ ErrorInfo = ~0ULL; // set by the loop; initialized so it can be read before.
+
+ // Rank a match status as (MatchResultOrder, MatchError):
+ // MnemonicFail < InvalidOperand < MissingFeature (Success lowest), ties
+ // broken by the failing operand's recorded OperandMatchError.
+ auto atLeastAsSpecific = [&](unsigned New, uint64_t NewIdx, unsigned Cur,
+ uint64_t CurIdx) {
+ auto rank = [&](unsigned M, uint64_t I) {
+ int MROrder = M == Match_MnemonicFail ? 1
+ : M == Match_InvalidOperand ? 2
+ : M == Match_MissingFeature ? 3
+ : 0; // Match_Success sentinel
+ OperandMatchError MatchError =
+ M == Match_InvalidOperand && I < Operands.size()
+ ? static_cast<const AMDGPUOperand &>(*Operands[I]).getMatchError()
+ : OperandMatchError::None;
+ return (MROrder << 16) | static_cast<int>(MatchError);
};
- return rank(New) >= rank(Cur);
+ return rank(New, NewIdx) >= rank(Cur, CurIdx);
};
for (auto Variant : getMatchedVariants()) {
uint64_t EI;
auto R =
MatchInstructionImpl(Operands, Inst, EI, MatchingInlineAsm, Variant);
- if (R == Match_Success || atLeastAsSpecific(R, Result)) {
+ if (R == Match_Success || atLeastAsSpecific(R, EI, Result, ErrorInfo)) {
Result = R;
ErrorInfo = EI;
}
@@ -5945,12 +6002,24 @@ bool AMDGPUAsmParser::matchAndEmitInstruction(SMLoc IDLoc, unsigned &Opcode,
if (ErrorInfo >= Operands.size()) {
return Error(IDLoc, "too few operands for instruction");
}
- ErrorLoc = ((AMDGPUOperand &)*Operands[ErrorInfo]).getStartLoc();
+ AMDGPUOperand &ErrorOp = (AMDGPUOperand &)*Operands[ErrorInfo];
+ ErrorLoc = ErrorOp.getStartLoc();
if (ErrorLoc == SMLoc())
ErrorLoc = IDLoc;
if (isInvalidVOPDY(Operands, ErrorInfo))
return Error(ErrorLoc, "invalid VOPDY instruction");
+
+ // A predicate may have recorded a more specific reason for rejecting the
+ // operand than the generic "invalid operand" below.
+ switch (ErrorOp.getMatchError()) {
+ case OperandMatchError::VGPRAlignMismatch:
+ return Error(
+ ErrorLoc,
+ "invalid register class: vgpr tuples must be 64 bit aligned");
+ case OperandMatchError::None:
+ break;
+ }
}
return Error(ErrorLoc, "invalid operand for instruction");
}
diff --git a/llvm/lib/Target/AMDGPU/Utils/AMDGPUBaseInfo.cpp b/llvm/lib/Target/AMDGPU/Utils/AMDGPUBaseInfo.cpp
index 340d569edb199..3a24d943323c1 100644
--- a/llvm/lib/Target/AMDGPU/Utils/AMDGPUBaseInfo.cpp
+++ b/llvm/lib/Target/AMDGPU/Utils/AMDGPUBaseInfo.cpp
@@ -3018,6 +3018,93 @@ unsigned getRegBitWidth(const MCRegisterClass &RC) {
return getRegBitWidth(RC.getID());
}
+int getUnalignedEquivalentRC(unsigned RCID) {
+ switch (RCID) {
+ case AMDGPU::VReg_64_Align2RegClassID:
+ return AMDGPU::VReg_64RegClassID;
+ case AMDGPU::VReg_96_Align2RegClassID:
+ return AMDGPU::VReg_96RegClassID;
+ case AMDGPU::VReg_128_Align2RegClassID:
+ return AMDGPU::VReg_128RegClassID;
+ case AMDGPU::VReg_160_Align2RegClassID:
+ return AMDGPU::VReg_160RegClassID;
+ case AMDGPU::VReg_192_Align2RegClassID:
+ return AMDGPU::VReg_192RegClassID;
+ case AMDGPU::VReg_224_Align2RegClassID:
+ return AMDGPU::VReg_224RegClassID;
+ case AMDGPU::VReg_256_Align2RegClassID:
+ return AMDGPU::VReg_256RegClassID;
+ case AMDGPU::VReg_288_Align2RegClassID:
+ return AMDGPU::VReg_288RegClassID;
+ case AMDGPU::VReg_320_Align2RegClassID:
+ return AMDGPU::VReg_320RegClassID;
+ case AMDGPU::VReg_352_Align2RegClassID:
+ return AMDGPU::VReg_352RegClassID;
+ case AMDGPU::VReg_384_Align2RegClassID:
+ return AMDGPU::VReg_384RegClassID;
+ case AMDGPU::VReg_512_Align2RegClassID:
+ return AMDGPU::VReg_512RegClassID;
+ case AMDGPU::VReg_1024_Align2RegClassID:
+ return AMDGPU::VReg_1024RegClassID;
+ case AMDGPU::AReg_64_Align2RegClassID:
+ return AMDGPU::AReg_64RegClassID;
+ case AMDGPU::AReg_96_Align2RegClassID:
+ return AMDGPU::AReg_96RegClassID;
+ case AMDGPU::AReg_128_Align2RegClassID:
+ return AMDGPU::AReg_128RegClassID;
+ case AMDGPU::AReg_160_Align2RegClassID:
+ return AMDGPU::AReg_160RegClassID;
+ case AMDGPU::AReg_192_Align2RegClassID:
+ return AMDGPU::AReg_192RegClassID;
+ case AMDGPU::AReg_224_Align2RegClassID:
+ return AMDGPU::AReg_224RegClassID;
+ case AMDGPU::AReg_256_Align2RegClassID:
+ return AMDGPU::AReg_256RegClassID;
+ case AMDGPU::AReg_288_Align2RegClassID:
+ return AMDGPU::AReg_288RegClassID;
+ case AMDGPU::AReg_320_Align2RegClassID:
+ return AMDGPU::AReg_320RegClassID;
+ case AMDGPU::AReg_352_Align2RegClassID:
+ return AMDGPU::AReg_352RegClassID;
+ case AMDGPU::AReg_384_Align2RegClassID:
+ return AMDGPU::AReg_384RegClassID;
+ case AMDGPU::AReg_512_Align2RegClassID:
+ return AMDGPU::AReg_512RegClassID;
+ case AMDGPU::AReg_1024_Align2RegClassID:
+ return AMDGPU::AReg_1024RegClassID;
+ case AMDGPU::AV_64_Align2RegClassID:
+ return AMDGPU::AV_64RegClassID;
+ case AMDGPU::AV_96_Align2RegClassID:
+ return AMDGPU::AV_96RegClassID;
+ case AMDGPU::AV_128_Align2RegClassID:
+ return AMDGPU::AV_128RegClassID;
+ case AMDGPU::AV_160_Align2RegClassID:
+ return AMDGPU::AV_160RegClassID;
+ case AMDGPU::AV_192_Align2RegClassID:
+ return AMDGPU::AV_192RegClassID;
+ case AMDGPU::AV_224_Align2RegClassID:
+ return AMDGPU::AV_224RegClassID;
+ case AMDGPU::AV_256_Align2RegClassID:
+ return AMDGPU::AV_256RegClassID;
+ case AMDGPU::AV_288_Align2RegClassID:
+ return AMDGPU::AV_288RegClassID;
+ case AMDGPU::AV_320_Align2RegClassID:
+ return AMDGPU::AV_320RegClassID;
+ case AMDGPU::AV_352_Align2RegClassID:
+ return AMDGPU::AV_352RegClassID;
+ case AMDGPU::AV_384_Align2RegClassID:
+ return AMDGPU::AV_384RegClassID;
+ case AMDGPU::AV_512_Align2RegClassID:
+ return AMDGPU::AV_512RegClassID;
+ case AMDGPU::AV_1024_Align2RegClassID:
+ return AMDGPU::AV_1024RegClassID;
+ case AMDGPU::VS_64_Align2RegClassID:
+ return AMDGPU::VS_64RegClassID;
+ default:
+ return -1;
+ }
+}
+
bool isInlinableLiteral64(int64_t Literal, bool HasInv2Pi) {
if (isInlinableIntLiteral(Literal))
return true;
diff --git a/llvm/lib/Target/AMDGPU/Utils/AMDGPUBaseInfo.h b/llvm/lib/Target/AMDGPU/Utils/AMDGPUBaseInfo.h
index 089b195f11097..e4ae372ea5f0b 100644
--- a/llvm/lib/Target/AMDGPU/Utils/AMDGPUBaseInfo.h
+++ b/llvm/lib/Target/AMDGPU/Utils/AMDGPUBaseInfo.h
@@ -1570,6 +1570,11 @@ unsigned getRegBitWidth(unsigned RCID);
/// Get the size in bits of a register from the register class \p RC.
unsigned getRegBitWidth(const MCRegisterClass &RC);
+/// Return the register class equivalent to the even-aligned VGPR/AGPR/AV tuple
+/// register class \p RCID with the alignment requirement relaxed, or -1 if
+/// \p RCID has no such equivalent.
+int getUnalignedEquivalentRC(unsigned RCID);
+
LLVM_READNONE
inline unsigned getOperandSize(const MCOperandInfo &OpInfo) {
switch (OpInfo.OperandType) {
diff --git a/llvm/test/MC/AMDGPU/gfx1250_asm_vopd_errs.s b/llvm/test/MC/AMDGPU/gfx1250_asm_vopd_errs.s
index 3014207ef2b92..f703510b1e198 100644
--- a/llvm/test/MC/AMDGPU/gfx1250_asm_vopd_errs.s
+++ b/llvm/test/MC/AMDGPU/gfx1250_asm_vopd_errs.s
@@ -330,17 +330,17 @@ v_dual_cndmask_b32 v28, -v15, v15, s46 :: v_dual_cndmask_b32 v29, sext(v13), -v1
//===----------------------------------------------------------------------===//
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: :[[@LINE-1]]:28: 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:{{^}}^
+// 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: :[[@LINE-1]]:36: 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:{{^}}^
+// 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: :[[@LINE-1]]:44: 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:{{^}}^
+// 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 4646517f9f4db..ac9a1c7c65ba8 100644
--- a/llvm/test/MC/AMDGPU/misaligned-vgpr-tuples-err.s
+++ b/llvm/test/MC/AMDGPU/misaligned-vgpr-tuples-err.s
@@ -116,7 +116,7 @@ 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
+// 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
More information about the llvm-branch-commits
mailing list