[llvm] [AMDGPU] Make custom AsmParser matchers alignment aware (PR #221988)
Valery Pykhtin via llvm-commits
llvm-commits at lists.llvm.org
Wed Sep 9 04:45:46 PDT 2026
https://github.com/vpykhtin updated https://github.com/llvm/llvm-project/pull/221988
>From 66219cb342d2e2cbf9a164defeae05bcf399d243 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 1/3] [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 50733d573b623..21a5f9fe71aac 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 {
@@ -1615,6 +1642,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 {
@@ -2258,6 +2291,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());
@@ -2266,8 +2316,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>
@@ -5894,25 +5944,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;
}
@@ -5951,12 +6008,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 0a6e98777953c..96e21d51a6da7 100644
--- a/llvm/lib/Target/AMDGPU/Utils/AMDGPUBaseInfo.cpp
+++ b/llvm/lib/Target/AMDGPU/Utils/AMDGPUBaseInfo.cpp
@@ -2949,6 +2949,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 44661a99da3e1..29eeb8854f524 100644
--- a/llvm/lib/Target/AMDGPU/Utils/AMDGPUBaseInfo.h
+++ b/llvm/lib/Target/AMDGPU/Utils/AMDGPUBaseInfo.h
@@ -1561,6 +1561,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
>From e5dceee9499efc4cf7d3f2d67e0ac1b934885268 Mon Sep 17 00:00:00 2001
From: Valery Pykhtin <valery.pykhtin at amd.com>
Date: Tue, 8 Sep 2026 16:58:27 +0000
Subject: [PATCH 2/3] [AMDGPU] Drop the misaligned VGPR-tuple AsmParser
diagnostic
The alignment-aware matchers already reject a misaligned tuple, so a
misaligned operand now fails as a plain invalid operand. Keeping the
dedicated "vgpr tuples must be 64 bit aligned" message only for that
case is not worth the extra match-error ranking, the per-operand
diagnostic, and getUnalignedEquivalentRC, so remove all of it.
---
.../AMDGPU/AsmParser/AMDGPUAsmParser.cpp | 72 +++------------
.../Target/AMDGPU/Utils/AMDGPUBaseInfo.cpp | 87 -------------------
llvm/lib/Target/AMDGPU/Utils/AMDGPUBaseInfo.h | 5 --
llvm/test/MC/AMDGPU/gfx1250_asm_vopd_errs.s | 6 +-
.../MC/AMDGPU/misaligned-vgpr-tuples-err.s | 2 +-
5 files changed, 18 insertions(+), 154 deletions(-)
diff --git a/llvm/lib/Target/AMDGPU/AsmParser/AMDGPUAsmParser.cpp b/llvm/lib/Target/AMDGPU/AsmParser/AMDGPUAsmParser.cpp
index 21a5f9fe71aac..d75ac36f4f623 100644
--- a/llvm/lib/Target/AMDGPU/AsmParser/AMDGPUAsmParser.cpp
+++ b/llvm/lib/Target/AMDGPU/AsmParser/AMDGPUAsmParser.cpp
@@ -66,24 +66,13 @@ 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_) {}
@@ -436,14 +425,9 @@ 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.
+ // Check the register against the HwMode-resolved operand class.
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 {
@@ -2295,17 +2279,7 @@ 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;
+ return RCID >= 0 && isRegClass(RCID);
}
bool AMDGPUOperand::isRegClass(unsigned RCID) const {
@@ -5944,32 +5918,25 @@ bool AMDGPUAsmParser::matchAndEmitInstruction(SMLoc IDLoc, unsigned &Opcode,
MCInst Inst;
Inst.setLoc(IDLoc);
unsigned Result = Match_Success;
- 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);
+
+ // 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
};
- return rank(New, NewIdx) >= rank(Cur, CurIdx);
+ return rank(New) >= rank(Cur);
};
for (auto Variant : getMatchedVariants()) {
uint64_t EI;
auto R =
MatchInstructionImpl(Operands, Inst, EI, MatchingInlineAsm, Variant);
- if (R == Match_Success || atLeastAsSpecific(R, EI, Result, ErrorInfo)) {
+ if (R == Match_Success || atLeastAsSpecific(R, Result)) {
Result = R;
ErrorInfo = EI;
}
@@ -6015,17 +5982,6 @@ bool AMDGPUAsmParser::matchAndEmitInstruction(SMLoc IDLoc, unsigned &Opcode,
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 96e21d51a6da7..0a6e98777953c 100644
--- a/llvm/lib/Target/AMDGPU/Utils/AMDGPUBaseInfo.cpp
+++ b/llvm/lib/Target/AMDGPU/Utils/AMDGPUBaseInfo.cpp
@@ -2949,93 +2949,6 @@ 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 29eeb8854f524..44661a99da3e1 100644
--- a/llvm/lib/Target/AMDGPU/Utils/AMDGPUBaseInfo.h
+++ b/llvm/lib/Target/AMDGPU/Utils/AMDGPUBaseInfo.h
@@ -1561,11 +1561,6 @@ 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 f703510b1e198..fd118610727cf 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]]:28: error: invalid register class: vgpr tuples must be 64 bit aligned
+// GFX12: :[[@LINE-1]]:28: error: invalid operand for instruction
// 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]]:36: error: invalid register class: vgpr tuples must be 64 bit aligned
+// GFX12: :[[@LINE-1]]:36: error: invalid operand for instruction
// 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]]:44: error: invalid register class: vgpr tuples must be 64 bit aligned
+// GFX12: :[[@LINE-1]]:44: error: invalid operand for instruction
// 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 ac9a1c7c65ba8..9297e98edacf5 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]]:24: error: invalid register class: vgpr tuples must be 64 bit aligned
+// GFX90A: :[[@LINE-1]]:24: error: invalid operand for instruction
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
>From cb84b466412d6da3ce68a6c96e6540d95698b458 Mon Sep 17 00:00:00 2001
From: Valery Pykhtin <valery.pykhtin at amd.com>
Date: Wed, 9 Sep 2026 11:17:47 +0000
Subject: [PATCH 3/3] [AMDGPU] Rename the alignment-aware operand predicates to
*Target, NFC
The predicates take a *_AlignTarget/_Target operand, so name them after
it. The HwMode resolution behind the operand's class is an implementation
detail that need not appear in the name.
---
.../AMDGPU/AsmParser/AMDGPUAsmParser.cpp | 28 +++++++++----------
1 file changed, 14 insertions(+), 14 deletions(-)
diff --git a/llvm/lib/Target/AMDGPU/AsmParser/AMDGPUAsmParser.cpp b/llvm/lib/Target/AMDGPU/AsmParser/AMDGPUAsmParser.cpp
index d75ac36f4f623..db5dd3a38b2a9 100644
--- a/llvm/lib/Target/AMDGPU/AsmParser/AMDGPUAsmParser.cpp
+++ b/llvm/lib/Target/AMDGPU/AsmParser/AMDGPUAsmParser.cpp
@@ -244,16 +244,16 @@ class AMDGPUOperand : public MCParsedAsmOperand {
return isRegClass(RCID) || isInlinableImm(type);
}
- bool isRegOrInlineByHwMode(unsigned RCByHwModeIdx, MVT type) const {
- return isRegClassByHwMode(RCByHwModeIdx) || isInlinableImm(type);
+ bool isRegOrInlineTarget(unsigned TargetRCIdx, MVT type) const {
+ return isRegClassTarget(TargetRCIdx) || 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 isRegOrImmWithInputModsTarget(unsigned TargetRCIdx, MVT type) const {
+ return isRegOrInlineTarget(TargetRCIdx, type) || isLiteralImm(type);
}
bool isRegOrImmWithInt16InputMods() const {
@@ -283,7 +283,7 @@ class AMDGPUOperand : public MCParsedAsmOperand {
}
bool isRegOrImmWithInt64InputMods() const {
- return isRegOrImmWithInputModsByHwMode(AMDGPU::VS_64_AlignTarget, MVT::i64);
+ return isRegOrImmWithInputModsTarget(AMDGPU::VS_64_AlignTarget, MVT::i64);
}
bool isRegOrImmWithFP16InputMods() const {
@@ -300,7 +300,7 @@ class AMDGPUOperand : public MCParsedAsmOperand {
}
bool isRegOrImmWithFP64InputMods() const {
- return isRegOrImmWithInputModsByHwMode(AMDGPU::VS_64_AlignTarget, MVT::f64);
+ return isRegOrImmWithInputModsTarget(AMDGPU::VS_64_AlignTarget, MVT::f64);
}
template <bool IsFake16> bool isRegOrInlineImmWithFP16InputMods() const {
@@ -313,7 +313,7 @@ class AMDGPUOperand : public MCParsedAsmOperand {
}
bool isRegOrInlineImmWithFP64InputMods() const {
- return isRegOrInlineByHwMode(AMDGPU::VS_64_AlignTarget, MVT::f64);
+ return isRegOrInlineTarget(AMDGPU::VS_64_AlignTarget, MVT::f64);
}
bool isVRegWithInputMods(unsigned RCID) const { return isRegClass(RCID); }
@@ -323,7 +323,7 @@ class AMDGPUOperand : public MCParsedAsmOperand {
}
bool isVRegWithFP64InputMods() const {
- return isRegClassByHwMode(AMDGPU::VReg_64_AlignTarget);
+ return isRegClassTarget(AMDGPU::VReg_64_AlignTarget);
}
bool isPackedFP16InputMods() const {
@@ -426,7 +426,7 @@ class AMDGPUOperand : public MCParsedAsmOperand {
bool isRegClass(unsigned RCID) const;
// Check the register against the HwMode-resolved operand class.
- bool isRegClassByHwMode(unsigned RCByHwModeIdx) const;
+ bool isRegClassTarget(unsigned TargetRCIdx) const;
bool isInlineValue() const;
@@ -1628,8 +1628,8 @@ class AMDGPUAsmParser : public MCTargetAsmParser {
// 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];
+ int16_t getTargetRegClass(unsigned TargetRCIdx) const {
+ return MII.getRegClassByHwModeTable(HwMode)[TargetRCIdx];
}
// FIXME: This should not be used. Instead, should use queries derived from
@@ -2275,10 +2275,10 @@ bool AMDGPUOperand::isLiteralImm(MVT type) const {
return canLosslesslyConvertToFPType(FPLiteral, ExpectedType);
}
-bool AMDGPUOperand::isRegClassByHwMode(unsigned RCByHwModeIdx) const {
+bool AMDGPUOperand::isRegClassTarget(unsigned TargetRCIdx) const {
if (!isRegKind())
return false;
- int16_t RCID = AsmParser->getRegClassByHwMode(RCByHwModeIdx);
+ int16_t RCID = AsmParser->getTargetRegClass(TargetRCIdx);
return RCID >= 0 && isRegClass(RCID);
}
@@ -2291,7 +2291,7 @@ bool AMDGPUOperand::isVRegWithInputMods() const {
return isRegClass(AMDGPU::VGPR_32RegClassID) ||
// GFX90A allows DPP on 64-bit operands.
(AsmParser->getFeatureBits()[AMDGPU::FeatureDPALU_DPP] &&
- isRegClassByHwMode(AMDGPU::VReg_64_AlignTarget));
+ isRegClassTarget(AMDGPU::VReg_64_AlignTarget));
}
template <bool IsFake16>
More information about the llvm-commits
mailing list