diff --git a/llvm/lib/Target/AMDGPU/AsmParser/AMDGPUAsmParser.cpp b/llvm/lib/Target/AMDGPU/AsmParser/AMDGPUAsmParser.cpp index e420f2ad676f9..fd9c045e06804 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,21 @@ 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; +} + +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 +8532,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 +10137,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..56e71b9cfcfd1 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 +// 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