[llvm] [AMDGPU][AsmParser] Simplify getting source locations of operands. (PR #158323)
Ivan Kosarev via llvm-commits
llvm-commits at lists.llvm.org
Fri Sep 12 09:36:13 PDT 2025
https://github.com/kosarev created https://github.com/llvm/llvm-project/pull/158323
Remember indexes of MCOperands in MCParsedAsmOperands as we add them to instructions. Then use the indexes to find locations by known MCOperands indexes.
Happens to fix some reported locations in tests. NFCI otherwise.
getImmLoc() is to be eliminated as well; there's enough work for another patch.
>From 65df61f2bbbd63af3f6a207188964e4bad95c77c Mon Sep 17 00:00:00 2001
From: Ivan Kosarev <ivan.kosarev at amd.com>
Date: Fri, 12 Sep 2025 15:50:20 +0100
Subject: [PATCH] [AMDGPU][AsmParser] Simplify getting source locations of
operands.
Remember indexes of MCOperands in MCParsedAsmOperands as we add
them to instructions. Then use the indexes to find locations by
known MCOperands indexes.
Happens to fix some reported locations in tests. NFCI otherwise.
getImmLoc() is to be eliminated as well; there's enough work for
another patch.
---
.../AMDGPU/AsmParser/AMDGPUAsmParser.cpp | 287 ++++++------------
llvm/test/MC/AMDGPU/gfx11_asm_vopd_err.s | 14 +-
llvm/test/MC/AMDGPU/gfx1250_asm_vopd_errs.s | 12 +-
llvm/test/MC/AMDGPU/gfx12_asm_vopd_errs.s | 10 +-
llvm/test/MC/AMDGPU/mai-gfx950-err.s | 4 +-
llvm/test/MC/AMDGPU/vop3-literal.s | 72 ++---
6 files changed, 158 insertions(+), 241 deletions(-)
diff --git a/llvm/lib/Target/AMDGPU/AsmParser/AMDGPUAsmParser.cpp b/llvm/lib/Target/AMDGPU/AsmParser/AMDGPUAsmParser.cpp
index e420f2ad676f9..15d03403e3868 100644
--- a/llvm/lib/Target/AMDGPU/AsmParser/AMDGPUAsmParser.cpp
+++ b/llvm/lib/Target/AMDGPU/AsmParser/AMDGPUAsmParser.cpp
@@ -188,20 +188,6 @@ class AMDGPUOperand : public MCParsedAsmOperand {
ImmTyByteSel,
};
- // Immediate operand kind.
- // It helps to identify the location of an offending operand after an error.
- // Note that regular literals and mandatory literals (KImm) must be handled
- // differently. When looking for an offending operand, we should usually
- // ignore mandatory literals because they are part of the instruction and
- // cannot be changed. Report location of mandatory operands only for VOPD,
- // when both OpX and OpY have a KImm and there are no other literals.
- enum ImmKindTy {
- ImmKindTyNone,
- ImmKindTyLiteral,
- ImmKindTyMandatoryLiteral,
- ImmKindTyConst,
- };
-
private:
struct TokOp {
const char *Data;
@@ -212,7 +198,6 @@ class AMDGPUOperand : public MCParsedAsmOperand {
int64_t Val;
ImmTy Type;
bool IsFPImm;
- mutable ImmKindTy Kind;
Modifiers Mods;
};
@@ -228,6 +213,9 @@ class AMDGPUOperand : public MCParsedAsmOperand {
const MCExpr *Expr;
};
+ // The index of the associated MCInst operand.
+ mutable int MCOpIdx = -1;
+
public:
bool isToken() const override { return Kind == Token; }
@@ -239,38 +227,6 @@ class AMDGPUOperand : public MCParsedAsmOperand {
return Kind == Immediate;
}
- void setImmKindNone() const {
- assert(isImm());
- Imm.Kind = ImmKindTyNone;
- }
-
- void setImmKindLiteral() const {
- assert(isImm());
- Imm.Kind = ImmKindTyLiteral;
- }
-
- void setImmKindMandatoryLiteral() const {
- assert(isImm());
- Imm.Kind = ImmKindTyMandatoryLiteral;
- }
-
- void setImmKindConst() const {
- assert(isImm());
- Imm.Kind = ImmKindTyConst;
- }
-
- bool IsImmKindLiteral() const {
- return isImm() && Imm.Kind == ImmKindTyLiteral;
- }
-
- bool IsImmKindMandatoryLiteral() const {
- return isImm() && Imm.Kind == ImmKindTyMandatoryLiteral;
- }
-
- bool isImmKindConst() const {
- return isImm() && Imm.Kind == ImmKindTyConst;
- }
-
bool isInlinableImm(MVT type) const;
bool isLiteralImm(MVT type) const;
@@ -1055,6 +1011,8 @@ class AMDGPUOperand : public MCParsedAsmOperand {
return SMRange(StartLoc, EndLoc);
}
+ int getMCOpIdx() const { return MCOpIdx; }
+
Modifiers getModifiers() const {
assert(isRegKind() || isImmTy(ImmTyNone));
return isRegKind() ? Reg.Mods : Imm.Mods;
@@ -1242,7 +1200,6 @@ class AMDGPUOperand : public MCParsedAsmOperand {
auto Op = std::make_unique<AMDGPUOperand>(Immediate, AsmParser);
Op->Imm.Val = Val;
Op->Imm.IsFPImm = IsFPImm;
- Op->Imm.Kind = ImmKindTyNone;
Op->Imm.Type = Type;
Op->Imm.Mods = Modifiers();
Op->StartLoc = Loc;
@@ -1836,25 +1793,24 @@ class AMDGPUAsmParser : public MCTargetAsmParser {
ParseStatus parseHwregFunc(OperandInfoTy &HwReg, OperandInfoTy &Offset,
OperandInfoTy &Width);
+ static SMLoc getLaterLoc(SMLoc a, SMLoc b);
+
SMLoc getFlatOffsetLoc(const OperandVector &Operands) const;
SMLoc getSMEMOffsetLoc(const OperandVector &Operands) const;
SMLoc getBLGPLoc(const OperandVector &Operands) const;
+ SMLoc getOperandLoc(const OperandVector &Operands, int MCOpIdx) const;
SMLoc getOperandLoc(std::function<bool(const AMDGPUOperand&)> Test,
const OperandVector &Operands) const;
- SMLoc getImmLoc(AMDGPUOperand::ImmTy Type, const OperandVector &Operands) const;
- SMLoc getRegLoc(MCRegister Reg, const OperandVector &Operands) const;
- SMLoc getLitLoc(const OperandVector &Operands,
- bool SearchMandatoryLiterals = false) const;
- SMLoc getMandatoryLitLoc(const OperandVector &Operands) const;
- SMLoc getConstLoc(const OperandVector &Operands) const;
+ SMLoc getImmLoc(AMDGPUOperand::ImmTy Type,
+ const OperandVector &Operands) const;
SMLoc getInstLoc(const OperandVector &Operands) const;
bool validateInstruction(const MCInst &Inst, const SMLoc &IDLoc, const OperandVector &Operands);
bool validateOffset(const MCInst &Inst, const OperandVector &Operands);
bool validateFlatOffset(const MCInst &Inst, const OperandVector &Operands);
bool validateSMEMOffset(const MCInst &Inst, const OperandVector &Operands);
- bool validateSOPLiteral(const MCInst &Inst) const;
+ bool validateSOPLiteral(const MCInst &Inst, const OperandVector &Operands);
bool validateConstantBusLimitations(const MCInst &Inst, const OperandVector &Operands);
std::optional<unsigned> checkVOPDRegBankConstraints(const MCInst &Inst,
bool AsVOPD3);
@@ -1895,7 +1851,7 @@ class AMDGPUAsmParser : public MCTargetAsmParser {
const unsigned CPol);
bool validateTFE(const MCInst &Inst, const OperandVector &Operands);
bool validateSetVgprMSB(const MCInst &Inst, const OperandVector &Operands);
- std::optional<StringRef> validateLdsDirect(const MCInst &Inst);
+ bool validateLdsDirect(const MCInst &Inst, const OperandVector &Operands);
bool validateWMMA(const MCInst &Inst, const OperandVector &Operands);
unsigned getConstantBusLimit(unsigned Opcode) const;
bool usesConstantBus(const MCInst &Inst, unsigned OpIdx);
@@ -2337,6 +2293,8 @@ uint64_t AMDGPUOperand::applyInputFPModifiers(uint64_t Val, unsigned Size) const
}
void AMDGPUOperand::addImmOperands(MCInst &Inst, unsigned N, bool ApplyModifiers) const {
+ MCOpIdx = Inst.getNumOperands();
+
if (isExpr()) {
Inst.addOperand(MCOperand::createExpr(Expr));
return;
@@ -2350,7 +2308,6 @@ void AMDGPUOperand::addImmOperands(MCInst &Inst, unsigned N, bool ApplyModifiers
} else {
assert(!isImmTy(ImmTyNone) || !hasModifiers());
Inst.addOperand(MCOperand::createImm(Imm.Val));
- setImmKindNone();
}
}
@@ -2379,7 +2336,6 @@ void AMDGPUOperand::addLiteralImmOperand(MCInst &Inst, int64_t Val, bool ApplyMo
if (AMDGPU::isInlinableLiteral64(Literal.getZExtValue(),
AsmParser->hasInv2PiInlineImm())) {
Inst.addOperand(MCOperand::createImm(Literal.getZExtValue()));
- setImmKindConst();
return;
}
@@ -2400,7 +2356,6 @@ void AMDGPUOperand::addLiteralImmOperand(MCInst &Inst, int64_t Val, bool ApplyMo
}
Inst.addOperand(MCOperand::createImm(Val));
- setImmKindLiteral();
return;
}
@@ -2411,7 +2366,6 @@ void AMDGPUOperand::addLiteralImmOperand(MCInst &Inst, int64_t Val, bool ApplyMo
case AMDGPU::OPERAND_KIMM64:
Inst.addOperand(MCOperand::createImm(Val));
- setImmKindMandatoryLiteral();
return;
case AMDGPU::OPERAND_REG_IMM_BF16:
@@ -2424,7 +2378,6 @@ void AMDGPUOperand::addLiteralImmOperand(MCInst &Inst, int64_t Val, bool ApplyMo
// 1/(2*pi) = 0.15915494 since bf16 is in fact fp32 with cleared low 16
// bits. Prevent rounding below.
Inst.addOperand(MCOperand::createImm(0x3e22));
- setImmKindLiteral();
return;
}
[[fallthrough]];
@@ -2459,11 +2412,6 @@ void AMDGPUOperand::addLiteralImmOperand(MCInst &Inst, int64_t Val, bool ApplyMo
uint64_t ImmVal = FPLiteral.bitcastToAPInt().getZExtValue();
Inst.addOperand(MCOperand::createImm(ImmVal));
- if (OpTy == AMDGPU::OPERAND_KIMM32 || OpTy == AMDGPU::OPERAND_KIMM16) {
- setImmKindMandatoryLiteral();
- } else {
- setImmKindLiteral();
- }
return;
}
default:
@@ -2492,7 +2440,6 @@ void AMDGPUOperand::addLiteralImmOperand(MCInst &Inst, int64_t Val, bool ApplyMo
AMDGPU::isInlinableLiteral32(static_cast<int32_t>(Val),
AsmParser->hasInv2PiInlineImm())) {
Inst.addOperand(MCOperand::createImm(Val));
- setImmKindConst();
return;
}
[[fallthrough]];
@@ -2500,14 +2447,12 @@ void AMDGPUOperand::addLiteralImmOperand(MCInst &Inst, int64_t Val, bool ApplyMo
case AMDGPU::OPERAND_REG_IMM_NOINLINE_V2FP16:
Inst.addOperand(MCOperand::createImm(Lo_32(Val)));
- setImmKindLiteral();
return;
case AMDGPU::OPERAND_REG_IMM_INT64:
case AMDGPU::OPERAND_REG_INLINE_C_INT64:
if (AMDGPU::isInlinableLiteral64(Val, AsmParser->hasInv2PiInlineImm())) {
Inst.addOperand(MCOperand::createImm(Val));
- setImmKindConst();
return;
}
@@ -2519,7 +2464,6 @@ void AMDGPUOperand::addLiteralImmOperand(MCInst &Inst, int64_t Val, bool ApplyMo
Val = Lo_32(Val);
Inst.addOperand(MCOperand::createImm(Val));
- setImmKindLiteral();
return;
case AMDGPU::OPERAND_REG_IMM_FP64:
@@ -2527,7 +2471,6 @@ void AMDGPUOperand::addLiteralImmOperand(MCInst &Inst, int64_t Val, bool ApplyMo
case AMDGPU::OPERAND_REG_INLINE_AC_FP64:
if (AMDGPU::isInlinableLiteral64(Val, AsmParser->hasInv2PiInlineImm())) {
Inst.addOperand(MCOperand::createImm(Val));
- setImmKindConst();
return;
}
@@ -2547,7 +2490,6 @@ void AMDGPUOperand::addLiteralImmOperand(MCInst &Inst, int64_t Val, bool ApplyMo
}
Inst.addOperand(MCOperand::createImm(Val));
- setImmKindLiteral();
return;
case AMDGPU::OPERAND_REG_IMM_INT16:
@@ -2555,12 +2497,10 @@ void AMDGPUOperand::addLiteralImmOperand(MCInst &Inst, int64_t Val, bool ApplyMo
if (isSafeTruncation(Val, 16) &&
AMDGPU::isInlinableIntLiteral(static_cast<int16_t>(Val))) {
Inst.addOperand(MCOperand::createImm(Lo_32(Val)));
- setImmKindConst();
return;
}
Inst.addOperand(MCOperand::createImm(Val & 0xffff));
- setImmKindLiteral();
return;
case AMDGPU::OPERAND_REG_INLINE_C_FP16:
@@ -2569,12 +2509,10 @@ void AMDGPUOperand::addLiteralImmOperand(MCInst &Inst, int64_t Val, bool ApplyMo
AMDGPU::isInlinableLiteralFP16(static_cast<int16_t>(Val),
AsmParser->hasInv2PiInlineImm())) {
Inst.addOperand(MCOperand::createImm(Val));
- setImmKindConst();
return;
}
Inst.addOperand(MCOperand::createImm(Val & 0xffff));
- setImmKindLiteral();
return;
case AMDGPU::OPERAND_REG_IMM_BF16:
@@ -2583,12 +2521,10 @@ void AMDGPUOperand::addLiteralImmOperand(MCInst &Inst, int64_t Val, bool ApplyMo
AMDGPU::isInlinableLiteralBF16(static_cast<int16_t>(Val),
AsmParser->hasInv2PiInlineImm())) {
Inst.addOperand(MCOperand::createImm(Val));
- setImmKindConst();
return;
}
Inst.addOperand(MCOperand::createImm(Val & 0xffff));
- setImmKindLiteral();
return;
case AMDGPU::OPERAND_REG_INLINE_C_V2INT16: {
@@ -2617,18 +2553,15 @@ void AMDGPUOperand::addLiteralImmOperand(MCInst &Inst, int64_t Val, bool ApplyMo
case AMDGPU::OPERAND_KIMM32:
Inst.addOperand(MCOperand::createImm(Literal.getLoBits(32).getZExtValue()));
- setImmKindMandatoryLiteral();
return;
case AMDGPU::OPERAND_KIMM16:
Inst.addOperand(MCOperand::createImm(Literal.getLoBits(16).getZExtValue()));
- setImmKindMandatoryLiteral();
return;
case AMDGPU::OPERAND_KIMM64:
if ((isInt<32>(Val) || isUInt<32>(Val)) && !getModifiers().Lit64)
Val <<= 32;
Inst.addOperand(MCOperand::createImm(Val));
- setImmKindMandatoryLiteral();
return;
default:
llvm_unreachable("invalid operand size");
@@ -2636,6 +2569,7 @@ void AMDGPUOperand::addLiteralImmOperand(MCInst &Inst, int64_t Val, bool ApplyMo
}
void AMDGPUOperand::addRegOperands(MCInst &Inst, unsigned N) const {
+ MCOpIdx = Inst.getNumOperands();
Inst.addOperand(MCOperand::createReg(AMDGPU::getMCReg(getReg(), AsmParser->getSTI())));
}
@@ -3942,6 +3876,8 @@ bool AMDGPUAsmParser::validateConstantBusLimitations(
OperandIndices OpIndices = getSrcOperandIndices(Opcode);
+ unsigned ConstantBusLimit = getConstantBusLimit(Opcode);
+
for (int OpIdx : OpIndices) {
if (OpIdx == -1)
continue;
@@ -3985,17 +3921,14 @@ bool AMDGPUAsmParser::validateConstantBusLimitations(
}
}
}
- }
- ConstantBusUseCount += NumLiterals;
- if (ConstantBusUseCount <= getConstantBusLimit(Opcode))
- return true;
-
- SMLoc LitLoc = getLitLoc(Operands);
- SMLoc RegLoc = getRegLoc(LastSGPR, Operands);
- SMLoc Loc = (LitLoc.getPointer() < RegLoc.getPointer()) ? RegLoc : LitLoc;
- Error(Loc, "invalid operand (violates constant bus restrictions)");
- return false;
+ if (ConstantBusUseCount + NumLiterals > ConstantBusLimit) {
+ Error(getOperandLoc(Operands, OpIdx),
+ "invalid operand (violates constant bus restrictions)");
+ return false;
+ }
+ }
+ return true;
}
std::optional<unsigned>
@@ -4408,19 +4341,15 @@ bool AMDGPUAsmParser::validateMovrels(const MCInst &Inst,
const int Src0Idx = AMDGPU::getNamedOperandIdx(Opc, AMDGPU::OpName::src0);
assert(Src0Idx != -1);
- SMLoc ErrLoc;
const MCOperand &Src0 = Inst.getOperand(Src0Idx);
if (Src0.isReg()) {
auto Reg = mc2PseudoReg(Src0.getReg());
const MCRegisterInfo *TRI = getContext().getRegisterInfo();
if (!isSGPR(Reg, TRI))
return true;
- ErrLoc = getRegLoc(Reg, Operands);
- } else {
- ErrLoc = getConstLoc(Operands);
}
- Error(ErrLoc, "source operand must be a VGPR");
+ Error(getOperandLoc(Operands, Src0Idx), "source operand must be a VGPR");
return false;
}
@@ -4442,7 +4371,7 @@ bool AMDGPUAsmParser::validateMAIAccWrite(const MCInst &Inst,
auto Reg = mc2PseudoReg(Src0.getReg());
const MCRegisterInfo *TRI = getContext().getRegisterInfo();
if (!isGFX90A() && isSGPR(Reg, TRI)) {
- Error(getRegLoc(Reg, Operands),
+ Error(getOperandLoc(Operands, Src0Idx),
"source operand must be either a VGPR or an inline constant");
return false;
}
@@ -4464,7 +4393,7 @@ bool AMDGPUAsmParser::validateMAISrc2(const MCInst &Inst,
return true;
if (Inst.getOperand(Src2Idx).isImm() && isInlineConstant(Inst, Src2Idx)) {
- Error(getConstLoc(Operands),
+ Error(getOperandLoc(Operands, Src2Idx),
"inline constants are not allowed for this operand");
return false;
}
@@ -4494,16 +4423,14 @@ bool AMDGPUAsmParser::validateMFMA(const MCInst &Inst,
bool Success = true;
if (Info->NumRegsSrcA != mfmaScaleF8F6F4FormatToNumRegs(CBSZ)) {
int Src0Idx = AMDGPU::getNamedOperandIdx(Opc, AMDGPU::OpName::src0);
- Error(getRegLoc(mc2PseudoReg(Inst.getOperand(Src0Idx).getReg()),
- Operands),
+ Error(getOperandLoc(Operands, Src0Idx),
"wrong register tuple size for cbsz value " + Twine(CBSZ));
Success = false;
}
if (Info->NumRegsSrcB != mfmaScaleF8F6F4FormatToNumRegs(BLGP)) {
int Src1Idx = AMDGPU::getNamedOperandIdx(Opc, AMDGPU::OpName::src1);
- Error(getRegLoc(mc2PseudoReg(Inst.getOperand(Src1Idx).getReg()),
- Operands),
+ Error(getOperandLoc(Operands, Src1Idx),
"wrong register tuple size for blgp value " + Twine(BLGP));
Success = false;
}
@@ -4530,7 +4457,7 @@ bool AMDGPUAsmParser::validateMFMA(const MCInst &Inst,
return true;
if (TRI->regsOverlap(Src2Reg, DstReg)) {
- Error(getRegLoc(mc2PseudoReg(Src2Reg), Operands),
+ Error(getOperandLoc(Operands, Src2Idx),
"source 2 operand must not partially overlap with dst");
return false;
}
@@ -4724,9 +4651,8 @@ static bool IsRevOpcode(const unsigned Opcode)
}
}
-std::optional<StringRef>
-AMDGPUAsmParser::validateLdsDirect(const MCInst &Inst) {
-
+bool AMDGPUAsmParser::validateLdsDirect(const MCInst &Inst,
+ const OperandVector &Operands) {
using namespace SIInstrFlags;
const unsigned Opcode = Inst.getOpcode();
const MCInstrDesc &Desc = MII.get(Opcode);
@@ -4735,7 +4661,7 @@ AMDGPUAsmParser::validateLdsDirect(const MCInst &Inst) {
// with 9-bit operands only. Ignore encodings which do not accept these.
const auto Enc = VOP1 | VOP2 | VOP3 | VOPC | VOP3P | SIInstrFlags::SDWA;
if ((Desc.TSFlags & Enc) == 0)
- return std::nullopt;
+ return true;
for (auto SrcName : {OpName::src0, OpName::src1, OpName::src2}) {
auto SrcIdx = getNamedOperandIdx(Opcode, SrcName);
@@ -4744,18 +4670,27 @@ AMDGPUAsmParser::validateLdsDirect(const MCInst &Inst) {
const auto &Src = Inst.getOperand(SrcIdx);
if (Src.isReg() && Src.getReg() == LDS_DIRECT) {
- if (isGFX90A() || isGFX11Plus())
- return StringRef("lds_direct is not supported on this GPU");
+ if (isGFX90A() || isGFX11Plus()) {
+ Error(getOperandLoc(Operands, SrcIdx),
+ "lds_direct is not supported on this GPU");
+ return false;
+ }
- if (IsRevOpcode(Opcode) || (Desc.TSFlags & SIInstrFlags::SDWA))
- return StringRef("lds_direct cannot be used with this instruction");
+ if (IsRevOpcode(Opcode) || (Desc.TSFlags & SIInstrFlags::SDWA)) {
+ Error(getOperandLoc(Operands, SrcIdx),
+ "lds_direct cannot be used with this instruction");
+ return false;
+ }
- if (SrcName != OpName::src0)
- return StringRef("lds_direct may be used as src0 only");
+ if (SrcName != OpName::src0) {
+ Error(getOperandLoc(Operands, SrcIdx),
+ "lds_direct may be used as src0 only");
+ return false;
+ }
}
}
- return std::nullopt;
+ return true;
}
SMLoc AMDGPUAsmParser::getFlatOffsetLoc(const OperandVector &Operands) const {
@@ -4881,7 +4816,8 @@ bool AMDGPUAsmParser::validateSMEMOffset(const MCInst &Inst,
return false;
}
-bool AMDGPUAsmParser::validateSOPLiteral(const MCInst &Inst) const {
+bool AMDGPUAsmParser::validateSOPLiteral(const MCInst &Inst,
+ const OperandVector &Operands) {
unsigned Opcode = Inst.getOpcode();
const MCInstrDesc &Desc = MII.get(Opcode);
if (!(Desc.TSFlags & (SIInstrFlags::SOP2 | SIInstrFlags::SOPC)))
@@ -4914,7 +4850,12 @@ bool AMDGPUAsmParser::validateSOPLiteral(const MCInst &Inst) const {
}
}
- return NumLiterals + NumExprs <= 1;
+ if (NumLiterals + NumExprs <= 1)
+ return true;
+
+ Error(getOperandLoc(Operands, Src1Idx),
+ "only one unique literal operand is allowed");
+ return false;
}
bool AMDGPUAsmParser::validateOpSel(const MCInst &Inst) {
@@ -5090,9 +5031,8 @@ bool AMDGPUAsmParser::validateDPP(const MCInst &Inst,
const MCOperand &Src1 = Inst.getOperand(Src1Idx);
const MCRegisterInfo *TRI = getContext().getRegisterInfo();
if (Src1.isReg() && isSGPR(mc2PseudoReg(Src1.getReg()), TRI)) {
- auto Reg = mc2PseudoReg(Inst.getOperand(Src1Idx).getReg());
- SMLoc S = getRegLoc(Reg, Operands);
- Error(S, "invalid operand for instruction");
+ Error(getOperandLoc(Operands, Src1Idx),
+ "invalid operand for instruction");
return false;
}
if (Src1.isImm()) {
@@ -5125,9 +5065,8 @@ bool AMDGPUAsmParser::validateVOPLiteral(const MCInst &Inst,
OperandIndices OpIndices = getSrcOperandIndices(Opcode, HasMandatoryLiteral);
- unsigned NumExprs = 0;
- unsigned NumLiterals = 0;
- uint64_t LiteralValue;
+ std::optional<unsigned> LiteralOpIdx;
+ std::optional<uint64_t> LiteralValue;
for (int OpIdx : OpIndices) {
if (OpIdx == -1)
@@ -5139,6 +5078,7 @@ bool AMDGPUAsmParser::validateVOPLiteral(const MCInst &Inst,
if (!isSISrcOperand(Desc, OpIdx))
continue;
+ bool IsAnotherLiteral = false;
if (MO.isImm() && !isInlineConstant(Inst, OpIdx)) {
uint64_t Value = static_cast<uint64_t>(MO.getImm());
bool IsForcedFP64 =
@@ -5151,34 +5091,37 @@ bool AMDGPUAsmParser::validateVOPLiteral(const MCInst &Inst,
if (!IsValid32Op && !isInt<32>(Value) && !isUInt<32>(Value) &&
!IsForcedFP64 && (!has64BitLiterals() || Desc.getSize() != 4)) {
- Error(getLitLoc(Operands), "invalid operand for instruction");
+ Error(getOperandLoc(Operands, OpIdx),
+ "invalid operand for instruction");
return false;
}
if (IsFP64 && IsValid32Op && !IsForcedFP64)
Value = Hi_32(Value);
- if (NumLiterals == 0 || LiteralValue != Value) {
- LiteralValue = Value;
- ++NumLiterals;
- }
+ IsAnotherLiteral = !LiteralValue || *LiteralValue != Value;
+ LiteralValue = Value;
} else if (MO.isExpr()) {
- ++NumExprs;
+ // Literal value not known, so we conservately assume it's different.
+ IsAnotherLiteral = true;
}
- }
- NumLiterals += NumExprs;
- if (!NumLiterals)
- return true;
+ if (IsAnotherLiteral && !HasMandatoryLiteral &&
+ !getFeatureBits()[FeatureVOP3Literal]) {
+ Error(getOperandLoc(Operands, OpIdx),
+ "literal operands are not supported");
+ return false;
+ }
- if (!HasMandatoryLiteral && !getFeatureBits()[FeatureVOP3Literal]) {
- Error(getLitLoc(Operands), "literal operands are not supported");
- return false;
- }
+ if (LiteralOpIdx && IsAnotherLiteral) {
+ Error(getLaterLoc(getOperandLoc(Operands, OpIdx),
+ getOperandLoc(Operands, *LiteralOpIdx)),
+ "only one unique literal operand is allowed");
+ return false;
+ }
- if (NumLiterals > 1) {
- Error(getLitLoc(Operands, true), "only one unique literal operand is allowed");
- return false;
+ if (IsAnotherLiteral)
+ LiteralOpIdx = OpIdx;
}
return true;
@@ -5352,8 +5295,7 @@ bool AMDGPUAsmParser::validateWaitCnt(const MCInst &Inst,
if (Reg == AMDGPU::SGPR_NULL)
return true;
- SMLoc RegLoc = getRegLoc(Reg, Operands);
- Error(RegLoc, "src0 must be null");
+ Error(getOperandLoc(Operands, Src0Idx), "src0 must be null");
return false;
}
@@ -5400,8 +5342,7 @@ bool AMDGPUAsmParser::validateGWS(const MCInst &Inst,
auto Reg = Inst.getOperand(Data0Pos).getReg();
auto RegIdx = Reg - (VGPR32.contains(Reg) ? AMDGPU::VGPR0 : AMDGPU::AGPR0);
if (RegIdx & 1) {
- SMLoc RegLoc = getRegLoc(Reg, Operands);
- Error(RegLoc, "vgpr must be even aligned");
+ Error(getOperandLoc(Operands, Data0Pos), "vgpr must be even aligned");
return false;
}
@@ -5598,7 +5539,7 @@ bool AMDGPUAsmParser::validateWMMA(const MCInst &Inst,
"MATRIX_FMT_FP6", "MATRIX_FMT_BF6",
"MATRIX_FMT_FP4"};
- Error(getRegLoc(mc2PseudoReg(Inst.getOperand(SrcIdx).getReg()), Operands),
+ Error(getOperandLoc(Operands, SrcIdx),
"wrong register tuple size for " + Twine(FmtNames[Fmt]));
return false;
};
@@ -5610,20 +5551,15 @@ bool AMDGPUAsmParser::validateWMMA(const MCInst &Inst,
bool AMDGPUAsmParser::validateInstruction(const MCInst &Inst,
const SMLoc &IDLoc,
const OperandVector &Operands) {
- if (auto ErrMsg = validateLdsDirect(Inst)) {
- Error(getRegLoc(LDS_DIRECT, Operands), *ErrMsg);
+ if (!validateLdsDirect(Inst, Operands))
return false;
- }
if (!validateTrue16OpSel(Inst)) {
Error(getImmLoc(AMDGPUOperand::ImmTyOpSel, Operands),
"op_sel operand conflicts with 16-bit operand suffix");
return false;
}
- if (!validateSOPLiteral(Inst)) {
- Error(getLitLoc(Operands),
- "only one unique literal operand is allowed");
+ if (!validateSOPLiteral(Inst, Operands))
return false;
- }
if (!validateVOPLiteral(Inst, Operands)) {
return false;
}
@@ -8563,6 +8499,20 @@ SMLoc AMDGPUAsmParser::getInstLoc(const OperandVector &Operands) const {
return ((AMDGPUOperand &)*Operands[0]).getStartLoc();
}
+SMLoc AMDGPUAsmParser::getLaterLoc(SMLoc a, SMLoc b) {
+ return a.getPointer() < b.getPointer() ? b : a;
+}
+
+SMLoc AMDGPUAsmParser::getOperandLoc(const OperandVector &Operands,
+ int MCOpIdx) const {
+ for (const auto &Op : Operands) {
+ const auto TargetOp = static_cast<AMDGPUOperand &>(*Op);
+ if (TargetOp.getMCOpIdx() == MCOpIdx)
+ return TargetOp.getStartLoc();
+ }
+ llvm_unreachable("No such MC operand!");
+}
+
SMLoc
AMDGPUAsmParser::getOperandLoc(std::function<bool(const AMDGPUOperand&)> Test,
const OperandVector &Operands) const {
@@ -8581,40 +8531,6 @@ AMDGPUAsmParser::getImmLoc(AMDGPUOperand::ImmTy Type,
return getOperandLoc(Test, Operands);
}
-SMLoc AMDGPUAsmParser::getRegLoc(MCRegister Reg,
- const OperandVector &Operands) const {
- auto Test = [=](const AMDGPUOperand& Op) {
- return Op.isRegKind() && Op.getReg() == Reg;
- };
- return getOperandLoc(Test, Operands);
-}
-
-SMLoc AMDGPUAsmParser::getLitLoc(const OperandVector &Operands,
- bool SearchMandatoryLiterals) const {
- auto Test = [](const AMDGPUOperand& Op) {
- return Op.IsImmKindLiteral() || Op.isExpr();
- };
- SMLoc Loc = getOperandLoc(Test, Operands);
- if (SearchMandatoryLiterals && Loc == getInstLoc(Operands))
- Loc = getMandatoryLitLoc(Operands);
- return Loc;
-}
-
-SMLoc AMDGPUAsmParser::getMandatoryLitLoc(const OperandVector &Operands) const {
- auto Test = [](const AMDGPUOperand &Op) {
- return Op.IsImmKindMandatoryLiteral();
- };
- return getOperandLoc(Test, Operands);
-}
-
-SMLoc
-AMDGPUAsmParser::getConstLoc(const OperandVector &Operands) const {
- auto Test = [](const AMDGPUOperand& Op) {
- return Op.isImmKindConst();
- };
- return getOperandLoc(Test, Operands);
-}
-
ParseStatus
AMDGPUAsmParser::parseStructuredOpFields(ArrayRef<StructuredOpField *> Fields) {
if (!trySkipToken(AsmToken::LCurly))
@@ -10220,7 +10136,6 @@ void AMDGPUAsmParser::cvtVOP3DPP(MCInst &Inst, const OperandVector &Operands,
Op.addRegOperands(Inst, 1);
} else if (Op.isImm() &&
Desc.operands()[Inst.getNumOperands()].RegClass != -1) {
- assert(!Op.IsImmKindLiteral() && "Cannot use literal with DPP");
Op.addImmOperands(Inst, 1);
} else if (Op.isImm()) {
OptionalIdx[Op.getImmTy()] = I;
diff --git a/llvm/test/MC/AMDGPU/gfx11_asm_vopd_err.s b/llvm/test/MC/AMDGPU/gfx11_asm_vopd_err.s
index 3c5905b14e06c..497fd86251fb4 100644
--- a/llvm/test/MC/AMDGPU/gfx11_asm_vopd_err.s
+++ b/llvm/test/MC/AMDGPU/gfx11_asm_vopd_err.s
@@ -22,7 +22,7 @@ v_dual_fmamk_f32 v122, v74, 0xa0172923, v161 :: v_dual_lshlrev_b32 v247, 0
v_dual_add_f32 v5, 0xaf123456, v2 :: v_dual_fmaak_f32 v6, v3, v1, 0xbabe
// GFX11: :[[@LINE-1]]:{{[0-9]+}}: error: only one unique literal operand is allowed
// GFX11-NEXT:{{^}}v_dual_add_f32 v5, 0xaf123456, v2 :: v_dual_fmaak_f32 v6, v3, v1, 0xbabe
-// GFX11-NEXT:{{^}} ^
+// GFX11-NEXT:{{^}} ^
v_dual_add_f32 v5, 0xaf123456, v2 :: v_dual_fmaak_f32 v6, 0xbabe, v1, 0xbabe
// GFX11: :[[@LINE-1]]:{{[0-9]+}}: error: only one unique literal operand is allowed
@@ -32,12 +32,12 @@ v_dual_add_f32 v5, 0xaf123456, v2 :: v_dual_fmaak_f32 v6, 0x
v_dual_fmamk_f32 v122, 0xdeadbeef, 0xdeadbeef, v161 :: v_dual_fmamk_f32 v123, 0xdeadbeef, 0x1234, v162
// GFX11: :[[@LINE-1]]:{{[0-9]+}}: error: only one unique literal operand is allowed
// GFX11-NEXT:{{^}}v_dual_fmamk_f32 v122, 0xdeadbeef, 0xdeadbeef, v161 :: v_dual_fmamk_f32 v123, 0xdeadbeef, 0x1234, v162
-// GFX11-NEXT:{{^}} ^
+// GFX11-NEXT:{{^}} ^
v_dual_fmamk_f32 v122, 0xdeadbeef, 0xdeadbeef, v161 :: v_dual_fmamk_f32 v123, s0, 0x1234, v162
// GFX11: :[[@LINE-1]]:{{[0-9]+}}: error: only one unique literal operand is allowed
// GFX11-NEXT:{{^}}v_dual_fmamk_f32 v122, 0xdeadbeef, 0xdeadbeef, v161 :: v_dual_fmamk_f32 v123, s0, 0x1234, v162
-// GFX11-NEXT:{{^}} ^
+// GFX11-NEXT:{{^}} ^
//===----------------------------------------------------------------------===//
// Check that KImm operands are counted as literals
@@ -52,12 +52,12 @@ v_dual_fmamk_f32 v122, v74, 0, v161 :: v_dual_lshlrev_b32 v247, 0
v_dual_add_f32 v5, 0xaf123456, v2 :: v_dual_fmaak_f32 v6, v3, v1, 1.0
// GFX11: :[[@LINE-1]]:{{[0-9]+}}: error: only one unique literal operand is allowed
// GFX11-NEXT:{{^}}v_dual_add_f32 v5, 0xaf123456, v2 :: v_dual_fmaak_f32 v6, v3, v1, 1.0
-// GFX11-NEXT:{{^}} ^
+// GFX11-NEXT:{{^}} ^
v_dual_fmamk_f32 v122, 0xdeadbeef, 2, v161 :: v_dual_fmamk_f32 v123, s0, 1, v162
// GFX11: :[[@LINE-1]]:{{[0-9]+}}: error: only one unique literal operand is allowed
// GFX11-NEXT:{{^}}v_dual_fmamk_f32 v122, 0xdeadbeef, 2, v161 :: v_dual_fmamk_f32 v123, s0, 1, v162
-// GFX11-NEXT:{{^}} ^
+// GFX11-NEXT:{{^}} ^
v_dual_fmamk_f32 v122, v1, 2, v161 :: v_dual_fmamk_f32 v123, s0, 1, v162
// GFX11: :[[@LINE-1]]:{{[0-9]+}}: error: only one unique literal operand is allowed
@@ -71,7 +71,7 @@ v_dual_fmamk_f32 v122, v1, 2, v161 :: v_dual_fmamk_f32 v123, s0,
v_dual_fmamk_f32 v122, 0xdeadbeef, 0xdeadbeef, v161 :: v_dual_fmamk_f32 v123, 0xdeadbeef, 0x1234, v162
// GFX11: :[[@LINE-1]]:{{[0-9]+}}: error: only one unique literal operand is allowed
// GFX11-NEXT:{{^}}v_dual_fmamk_f32 v122, 0xdeadbeef, 0xdeadbeef, v161 :: v_dual_fmamk_f32 v123, 0xdeadbeef, 0x1234, v162
-// GFX11-NEXT:{{^}} ^
+// GFX11-NEXT:{{^}} ^
v_dual_fmamk_f32 v122, 0xdeadbeef, 0xdeadbeef, v161 :: v_dual_fmamk_f32 v123, 0x1234, 0xdeadbeef, v162
// GFX11: :[[@LINE-1]]:{{[0-9]+}}: error: only one unique literal operand is allowed
@@ -81,7 +81,7 @@ v_dual_fmamk_f32 v122, 0xdeadbeef, 0xdeadbeef, v161 :: v_dual_fmamk_f32 v12
v_dual_fmamk_f32 v122, 0xdeadbeef, 0x1234, v161 :: v_dual_fmamk_f32 v123, 0xdeadbeef, 0xdeadbeef, v162
// GFX11: :[[@LINE-1]]:{{[0-9]+}}: error: only one unique literal operand is allowed
// GFX11-NEXT:{{^}}v_dual_fmamk_f32 v122, 0xdeadbeef, 0x1234, v161 :: v_dual_fmamk_f32 v123, 0xdeadbeef, 0xdeadbeef, v162
-// GFX11-NEXT:{{^}} ^
+// GFX11-NEXT:{{^}} ^
v_dual_fmamk_f32 v122, 0x1234, 0xdeadbeef, v161 :: v_dual_fmamk_f32 v123, 0xdeadbeef, 0xdeadbeef, v162
// GFX11: :[[@LINE-1]]:{{[0-9]+}}: error: only one unique literal operand is allowed
diff --git a/llvm/test/MC/AMDGPU/gfx1250_asm_vopd_errs.s b/llvm/test/MC/AMDGPU/gfx1250_asm_vopd_errs.s
index 81b79cb8c28da..f8cdf31c48e17 100644
--- a/llvm/test/MC/AMDGPU/gfx1250_asm_vopd_errs.s
+++ b/llvm/test/MC/AMDGPU/gfx1250_asm_vopd_errs.s
@@ -22,7 +22,7 @@ v_dual_fmamk_f32 v122, v74, 0xa0172923, v161 :: v_dual_lshlrev_b32 v247, 0
v_dual_add_f32 v5, 0xaf123456, v2 :: v_dual_fmaak_f32 v6, v3, v1, 0xbabe
// GFX12: :[[@LINE-1]]:{{[0-9]+}}: error: only one unique literal operand is allowed
// GFX12-NEXT:{{^}}v_dual_add_f32 v5, 0xaf123456, v2 :: v_dual_fmaak_f32 v6, v3, v1, 0xbabe
-// GFX12-NEXT:{{^}} ^
+// GFX12-NEXT:{{^}} ^
v_dual_add_f32 v5, 0xaf123456, v2 :: v_dual_fmaak_f32 v6, 0xbabe, v1, 0xbabe
// GFX12: :[[@LINE-1]]:{{[0-9]+}}: error: only one unique literal operand is allowed
@@ -32,12 +32,12 @@ v_dual_add_f32 v5, 0xaf123456, v2 :: v_dual_fmaak_f32 v6, 0x
v_dual_fmamk_f32 v122, 0xdeadbeef, 0xdeadbeef, v161 :: v_dual_fmamk_f32 v123, 0xdeadbeef, 0x1234, v162
// GFX12: :[[@LINE-1]]:{{[0-9]+}}: error: only one unique literal operand is allowed
// GFX12-NEXT:{{^}}v_dual_fmamk_f32 v122, 0xdeadbeef, 0xdeadbeef, v161 :: v_dual_fmamk_f32 v123, 0xdeadbeef, 0x1234, v162
-// GFX12-NEXT:{{^}} ^
+// GFX12-NEXT:{{^}} ^
v_dual_fmamk_f32 v122, 0xdeadbeef, 0xdeadbeef, v161 :: v_dual_fmamk_f32 v123, s0, 0x1234, v162
// GFX12: :[[@LINE-1]]:{{[0-9]+}}: error: only one unique literal operand is allowed
// GFX12-NEXT:{{^}}v_dual_fmamk_f32 v122, 0xdeadbeef, 0xdeadbeef, v161 :: v_dual_fmamk_f32 v123, s0, 0x1234, v162
-// GFX12-NEXT:{{^}} ^
+// GFX12-NEXT:{{^}} ^
//===----------------------------------------------------------------------===//
// Check that assembler detects a different literal regardless of its location.
@@ -46,7 +46,7 @@ v_dual_fmamk_f32 v122, 0xdeadbeef, 0xdeadbeef, v161 :: v_dual_fmamk_f32 v12
v_dual_fmamk_f32 v122, 0xdeadbeef, 0xdeadbeef, v161 :: v_dual_fmamk_f32 v123, 0xdeadbeef, 0x1234, v162
// GFX12: :[[@LINE-1]]:{{[0-9]+}}: error: only one unique literal operand is allowed
// GFX12-NEXT:{{^}}v_dual_fmamk_f32 v122, 0xdeadbeef, 0xdeadbeef, v161 :: v_dual_fmamk_f32 v123, 0xdeadbeef, 0x1234, v162
-// GFX12-NEXT:{{^}} ^
+// GFX12-NEXT:{{^}} ^
v_dual_fmamk_f32 v122, 0xdeadbeef, 0xdeadbeef, v161 :: v_dual_fmamk_f32 v123, 0x1234, 0xdeadbeef, v162
// GFX12: :[[@LINE-1]]:{{[0-9]+}}: error: only one unique literal operand is allowed
@@ -56,7 +56,7 @@ v_dual_fmamk_f32 v122, 0xdeadbeef, 0xdeadbeef, v161 :: v_dual_fmamk_f32 v12
v_dual_fmamk_f32 v122, 0xdeadbeef, 0x1234, v161 :: v_dual_fmamk_f32 v123, 0xdeadbeef, 0xdeadbeef, v162
// GFX12: :[[@LINE-1]]:{{[0-9]+}}: error: only one unique literal operand is allowed
// GFX12-NEXT:{{^}}v_dual_fmamk_f32 v122, 0xdeadbeef, 0x1234, v161 :: v_dual_fmamk_f32 v123, 0xdeadbeef, 0xdeadbeef, v162
-// GFX12-NEXT:{{^}} ^
+// GFX12-NEXT:{{^}} ^
v_dual_fmamk_f32 v122, 0x1234, 0xdeadbeef, v161 :: v_dual_fmamk_f32 v123, 0xdeadbeef, 0xdeadbeef, v162
// GFX12: :[[@LINE-1]]:{{[0-9]+}}: error: only one unique literal operand is allowed
@@ -114,7 +114,7 @@ v_dual_cndmask_b32 v255, s1, v2 :: v_dual_cndmask_b32 v6, s2
v_dual_cndmask_b32 v1, s2, v3, vcc_lo :: v_dual_cndmask_b32 v2, s3, v4, vcc_lo
// GFX12: :[[@LINE-1]]:{{[0-9]+}}: error: invalid operand (violates constant bus restrictions)
// GFX12-NEXT:{{^}}v_dual_cndmask_b32 v1, s2, v3, vcc_lo :: v_dual_cndmask_b32 v2, s3, v4, vcc_lo
-// GFX12-NEXT:{{^}} ^
+// GFX12-NEXT:{{^}} ^
// SGPR + LITERAL + VCC
diff --git a/llvm/test/MC/AMDGPU/gfx12_asm_vopd_errs.s b/llvm/test/MC/AMDGPU/gfx12_asm_vopd_errs.s
index 5751258fe85d8..aa71e4d2969fc 100644
--- a/llvm/test/MC/AMDGPU/gfx12_asm_vopd_errs.s
+++ b/llvm/test/MC/AMDGPU/gfx12_asm_vopd_errs.s
@@ -22,7 +22,7 @@ v_dual_fmamk_f32 v122, v74, 0xa0172923, v161 :: v_dual_lshlrev_b32 v247, 0
v_dual_add_f32 v5, 0xaf123456, v2 :: v_dual_fmaak_f32 v6, v3, v1, 0xbabe
// GFX12: :[[@LINE-1]]:{{[0-9]+}}: error: only one unique literal operand is allowed
// GFX12-NEXT:{{^}}v_dual_add_f32 v5, 0xaf123456, v2 :: v_dual_fmaak_f32 v6, v3, v1, 0xbabe
-// GFX12-NEXT:{{^}} ^
+// GFX12-NEXT:{{^}} ^
v_dual_add_f32 v5, 0xaf123456, v2 :: v_dual_fmaak_f32 v6, 0xbabe, v1, 0xbabe
// GFX12: :[[@LINE-1]]:{{[0-9]+}}: error: only one unique literal operand is allowed
@@ -32,12 +32,12 @@ v_dual_add_f32 v5, 0xaf123456, v2 :: v_dual_fmaak_f32 v6, 0x
v_dual_fmamk_f32 v122, 0xdeadbeef, 0xdeadbeef, v161 :: v_dual_fmamk_f32 v123, 0xdeadbeef, 0x1234, v162
// GFX12: :[[@LINE-1]]:{{[0-9]+}}: error: only one unique literal operand is allowed
// GFX12-NEXT:{{^}}v_dual_fmamk_f32 v122, 0xdeadbeef, 0xdeadbeef, v161 :: v_dual_fmamk_f32 v123, 0xdeadbeef, 0x1234, v162
-// GFX12-NEXT:{{^}} ^
+// GFX12-NEXT:{{^}} ^
v_dual_fmamk_f32 v122, 0xdeadbeef, 0xdeadbeef, v161 :: v_dual_fmamk_f32 v123, s0, 0x1234, v162
// GFX12: :[[@LINE-1]]:{{[0-9]+}}: error: only one unique literal operand is allowed
// GFX12-NEXT:{{^}}v_dual_fmamk_f32 v122, 0xdeadbeef, 0xdeadbeef, v161 :: v_dual_fmamk_f32 v123, s0, 0x1234, v162
-// GFX12-NEXT:{{^}} ^
+// GFX12-NEXT:{{^}} ^
//===----------------------------------------------------------------------===//
// Check that assembler detects a different literal regardless of its location.
@@ -46,7 +46,7 @@ v_dual_fmamk_f32 v122, 0xdeadbeef, 0xdeadbeef, v161 :: v_dual_fmamk_f32 v12
v_dual_fmamk_f32 v122, 0xdeadbeef, 0xdeadbeef, v161 :: v_dual_fmamk_f32 v123, 0xdeadbeef, 0x1234, v162
// GFX12: :[[@LINE-1]]:{{[0-9]+}}: error: only one unique literal operand is allowed
// GFX12-NEXT:{{^}}v_dual_fmamk_f32 v122, 0xdeadbeef, 0xdeadbeef, v161 :: v_dual_fmamk_f32 v123, 0xdeadbeef, 0x1234, v162
-// GFX12-NEXT:{{^}} ^
+// GFX12-NEXT:{{^}} ^
v_dual_fmamk_f32 v122, 0xdeadbeef, 0xdeadbeef, v161 :: v_dual_fmamk_f32 v123, 0x1234, 0xdeadbeef, v162
// GFX12: :[[@LINE-1]]:{{[0-9]+}}: error: only one unique literal operand is allowed
@@ -56,7 +56,7 @@ v_dual_fmamk_f32 v122, 0xdeadbeef, 0xdeadbeef, v161 :: v_dual_fmamk_f32 v12
v_dual_fmamk_f32 v122, 0xdeadbeef, 0x1234, v161 :: v_dual_fmamk_f32 v123, 0xdeadbeef, 0xdeadbeef, v162
// GFX12: :[[@LINE-1]]:{{[0-9]+}}: error: only one unique literal operand is allowed
// GFX12-NEXT:{{^}}v_dual_fmamk_f32 v122, 0xdeadbeef, 0x1234, v161 :: v_dual_fmamk_f32 v123, 0xdeadbeef, 0xdeadbeef, v162
-// GFX12-NEXT:{{^}} ^
+// GFX12-NEXT:{{^}} ^
v_dual_fmamk_f32 v122, 0x1234, 0xdeadbeef, v161 :: v_dual_fmamk_f32 v123, 0xdeadbeef, 0xdeadbeef, v162
// GFX12: :[[@LINE-1]]:{{[0-9]+}}: error: only one unique literal operand is allowed
diff --git a/llvm/test/MC/AMDGPU/mai-gfx950-err.s b/llvm/test/MC/AMDGPU/mai-gfx950-err.s
index 5c9dbd7f7636f..747deab3bfcae 100644
--- a/llvm/test/MC/AMDGPU/mai-gfx950-err.s
+++ b/llvm/test/MC/AMDGPU/mai-gfx950-err.s
@@ -7,7 +7,9 @@ v_mfma_ld_scale_b32 65, v0
// CHECK: :[[@LINE-1]]:21: error: literal operands are not supported
v_mfma_ld_scale_b32 65, 65
-// CHECK: :[[@LINE-1]]:25: error: literal operands are not supported
+// CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: literal operands are not supported
+// CHECK-NEXT:{{^}}v_mfma_ld_scale_b32 65, 65
+// CHECK-NEXT:{{^}} ^
v_mfma_ld_scale_b32 s0, s1
// CHECK: :[[@LINE-1]]:25: error: invalid operand (violates constant bus restrictions)
diff --git a/llvm/test/MC/AMDGPU/vop3-literal.s b/llvm/test/MC/AMDGPU/vop3-literal.s
index ba683e5423c72..c1f1fd69042c1 100644
--- a/llvm/test/MC/AMDGPU/vop3-literal.s
+++ b/llvm/test/MC/AMDGPU/vop3-literal.s
@@ -26,25 +26,25 @@ v_bfe_u32 v0, v1, s1, 0x3039
v_bfe_u32 v0, 0x3039, 0x3039, s1
// GFX10: v_bfe_u32 v0, 0x3039, 0x3039, s1 ; encoding: [0x00,0x00,0x48,0xd5,0xff,0xfe,0x05,0x00,0x39,0x30,0x00,0x00]
// GFX1250: v_bfe_u32 v0, 0x3039, 0x3039, s1 ; encoding: [0x00,0x00,0x10,0xd6,0xff,0xfe,0x05,0x00,0x39,0x30,0x00,0x00]
-// GFX9-ERR: :[[@LINE-3]]:23: error: literal operands are not supported
+// GFX9-ERR: :[[@LINE-3]]:15: error: literal operands are not supported
v_bfe_u32 v0, 0x3039, s1, 0x3039
// GFX10: v_bfe_u32 v0, 0x3039, s1, 0x3039 ; encoding: [0x00,0x00,0x48,0xd5,0xff,0x02,0xfc,0x03,0x39,0x30,0x00,0x00]
// GFX1250: v_bfe_u32 v0, 0x3039, s1, 0x3039 ; encoding: [0x00,0x00,0x10,0xd6,0xff,0x02,0xfc,0x03,0x39,0x30,0x00,0x00]
-// GFX9-ERR: :[[@LINE-3]]:27: error: literal operands are not supported
+// GFX9-ERR: :[[@LINE-3]]:15: error: literal operands are not supported
v_bfe_u32 v0, v1, 0x3039, 0x3039
// GFX10: v_bfe_u32 v0, v1, 0x3039, 0x3039 ; encoding: [0x00,0x00,0x48,0xd5,0x01,0xff,0xfd,0x03,0x39,0x30,0x00,0x00]
// GFX1250: v_bfe_u32 v0, v1, 0x3039, 0x3039 ; encoding: [0x00,0x00,0x10,0xd6,0x01,0xff,0xfd,0x03,0x39,0x30,0x00,0x00]
-// GFX9-ERR: :[[@LINE-3]]:27: error: literal operands are not supported
+// GFX9-ERR: :[[@LINE-3]]:19: error: literal operands are not supported
v_bfe_u32 v0, 0x3039, 0x3039, 0x3039
// GFX10: v_bfe_u32 v0, 0x3039, 0x3039, 0x3039 ; encoding: [0x00,0x00,0x48,0xd5,0xff,0xfe,0xfd,0x03,0x39,0x30,0x00,0x00]
// GFX1250: v_bfe_u32 v0, 0x3039, 0x3039, 0x3039 ; encoding: [0x00,0x00,0x10,0xd6,0xff,0xfe,0xfd,0x03,0x39,0x30,0x00,0x00]
-// GFX9-ERR: :[[@LINE-3]]:31: error: literal operands are not supported
+// GFX9-ERR: :[[@LINE-3]]:15: error: literal operands are not supported
v_bfe_u32 v0, 0x3039, s1, 0x3038
-// GFX9-ERR: :[[@LINE-1]]:27: error: literal operands are not supported
+// GFX9-ERR: :[[@LINE-1]]:15: error: literal operands are not supported
// GFX10-ERR: :[[@LINE-2]]:27: error: only one unique literal operand is allowed
// GFX1250-ERR: :[[@LINE-3]]:27: error: only one unique literal operand is allowed
@@ -54,7 +54,7 @@ v_bfe_u32 v0, 0x3039, v1, v2
// GFX9-ERR: :[[@LINE-3]]:15: error: literal operands are not supported
v_bfe_u32 v0, 0x3039, 0x12345, v2
-// GFX9-ERR: :[[@LINE-1]]:23: error: literal operands are not supported
+// GFX9-ERR: :[[@LINE-1]]:15: error: literal operands are not supported
// GFX10-ERR: :[[@LINE-2]]:23: error: only one unique literal operand is allowed
// GFX1250-ERR: :[[@LINE-3]]:23: error: only one unique literal operand is allowed
@@ -81,10 +81,10 @@ v_bfm_b32_e64 v0, 0x3039, v1
v_bfm_b32_e64 v0, 0x3039, 0x3039
// GFX10: v_bfm_b32 v0, 0x3039, 0x3039 ; encoding: [0x00,0x00,0x63,0xd7,0xff,0xfe,0x01,0x00,0x39,0x30,0x00,0x00]
// GFX1250: v_bfm_b32 v0, 0x3039, 0x3039 ; encoding: [0x00,0x00,0x1d,0xd7,0xff,0xfe,0x01,0x00,0x39,0x30,0x00,0x00]
-// GFX9-ERR: :[[@LINE-3]]:27: error: literal operands are not supported
+// GFX9-ERR: :[[@LINE-3]]:19: error: literal operands are not supported
v_bfm_b32_e64 v0, 0x3039, 0x3038
-// GFX9-ERR: :[[@LINE-1]]:27: error: literal operands are not supported
+// GFX9-ERR: :[[@LINE-1]]:19: error: literal operands are not supported
// GFX10-ERR: :[[@LINE-2]]:27: error: only one unique literal operand is allowed
// GFX1250-ERR: :[[@LINE-3]]:27: error: only one unique literal operand is allowed
@@ -106,10 +106,10 @@ v_pk_add_f16 v1, -200, v2
v_pk_add_f16 v1, 25.0, 25.0
// GFX10: v_pk_add_f16 v1, 0x4e40, 0x4e40 ; encoding: [0x01,0x40,0x0f,0xcc,0xff,0xfe,0x01,0x18,0x40,0x4e,0x00,0x00]
// GFX1250: v_pk_add_f16 v1, 0x4e40, 0x4e40 ; encoding: [0x01,0x40,0x0f,0xcc,0xff,0xfe,0x01,0x18,0x40,0x4e,0x00,0x00]
-// GFX9-ERR: :[[@LINE-3]]:24: error: literal operands are not supported
+// GFX9-ERR: :[[@LINE-3]]:18: error: literal operands are not supported
v_pk_add_f16 v1, 25.0, 25.1
-// GFX9-ERR: :[[@LINE-1]]:24: error: literal operands are not supported
+// GFX9-ERR: :[[@LINE-1]]:18: error: literal operands are not supported
// GFX10-ERR: :[[@LINE-2]]:24: error: only one unique literal operand is allowed
// GFX1250-ERR: :[[@LINE-3]]:24: error: only one unique literal operand is allowed
@@ -146,7 +146,7 @@ v_pk_add_u16 v1, -100, v2
v_pk_add_u16 v1, -100, -100
// GFX10: v_pk_add_u16 v1, 0xffffff9c, 0xffffff9c ; encoding: [0x01,0x40,0x0a,0xcc,0xff,0xfe,0x01,0x18,0x9c,0xff,0xff,0xff]
// GFX1250: v_pk_add_u16 v1, 0xffffff9c, 0xffffff9c ; encoding: [0x01,0x40,0x0a,0xcc,0xff,0xfe,0x01,0x18,0x9c,0xff,0xff,0xff]
-// GFX9-ERR: :[[@LINE-3]]:24: error: literal operands are not supported
+// GFX9-ERR: :[[@LINE-3]]:18: error: literal operands are not supported
v_add_f32_e64 v1, neg(abs(0x123)), v3
// GFX10: v_add_f32_e64 v1, -|0x123|, v3 ; encoding: [0x01,0x01,0x03,0xd5,0xff,0x06,0x02,0x20,0x23,0x01,0x00,0x00]
@@ -161,7 +161,7 @@ v_add_f32_e64 v1, v3, neg(0x123)
v_add_f32_e64 v1, neg(abs(0x12345678)), neg(0x12345678)
// GFX10: v_add_f32_e64 v1, -|0x12345678|, neg(0x12345678) ; encoding: [0x01,0x01,0x03,0xd5,0xff,0xfe,0x01,0x60,0x78,0x56,0x34,0x12]
// GFX1250: v_add_f32_e64 v1, -|0x12345678|, neg(0x12345678) ; encoding: [0x01,0x01,0x03,0xd5,0xff,0xfe,0x01,0x60,0x78,0x56,0x34,0x12]
-// GFX9-ERR: :[[@LINE-3]]:45: error: literal operands are not supported
+// GFX9-ERR: :[[@LINE-3]]:27: error: literal operands are not supported
v_add_f16_e64 v0, v0, 0xfe0b
// GFX10: v_add_f16_e64 v0, v0, 0xfe0b ; encoding: [0x00,0x00,0x32,0xd5,0x00,0xff,0x01,0x00,0x0b,0xfe,0x00,0x00]
@@ -181,7 +181,7 @@ v_add_f16_e64 v0, 0x3456, v0
v_add_f16_e64 v0, 0xfe0b, neg(0xfe0b)
// GFX10: v_add_f16_e64 v0, 0xfe0b, neg(0xfe0b) ; encoding: [0x00,0x00,0x32,0xd5,0xff,0xfe,0x01,0x40,0x0b,0xfe,0x00,0x00]
// GFX1250: v_add_f16_e64 v0, 0xfe0b, neg(0xfe0b) ; encoding: [0x00,0x00,0x32,0xd5,0xff,0xfe,0x01,0x40,0x0b,0xfe,0x00,0x00]
-// GFX9-ERR: :[[@LINE-3]]:31: error: literal operands are not supported
+// GFX9-ERR: :[[@LINE-3]]:19: error: literal operands are not supported
v_add_f64 v[0:1], 1.23456, v[0:1]
// GFX10: v_add_f64 v[0:1], 0x3ff3c0c1, v[0:1] ; encoding: [0x00,0x00,0x64,0xd5,0xff,0x00,0x02,0x00,0xc1,0xc0,0xf3,0x3f]
@@ -196,10 +196,10 @@ v_add_f64 v[0:1], v[0:1], -abs(1.23456)
v_add_f64 v[0:1], 1.23456, -abs(1.23456)
// GFX10: v_add_f64 v[0:1], 0x3ff3c0c1, -|0x3ff3c0c1| ; encoding: [0x00,0x02,0x64,0xd5,0xff,0xfe,0x01,0x40,0xc1,0xc0,0xf3,0x3f]
// GFX1250: v_add_f64_e64 v[0:1], 0x3ff3c0c1, -|0x3ff3c0c1| ; encoding: [0x00,0x02,0x02,0xd5,0xff,0xfe,0x01,0x40,0xc1,0xc0,0xf3,0x3f]
-// GFX9-ERR: :[[@LINE-3]]:33: error: literal operands are not supported
+// GFX9-ERR: :[[@LINE-3]]:19: error: literal operands are not supported
v_add_f64 v[0:1], 1.23456, -abs(1.2345)
-// GFX9-ERR: :[[@LINE-1]]:33: error: literal operands are not supported
+// GFX9-ERR: :[[@LINE-1]]:19: error: literal operands are not supported
// GFX10-ERR: :[[@LINE-2]]:33: error: only one unique literal operand is allowed
// GFX1250-ERR: :[[@LINE-3]]:33: error: only one unique literal operand is allowed
@@ -216,7 +216,7 @@ v_max_i16_e64 v5, v1, 0x123
v_max_i16_e64 v5, 0x1234, 0x1234
// GFX10: v_max_i16 v5, 0x1234, 0x1234 ; encoding: [0x05,0x00,0x0a,0xd7,0xff,0xfe,0x01,0x00,0x34,0x12,0x00,0x00]
// GFX1250: v_max_i16 v5, 0x1234, 0x1234 ; encoding: [0x05,0x00,0x0a,0xd7,0xff,0xfe,0x01,0x00,0x34,0x12,0x00,0x00]
-// GFX9-ERR: :[[@LINE-3]]:27: error: literal operands are not supported
+// GFX9-ERR: :[[@LINE-3]]:19: error: literal operands are not supported
v_min3_i16 v5, 0xfe0b, v2, v3
// GFX10: v_min3_i16 v5, 0xfe0b, v2, v3 ; encoding: [0x05,0x00,0x52,0xd7,0xff,0x04,0x0e,0x04,0x0b,0xfe,0x00,0x00]
@@ -236,15 +236,15 @@ v_min3_i16 v5, v1, v2, 0x5678
v_min3_i16 v5, 0x5678, 0x5678, 0x5678
// GFX10: v_min3_i16 v5, 0x5678, 0x5678, 0x5678 ; encoding: [0x05,0x00,0x52,0xd7,0xff,0xfe,0xfd,0x03,0x78,0x56,0x00,0x00]
// GFX1250: v_min3_i16 v5, 0x5678, 0x5678, 0x5678 ; encoding: [0x05,0x00,0x4a,0xd6,0xff,0xfe,0xfd,0x03,0x78,0x56,0x00,0x00]
-// GFX9-ERR: :[[@LINE-3]]:32: error: literal operands are not supported
+// GFX9-ERR: :[[@LINE-3]]:16: error: literal operands are not supported
v_min3_i16 v5, 0x5678, 0x5679, 0x5678
-// GFX9-ERR: :[[@LINE-1]]:32: error: literal operands are not supported
-// GFX10-ERR: :[[@LINE-2]]:32: error: only one unique literal operand is allowed
-// GFX1250-ERR: :[[@LINE-3]]:32: error: only one unique literal operand is allowed
+// GFX9-ERR: :[[@LINE-1]]:16: error: literal operands are not supported
+// GFX10-ERR: :[[@LINE-2]]:24: error: only one unique literal operand is allowed
+// GFX1250-ERR: :[[@LINE-3]]:24: error: only one unique literal operand is allowed
v_min3_i16 v5, 0x5678, 0x5678, 0x5679
-// GFX9-ERR: :[[@LINE-1]]:32: error: literal operands are not supported
+// GFX9-ERR: :[[@LINE-1]]:16: error: literal operands are not supported
// GFX10-ERR: :[[@LINE-2]]:32: error: only one unique literal operand is allowed
// GFX1250-ERR: :[[@LINE-3]]:32: error: only one unique literal operand is allowed
@@ -286,7 +286,7 @@ v_mad_u16 v5, v1, v2, 0x5678
v_mad_u16 v5, 0x5678, 0x5678, 0x5678
// GFX10: v_mad_u16 v5, 0x5678, 0x5678, 0x5678 ; encoding: [0x05,0x00,0x40,0xd7,0xff,0xfe,0xfd,0x03,0x78,0x56,0x00,0x00]
// GFX1250: v_mad_u16 v5, 0x5678, 0x5678, 0x5678 ; encoding: [0x05,0x00,0x41,0xd6,0xff,0xfe,0xfd,0x03,0x78,0x56,0x00,0x00]
-// GFX9-ERR: :[[@LINE-3]]:31: error: literal operands are not supported
+// GFX9-ERR: :[[@LINE-3]]:15: error: literal operands are not supported
v_mad_legacy_f32 v5, 0xaf123456, v2, v3
// GFX10: v_mad_legacy_f32 v5, 0xaf123456, v2, v3 ; encoding: [0x05,0x00,0x40,0xd5,0xff,0x04,0x0e,0x04,0x56,0x34,0x12,0xaf]
@@ -305,7 +305,7 @@ v_mad_legacy_f32 v5, v1, v2, 0xaf123456
v_mad_legacy_f32 v5, 0xaf123456, 0xaf123456, 0xaf123456
// GFX10: v_mad_legacy_f32 v5, 0xaf123456, 0xaf123456, 0xaf123456 ; encoding: [0x05,0x00,0x40,0xd5,0xff,0xfe,0xfd,0x03,0x56,0x34,0x12,0xaf]
-// GFX9-ERR: :[[@LINE-2]]:46: error: literal operands are not supported
+// GFX9-ERR: :[[@LINE-2]]:22: error: literal operands are not supported
// GFX1250-ERR: :[[@LINE-3]]:1: error: instruction not supported on this GPU
v_cmp_eq_i32_e64 s[10:11], 0xaf123456, v2
@@ -321,10 +321,10 @@ v_cmp_eq_i32_e64 s[10:11], v1, 0xaf123456
v_cmp_eq_i32_e64 s[10:11], 0xaf123456, 0xaf123456
// GFX10: v_cmp_eq_i32_e64 s[10:11], 0xaf123456, 0xaf123456 ; encoding: [0x0a,0x00,0x82,0xd4,0xff,0xfe,0x01,0x00,0x56,0x34,0x12,0xaf]
// GFX1250: v_cmp_eq_i32_e64 s[10:11], 0xaf123456, 0xaf123456 ; encoding: [0x0a,0x00,0x42,0xd4,0xff,0xfe,0x01,0x00,0x56,0x34,0x12,0xaf]
-// GFX9-ERR: :[[@LINE-3]]:40: error: literal operands are not supported
+// GFX9-ERR: :[[@LINE-3]]:28: error: literal operands are not supported
v_cmp_eq_i32_e64 s[10:11], 0xaf123456, 0xaf123455
-// GFX9-ERR: :[[@LINE-1]]:40: error: literal operands are not supported
+// GFX9-ERR: :[[@LINE-1]]:28: error: literal operands are not supported
// GFX10-ERR: :[[@LINE-2]]:40: error: only one unique literal operand is allowed
// GFX1250-ERR: :[[@LINE-3]]:40: error: only one unique literal operand is allowed
@@ -341,7 +341,7 @@ v_cmp_eq_u64_e64 s[10:11], v[2:3], 0x3f717273
v_cmp_eq_u64_e64 s[10:11], 0x3f717273, 0x3f717273
// GFX10: v_cmp_eq_u64_e64 s[10:11], 0x3f717273, 0x3f717273 ; encoding: [0x0a,0x00,0xe2,0xd4,0xff,0xfe,0x01,0x00,0x73,0x72,0x71,0x3f]
// GFX1250: v_cmp_eq_u64_e64 s[10:11], 0x3f717273, 0x3f717273 ; encoding: [0x0a,0x00,0x5a,0xd4,0xff,0xfe,0x01,0x00,0x73,0x72,0x71,0x3f]
-// GFX9-ERR: :[[@LINE-3]]:40: error: literal operands are not supported
+// GFX9-ERR: :[[@LINE-3]]:28: error: literal operands are not supported
v_cmpx_class_f32_e64 0xaf123456, v2
// GFX10: v_cmpx_class_f32_e64 0xaf123456, v2 ; encoding: [0x7e,0x00,0x98,0xd4,0xff,0x04,0x02,0x00,0x56,0x34,0x12,0xaf]
@@ -441,7 +441,7 @@ v_pk_add_f16 v5, v1, 0xbf717273
v_pk_add_f16 v5, 0x3f717273, 0x3f717273
// GFX10: v_pk_add_f16 v5, 0x3f717273, 0x3f717273 ; encoding: [0x05,0x40,0x0f,0xcc,0xff,0xfe,0x01,0x18,0x73,0x72,0x71,0x3f]
// GFX1250: v_pk_add_f16 v5, 0x3f717273, 0x3f717273 ; encoding: [0x05,0x40,0x0f,0xcc,0xff,0xfe,0x01,0x18,0x73,0x72,0x71,0x3f]
-// GFX9-ERR: :[[@LINE-3]]:30: error: literal operands are not supported
+// GFX9-ERR: :[[@LINE-3]]:18: error: literal operands are not supported
v_pk_add_i16 v5, 0x7b, v2
// GFX10: v_pk_add_i16 v5, 0x7b, v2 ; encoding: [0x05,0x40,0x02,0xcc,0xff,0x04,0x02,0x18,0x7b,0x00,0x00,0x00]
@@ -456,10 +456,10 @@ v_pk_add_i16 v5, v1, 0x7b
v_pk_add_i16 v5, 0xab7b, 0xab7b
// GFX10: v_pk_add_i16 v5, 0xab7b, 0xab7b ; encoding: [0x05,0x40,0x02,0xcc,0xff,0xfe,0x01,0x18,0x7b,0xab,0x00,0x00]
// GFX1250: v_pk_add_i16 v5, 0xab7b, 0xab7b ; encoding: [0x05,0x40,0x02,0xcc,0xff,0xfe,0x01,0x18,0x7b,0xab,0x00,0x00]
-// GFX9-ERR: :[[@LINE-3]]:26: error: literal operands are not supported
+// GFX9-ERR: :[[@LINE-3]]:18: error: literal operands are not supported
v_pk_add_i16 v5, 0xab7b, 0xab7a
-// GFX9-ERR: :[[@LINE-1]]:26: error: literal operands are not supported
+// GFX9-ERR: :[[@LINE-1]]:18: error: literal operands are not supported
// GFX10-ERR: :[[@LINE-2]]:26: error: only one unique literal operand is allowed
// GFX1250-ERR: :[[@LINE-3]]:26: error: only one unique literal operand is allowed
@@ -471,12 +471,12 @@ v_div_fmas_f32 v5, v1, 0x123, v3
v_div_fmas_f32 v5, v1, 0x123, 0x123
// GFX10: v_div_fmas_f32 v5, v1, 0x123, 0x123 ; encoding: [0x05,0x00,0x6f,0xd5,0x01,0xff,0xfd,0x03,0x23,0x01,0x00,0x00]
// GFX1250: v_div_fmas_f32 v5, v1, 0x123, 0x123 ; encoding: [0x05,0x00,0x37,0xd6,0x01,0xff,0xfd,0x03,0x23,0x01,0x00,0x00]
-// GFX9-ERR: :[[@LINE-3]]:31: error: literal operands are not supported
+// GFX9-ERR: :[[@LINE-3]]:24: error: literal operands are not supported
v_div_fmas_f32 v5, 0x123, 0x123, 0x123
// GFX10: v_div_fmas_f32 v5, 0x123, 0x123, 0x123 ; encoding: [0x05,0x00,0x6f,0xd5,0xff,0xfe,0xfd,0x03,0x23,0x01,0x00,0x00]
// GFX1250: v_div_fmas_f32 v5, 0x123, 0x123, 0x123 ; encoding: [0x05,0x00,0x37,0xd6,0xff,0xfe,0xfd,0x03,0x23,0x01,0x00,0x00]
-// GFX9-ERR: :[[@LINE-3]]:34: error: literal operands are not supported
+// GFX9-ERR: :[[@LINE-3]]:20: error: literal operands are not supported
v_div_fmas_f64 v[4:5], 0x12345678, v[2:3], v[4:5]
// GFX10: v_div_fmas_f64 v[4:5], 0x12345678, v[2:3], v[4:5] ; encoding: [0x04,0x00,0x70,0xd5,0xff,0x04,0x12,0x04,0x78,0x56,0x34,0x12]
@@ -486,10 +486,10 @@ v_div_fmas_f64 v[4:5], 0x12345678, v[2:3], v[4:5]
v_div_fmas_f64 v[6:7], 0x12345678, 0x12345678, 0x12345678
// GFX10: v_div_fmas_f64 v[6:7], 0x12345678, 0x12345678, 0x12345678 ; encoding: [0x06,0x00,0x70,0xd5,0xff,0xfe,0xfd,0x03,0x78,0x56,0x34,0x12]
// GFX1250: v_div_fmas_f64 v[6:7], 0x12345678, 0x12345678, 0x12345678 ; encoding: [0x06,0x00,0x38,0xd6,0xff,0xfe,0xfd,0x03,0x78,0x56,0x34,0x12]
-// GFX9-ERR: :[[@LINE-3]]:48: error: literal operands are not supported
+// GFX9-ERR: :[[@LINE-3]]:24: error: literal operands are not supported
-v_div_fmas_f64 v[4:5], v[2:3], 0x123457, 0x123456
-// GFX9-ERR: :[[@LINE-1]]:42: error: literal operands are not supported
+v_div_fmas_f64 v[4:5], v[1:2], 0x123457, 0x123456
+// GFX9-ERR: :[[@LINE-1]]:32: error: literal operands are not supported
// GFX10-ERR: :[[@LINE-2]]:42: error: only one unique literal operand is allowed
// GFX1250-ERR: :[[@LINE-3]]:42: error: only one unique literal operand is allowed
@@ -501,9 +501,9 @@ v_ldexp_f64 v[4:5], 0.12345, v2
v_ldexp_f64 v[6:7], 0.12345, 0x3fbf9a6b
// GFX10: v_ldexp_f64 v[6:7], 0x3fbf9a6b, 0x3fbf9a6b ; encoding: [0x06,0x00,0x68,0xd5,0xff,0xfe,0x01,0x00,0x6b,0x9a,0xbf,0x3f]
// GFX1250: v_ldexp_f64 v[6:7], 0x3fbf9a6b, 0x3fbf9a6b ; encoding: [0x06,0x00,0x2b,0xd7,0xff,0xfe,0x01,0x00,0x6b,0x9a,0xbf,0x3f]
-// GFX9-ERR: :[[@LINE-3]]:30: error: literal operands are not supported
+// GFX9-ERR: :[[@LINE-3]]:21: error: literal operands are not supported
v_ldexp_f64 v[4:5], 0.12345, 0x3fbf9a6c
-// GFX9-ERR: :[[@LINE-1]]:30: error: literal operands are not supported
+// GFX9-ERR: :[[@LINE-1]]:21: error: literal operands are not supported
// GFX10-ERR: :[[@LINE-2]]:30: error: only one unique literal operand is allowed
// GFX1250-ERR: :[[@LINE-3]]:30: error: only one unique literal operand is allowed
More information about the llvm-commits
mailing list