From 65df61f2bbbd63af3f6a207188964e4bad95c77c Mon Sep 17 00:00:00 2001 From: Ivan Kosarev Date: Fri, 12 Sep 2025 15:50:20 +0100 Subject: [PATCH 1/2] [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(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 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 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 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(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(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(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(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 @@ -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 -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 LiteralOpIdx; + std::optional 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(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(*Op); + if (TargetOp.getMCOpIdx() == MCOpIdx) + return TargetOp.getStartLoc(); + } + llvm_unreachable("No such MC operand!"); +} + SMLoc AMDGPUAsmParser::getOperandLoc(std::function 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 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 From 0747d0771b9b9e54167de4ec81da2fdfddaab223 Mon Sep 17 00:00:00 2001 From: Ivan Kosarev Date: Mon, 15 Sep 2025 11:41:14 +0100 Subject: [PATCH 2/2] Give getLaterLoc() a comment and remove an unwanted test change. --- llvm/lib/Target/AMDGPU/AsmParser/AMDGPUAsmParser.cpp | 1 + llvm/test/MC/AMDGPU/vop3-literal.s | 2 +- 2 files changed, 2 insertions(+), 1 deletion(-) diff --git a/llvm/lib/Target/AMDGPU/AsmParser/AMDGPUAsmParser.cpp b/llvm/lib/Target/AMDGPU/AsmParser/AMDGPUAsmParser.cpp index 15d03403e3868..fd9c045e06804 100644 --- a/llvm/lib/Target/AMDGPU/AsmParser/AMDGPUAsmParser.cpp +++ b/llvm/lib/Target/AMDGPU/AsmParser/AMDGPUAsmParser.cpp @@ -8499,6 +8499,7 @@ SMLoc AMDGPUAsmParser::getInstLoc(const OperandVector &Operands) const { return ((AMDGPUOperand &)*Operands[0]).getStartLoc(); } +// Returns one of the given locations that comes later in the source. SMLoc AMDGPUAsmParser::getLaterLoc(SMLoc a, SMLoc b) { return a.getPointer() < b.getPointer() ? b : a; } diff --git a/llvm/test/MC/AMDGPU/vop3-literal.s b/llvm/test/MC/AMDGPU/vop3-literal.s index c1f1fd69042c1..56e71b9cfcfd1 100644 --- a/llvm/test/MC/AMDGPU/vop3-literal.s +++ b/llvm/test/MC/AMDGPU/vop3-literal.s @@ -488,7 +488,7 @@ v_div_fmas_f64 v[6:7], 0x12345678, 0x12345678, 0x12345678 // 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]]:24: error: literal operands are not supported -v_div_fmas_f64 v[4:5], v[1:2], 0x123457, 0x123456 +v_div_fmas_f64 v[4:5], v[2:3], 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