diff --git a/llvm/lib/Target/AMDGPU/AsmParser/AMDGPUAsmParser.cpp b/llvm/lib/Target/AMDGPU/AsmParser/AMDGPUAsmParser.cpp --- a/llvm/lib/Target/AMDGPU/AsmParser/AMDGPUAsmParser.cpp +++ b/llvm/lib/Target/AMDGPU/AsmParser/AMDGPUAsmParser.cpp @@ -188,6 +188,12 @@ ImmTyEndpgm, }; + enum ImmKindTy { + ImmKindTyNone, + ImmKindTyLiteral, + ImmKindTyConst, + }; + private: struct TokOp { const char *Data; @@ -198,6 +204,7 @@ int64_t Val; ImmTy Type; bool IsFPImm; + mutable ImmKindTy Kind; Modifiers Mods; }; @@ -233,6 +240,29 @@ return Kind == Immediate; } + void setImmKindNone() const { + assert(isImm()); + Imm.Kind = ImmKindTyNone; + } + + void setImmKindLiteral() const { + assert(isImm()); + Imm.Kind = ImmKindTyLiteral; + } + + void setImmKindConst() const { + assert(isImm()); + Imm.Kind = ImmKindTyConst; + } + + bool IsImmKindLiteral() const { + return isImm() && Imm.Kind == ImmKindTyLiteral; + } + + bool isImmKindConst() const { + return isImm() && Imm.Kind == ImmKindTyConst; + } + bool isInlinableImm(MVT type) const; bool isLiteralImm(MVT type) const; @@ -911,6 +941,7 @@ 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; @@ -1358,6 +1389,8 @@ const OperandVector &Operands) const; SMLoc getImmLoc(AMDGPUOperand::ImmTy Type, const OperandVector &Operands) const; SMLoc getRegLoc(unsigned Reg, const OperandVector &Operands) const; + SMLoc getLitLoc(const OperandVector &Operands) const; + SMLoc getConstLoc(const OperandVector &Operands) const; bool validateInstruction(const MCInst &Inst, const SMLoc &IDLoc, const OperandVector &Operands); bool validateFlatOffset(const MCInst &Inst, const OperandVector &Operands); @@ -1376,7 +1409,7 @@ bool validateLdsDirect(const MCInst &Inst); bool validateOpSel(const MCInst &Inst); bool validateVccOperand(unsigned Reg) const; - bool validateVOP3Literal(const MCInst &Inst) const; + bool validateVOP3Literal(const MCInst &Inst, const OperandVector &Operands); bool validateMAIAccWrite(const MCInst &Inst); bool validateDivScale(const MCInst &Inst); bool validateCoherencyBits(const MCInst &Inst, const OperandVector &Operands, @@ -1772,6 +1805,7 @@ } else { assert(!isImmTy(ImmTyNone) || !hasModifiers()); Inst.addOperand(MCOperand::createImm(Imm.Val)); + setImmKindNone(); } } @@ -1799,6 +1833,7 @@ if (AMDGPU::isInlinableLiteral64(Literal.getZExtValue(), AsmParser->hasInv2PiInlineImm())) { Inst.addOperand(MCOperand::createImm(Literal.getZExtValue())); + setImmKindConst(); return; } @@ -1812,6 +1847,7 @@ } Inst.addOperand(MCOperand::createImm(Literal.lshr(32).getZExtValue())); + setImmKindLiteral(); return; } @@ -1848,6 +1884,7 @@ uint64_t ImmVal = FPLiteral.bitcastToAPInt().getZExtValue(); Inst.addOperand(MCOperand::createImm(ImmVal)); + setImmKindLiteral(); return; } default: @@ -1872,10 +1909,12 @@ AMDGPU::isInlinableLiteral32(static_cast(Val), AsmParser->hasInv2PiInlineImm())) { Inst.addOperand(MCOperand::createImm(Val)); + setImmKindConst(); return; } Inst.addOperand(MCOperand::createImm(Val & 0xffffffff)); + setImmKindLiteral(); return; case AMDGPU::OPERAND_REG_IMM_INT64: @@ -1884,10 +1923,12 @@ case AMDGPU::OPERAND_REG_INLINE_C_FP64: if (AMDGPU::isInlinableLiteral64(Val, AsmParser->hasInv2PiInlineImm())) { Inst.addOperand(MCOperand::createImm(Val)); + setImmKindConst(); return; } Inst.addOperand(MCOperand::createImm(Lo_32(Val))); + setImmKindLiteral(); return; case AMDGPU::OPERAND_REG_IMM_INT16: @@ -1900,10 +1941,12 @@ AMDGPU::isInlinableLiteral16(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: @@ -1925,6 +1968,7 @@ template void AMDGPUOperand::addKImmFPOperands(MCInst &Inst, unsigned N) const { APInt Literal(64, Imm.Val); + setImmKindNone(); if (!Imm.IsFPImm) { // We got int literal token. @@ -3699,7 +3743,8 @@ } // VOP3 literal is only allowed in GFX10+ and only one can be used -bool AMDGPUAsmParser::validateVOP3Literal(const MCInst &Inst) const { +bool AMDGPUAsmParser::validateVOP3Literal(const MCInst &Inst, + const OperandVector &Operands) { unsigned Opcode = Inst.getOpcode(); const MCInstrDesc &Desc = MII.get(Opcode); if (!(Desc.TSFlags & (SIInstrFlags::VOP3 | SIInstrFlags::VOP3P))) @@ -3725,8 +3770,11 @@ continue; if (OpIdx == Src2Idx && (Desc.TSFlags & SIInstrFlags::IsMAI) && - getFeatureBits()[AMDGPU::FeatureMFMAInlineLiteralBug]) + getFeatureBits()[AMDGPU::FeatureMFMAInlineLiteralBug]) { + Error(getConstLoc(Operands), + "inline constants are not allowed for this operand"); return false; + } if (MO.isImm() && !isInlineConstant(Inst, OpIdx)) { uint32_t Value = static_cast(MO.getImm()); @@ -3740,8 +3788,20 @@ } NumLiterals += NumExprs; - return !NumLiterals || - (NumLiterals == 1 && getFeatureBits()[AMDGPU::FeatureVOP3Literal]); + if (!NumLiterals) + return true; + + if (!getFeatureBits()[AMDGPU::FeatureVOP3Literal]) { + Error(getLitLoc(Operands), "literal operands are not supported"); + return false; + } + + if (NumLiterals > 1) { + Error(getLitLoc(Operands), "only one literal operand is allowed"); + return false; + } + + return true; } bool AMDGPUAsmParser::validateCoherencyBits(const MCInst &Inst, @@ -3770,13 +3830,11 @@ return false; } if (!validateSOPLiteral(Inst)) { - Error(IDLoc, + Error(getLitLoc(Operands), "only one literal operand is allowed"); return false; } - if (!validateVOP3Literal(Inst)) { - Error(IDLoc, - "invalid literal operand"); + if (!validateVOP3Literal(Inst, Operands)) { return false; } if (!validateConstantBusLimitations(Inst)) { @@ -6112,6 +6170,22 @@ return getOperandLoc(Test, Operands); } +SMLoc +AMDGPUAsmParser::getLitLoc(const OperandVector &Operands) const { + auto Test = [](const AMDGPUOperand& Op) { + return Op.IsImmKindLiteral() || Op.isExpr(); + }; + return getOperandLoc(Test, Operands); +} + +SMLoc +AMDGPUAsmParser::getConstLoc(const OperandVector &Operands) const { + auto Test = [](const AMDGPUOperand& Op) { + return Op.isImmKindConst(); + }; + return getOperandLoc(Test, Operands); +} + //===----------------------------------------------------------------------===// // swizzle //===----------------------------------------------------------------------===// diff --git a/llvm/test/MC/AMDGPU/expressions-gfx10.s b/llvm/test/MC/AMDGPU/expressions-gfx10.s --- a/llvm/test/MC/AMDGPU/expressions-gfx10.s +++ b/llvm/test/MC/AMDGPU/expressions-gfx10.s @@ -64,10 +64,10 @@ // NOGFX10: error: only one literal operand is allowed v_bfe_u32 v0, v2, 123, u -// NOGFX10: error: invalid literal operand +// NOGFX10: error: only one literal operand is allowed v_bfe_u32 v0, v2, u, u -// NOGFX10: error: invalid literal operand +// NOGFX10: error: only one literal operand is allowed v_bfe_u32 v0, v2, u, u1 -// NOGFX10: error: invalid literal operand +// NOGFX10: error: only one literal operand is allowed diff --git a/llvm/test/MC/AMDGPU/expressions-gfx9.s b/llvm/test/MC/AMDGPU/expressions-gfx9.s --- a/llvm/test/MC/AMDGPU/expressions-gfx9.s +++ b/llvm/test/MC/AMDGPU/expressions-gfx9.s @@ -23,7 +23,7 @@ // NOGFX9: error: only one literal operand is allowed v_bfe_u32 v0, v2, v3, u -// NOGFX9: error: invalid literal operand +// NOGFX9: error: literal operands are not supported //===----------------------------------------------------------------------===// // Constant expressions may be used with 'sext' modifier diff --git a/llvm/test/MC/AMDGPU/gfx10_err_pos.s b/llvm/test/MC/AMDGPU/gfx10_err_pos.s --- a/llvm/test/MC/AMDGPU/gfx10_err_pos.s +++ b/llvm/test/MC/AMDGPU/gfx10_err_pos.s @@ -545,24 +545,6 @@ // CHECK-NEXT:{{^}}v_dot_f32_f16 v0, v1, v2 // CHECK-NEXT:{{^}}^ -//============================================================================== -// invalid literal operand - -v_add_f64 v[0:1], 1.23456, -abs(1.2345) -// CHECK: error: invalid literal operand -// CHECK-NEXT:{{^}}v_add_f64 v[0:1], 1.23456, -abs(1.2345) -// CHECK-NEXT:{{^}}^ - -v_min3_i16 v5, 0x5678, 0x5678, 0x5679 -// CHECK: error: invalid literal operand -// CHECK-NEXT:{{^}}v_min3_i16 v5, 0x5678, 0x5678, 0x5679 -// CHECK-NEXT:{{^}}^ - -v_pk_add_f16 v1, 25.0, 25.1 -// CHECK: error: invalid literal operand -// CHECK-NEXT:{{^}}v_pk_add_f16 v1, 25.0, 25.1 -// CHECK-NEXT:{{^}}^ - //============================================================================== // invalid mask @@ -777,7 +759,47 @@ s_and_b32 s2, 0x12345678, 0x12345679 // CHECK: error: only one literal operand is allowed // CHECK-NEXT:{{^}}s_and_b32 s2, 0x12345678, 0x12345679 -// CHECK-NEXT:{{^}}^ +// CHECK-NEXT:{{^}} ^ + +v_add_f64 v[0:1], 1.23456, -abs(1.2345) +// CHECK: error: only one literal operand is allowed +// CHECK-NEXT:{{^}}v_add_f64 v[0:1], 1.23456, -abs(1.2345) +// CHECK-NEXT:{{^}} ^ + +v_min3_i16 v5, 0x5678, 0x5678, 0x5679 +// CHECK: error: only one literal operand is allowed +// CHECK-NEXT:{{^}}v_min3_i16 v5, 0x5678, 0x5678, 0x5679 +// CHECK-NEXT:{{^}} ^ + +v_pk_add_f16 v1, 25.0, 25.1 +// CHECK: error: only one literal operand is allowed +// CHECK-NEXT:{{^}}v_pk_add_f16 v1, 25.0, 25.1 +// CHECK-NEXT:{{^}} ^ + +v_fma_mix_f32 v5, 0x7c, 0x7b, 1 +// CHECK: error: only one literal operand is allowed +// CHECK-NEXT:{{^}}v_fma_mix_f32 v5, 0x7c, 0x7b, 1 +// CHECK-NEXT:{{^}} ^ + +v_pk_add_i16 v5, 0x7c, 0x4000 +// CHECK: error: only one literal operand is allowed +// CHECK-NEXT:{{^}}v_pk_add_i16 v5, 0x7c, 0x4000 +// CHECK-NEXT:{{^}} ^ + +v_pk_add_i16 v5, 0x4400, 0x4000 +// CHECK: error: only one literal operand is allowed +// CHECK-NEXT:{{^}}v_pk_add_i16 v5, 0x4400, 0x4000 +// CHECK-NEXT:{{^}} ^ + +v_bfe_u32 v0, v2, 123, undef +// CHECK: error: only one literal operand is allowed +// CHECK-NEXT:{{^}}v_bfe_u32 v0, v2, 123, undef +// CHECK-NEXT:{{^}} ^ + +v_bfe_u32 v0, v2, undef, 123 +// CHECK: error: only one literal operand is allowed +// CHECK-NEXT:{{^}}v_bfe_u32 v0, v2, undef, 123 +// CHECK-NEXT:{{^}} ^ //============================================================================== // out of bounds attr diff --git a/llvm/test/MC/AMDGPU/gfx7_err_pos.s b/llvm/test/MC/AMDGPU/gfx7_err_pos.s --- a/llvm/test/MC/AMDGPU/gfx7_err_pos.s +++ b/llvm/test/MC/AMDGPU/gfx7_err_pos.s @@ -17,9 +17,9 @@ // CHECK-NEXT:{{^}} ^ //============================================================================== -// invalid literal operand +// literal operands are not supported v_and_b32_e64 v0, 0.159154943091895317852646485335, v1 -// CHECK: error: invalid literal operand +// CHECK: error: literal operands are not supported // CHECK-NEXT:{{^}}v_and_b32_e64 v0, 0.159154943091895317852646485335, v1 -// CHECK-NEXT:{{^}}^ +// CHECK-NEXT:{{^}} ^ diff --git a/llvm/test/MC/AMDGPU/gfx8_asm_all.s b/llvm/test/MC/AMDGPU/gfx8_asm_all.s --- a/llvm/test/MC/AMDGPU/gfx8_asm_all.s +++ b/llvm/test/MC/AMDGPU/gfx8_asm_all.s @@ -31646,10 +31646,10 @@ // CHECK: [0x05,0x00,0x79,0xd1,0xc1,0x00,0x00,0x00] v_cvt_f16_u16_e64 v5, 0.5 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_cvt_f16_u16_e64 v5, -4.0 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_cvt_f16_u16_e64 v5, v1 clamp // CHECK: [0x05,0x80,0x79,0xd1,0x01,0x01,0x00,0x00] @@ -31781,10 +31781,10 @@ // CHECK: [0x05,0x00,0x7a,0xd1,0xc1,0x00,0x00,0x00] v_cvt_f16_i16_e64 v5, 0.5 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_cvt_f16_i16_e64 v5, -4.0 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_cvt_f16_i16_e64 v5, v1 clamp // CHECK: [0x05,0x80,0x7a,0xd1,0x01,0x01,0x00,0x00] @@ -41129,10 +41129,10 @@ // CHECK: [0x05,0x00,0x26,0xd1,0xc1,0x04,0x02,0x00] v_add_u16_e64 v5, 0.5, v2 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_add_u16_e64 v5, -4.0, v2 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_add_u16_e64 v5, v1, v255 // CHECK: [0x05,0x00,0x26,0xd1,0x01,0xff,0x03,0x00] @@ -41186,10 +41186,10 @@ // CHECK: [0x05,0x00,0x26,0xd1,0x01,0x83,0x01,0x00] v_add_u16_e64 v5, v1, 0.5 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_add_u16_e64 v5, v1, -4.0 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_sub_u16 v5, v1, v2 // CHECK: [0x01,0x05,0x0a,0x4e] @@ -41321,10 +41321,10 @@ // CHECK: [0x05,0x00,0x27,0xd1,0xc1,0x04,0x02,0x00] v_sub_u16_e64 v5, 0.5, v2 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_sub_u16_e64 v5, -4.0, v2 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_sub_u16_e64 v5, v1, v255 // CHECK: [0x05,0x00,0x27,0xd1,0x01,0xff,0x03,0x00] @@ -41378,10 +41378,10 @@ // CHECK: [0x05,0x00,0x27,0xd1,0x01,0x83,0x01,0x00] v_sub_u16_e64 v5, v1, 0.5 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_sub_u16_e64 v5, v1, -4.0 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_subrev_u16 v5, v1, v2 // CHECK: [0x01,0x05,0x0a,0x50] @@ -41513,10 +41513,10 @@ // CHECK: [0x05,0x00,0x28,0xd1,0xc1,0x04,0x02,0x00] v_subrev_u16_e64 v5, 0.5, v2 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_subrev_u16_e64 v5, -4.0, v2 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_subrev_u16_e64 v5, v1, v255 // CHECK: [0x05,0x00,0x28,0xd1,0x01,0xff,0x03,0x00] @@ -41570,10 +41570,10 @@ // CHECK: [0x05,0x00,0x28,0xd1,0x01,0x83,0x01,0x00] v_subrev_u16_e64 v5, v1, 0.5 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_subrev_u16_e64 v5, v1, -4.0 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_mul_lo_u16 v5, v1, v2 // CHECK: [0x01,0x05,0x0a,0x52] @@ -41705,10 +41705,10 @@ // CHECK: [0x05,0x00,0x29,0xd1,0xc1,0x04,0x02,0x00] v_mul_lo_u16_e64 v5, 0.5, v2 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_mul_lo_u16_e64 v5, -4.0, v2 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_mul_lo_u16_e64 v5, v1, v255 // CHECK: [0x05,0x00,0x29,0xd1,0x01,0xff,0x03,0x00] @@ -41762,10 +41762,10 @@ // CHECK: [0x05,0x00,0x29,0xd1,0x01,0x83,0x01,0x00] v_mul_lo_u16_e64 v5, v1, 0.5 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_mul_lo_u16_e64 v5, v1, -4.0 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_lshlrev_b16 v5, v1, v2 // CHECK: [0x01,0x05,0x0a,0x54] @@ -41897,10 +41897,10 @@ // CHECK: [0x05,0x00,0x2a,0xd1,0xc1,0x04,0x02,0x00] v_lshlrev_b16_e64 v5, 0.5, v2 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_lshlrev_b16_e64 v5, -4.0, v2 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_lshlrev_b16_e64 v5, v1, v255 // CHECK: [0x05,0x00,0x2a,0xd1,0x01,0xff,0x03,0x00] @@ -41954,10 +41954,10 @@ // CHECK: [0x05,0x00,0x2a,0xd1,0x01,0x83,0x01,0x00] v_lshlrev_b16_e64 v5, v1, 0.5 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_lshlrev_b16_e64 v5, v1, -4.0 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_lshrrev_b16 v5, v1, v2 // CHECK: [0x01,0x05,0x0a,0x56] @@ -42089,10 +42089,10 @@ // CHECK: [0x05,0x00,0x2b,0xd1,0xc1,0x04,0x02,0x00] v_lshrrev_b16_e64 v5, 0.5, v2 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_lshrrev_b16_e64 v5, -4.0, v2 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_lshrrev_b16_e64 v5, v1, v255 // CHECK: [0x05,0x00,0x2b,0xd1,0x01,0xff,0x03,0x00] @@ -42146,10 +42146,10 @@ // CHECK: [0x05,0x00,0x2b,0xd1,0x01,0x83,0x01,0x00] v_lshrrev_b16_e64 v5, v1, 0.5 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_lshrrev_b16_e64 v5, v1, -4.0 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_ashrrev_i16 v5, v1, v2 // CHECK: [0x01,0x05,0x0a,0x58] @@ -42281,10 +42281,10 @@ // CHECK: [0x05,0x00,0x2c,0xd1,0xc1,0x04,0x02,0x00] v_ashrrev_i16_e64 v5, 0.5, v2 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_ashrrev_i16_e64 v5, -4.0, v2 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_ashrrev_i16_e64 v5, v1, v255 // CHECK: [0x05,0x00,0x2c,0xd1,0x01,0xff,0x03,0x00] @@ -42338,10 +42338,10 @@ // CHECK: [0x05,0x00,0x2c,0xd1,0x01,0x83,0x01,0x00] v_ashrrev_i16_e64 v5, v1, 0.5 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_ashrrev_i16_e64 v5, v1, -4.0 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_max_f16 v5, v1, v2 // CHECK: [0x01,0x05,0x0a,0x5a] @@ -42899,10 +42899,10 @@ // CHECK: [0x05,0x00,0x2f,0xd1,0xc1,0x04,0x02,0x00] v_max_u16_e64 v5, 0.5, v2 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_max_u16_e64 v5, -4.0, v2 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_max_u16_e64 v5, v1, v255 // CHECK: [0x05,0x00,0x2f,0xd1,0x01,0xff,0x03,0x00] @@ -42956,10 +42956,10 @@ // CHECK: [0x05,0x00,0x2f,0xd1,0x01,0x83,0x01,0x00] v_max_u16_e64 v5, v1, 0.5 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_max_u16_e64 v5, v1, -4.0 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_max_i16 v5, v1, v2 // CHECK: [0x01,0x05,0x0a,0x60] @@ -43091,10 +43091,10 @@ // CHECK: [0x05,0x00,0x30,0xd1,0xc1,0x04,0x02,0x00] v_max_i16_e64 v5, 0.5, v2 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_max_i16_e64 v5, -4.0, v2 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_max_i16_e64 v5, v1, v255 // CHECK: [0x05,0x00,0x30,0xd1,0x01,0xff,0x03,0x00] @@ -43148,10 +43148,10 @@ // CHECK: [0x05,0x00,0x30,0xd1,0x01,0x83,0x01,0x00] v_max_i16_e64 v5, v1, 0.5 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_max_i16_e64 v5, v1, -4.0 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_min_u16 v5, v1, v2 // CHECK: [0x01,0x05,0x0a,0x62] @@ -43283,10 +43283,10 @@ // CHECK: [0x05,0x00,0x31,0xd1,0xc1,0x04,0x02,0x00] v_min_u16_e64 v5, 0.5, v2 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_min_u16_e64 v5, -4.0, v2 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_min_u16_e64 v5, v1, v255 // CHECK: [0x05,0x00,0x31,0xd1,0x01,0xff,0x03,0x00] @@ -43340,10 +43340,10 @@ // CHECK: [0x05,0x00,0x31,0xd1,0x01,0x83,0x01,0x00] v_min_u16_e64 v5, v1, 0.5 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_min_u16_e64 v5, v1, -4.0 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_min_i16 v5, v1, v2 // CHECK: [0x01,0x05,0x0a,0x64] @@ -43475,10 +43475,10 @@ // CHECK: [0x05,0x00,0x32,0xd1,0xc1,0x04,0x02,0x00] v_min_i16_e64 v5, 0.5, v2 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_min_i16_e64 v5, -4.0, v2 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_min_i16_e64 v5, v1, v255 // CHECK: [0x05,0x00,0x32,0xd1,0x01,0xff,0x03,0x00] @@ -43532,10 +43532,10 @@ // CHECK: [0x05,0x00,0x32,0xd1,0x01,0x83,0x01,0x00] v_min_i16_e64 v5, v1, 0.5 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_min_i16_e64 v5, v1, -4.0 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_ldexp_f16 v5, v1, v2 // CHECK: [0x01,0x05,0x0a,0x66] @@ -51509,10 +51509,10 @@ // CHECK: [0x05,0x00,0xeb,0xd1,0xc1,0x04,0x0e,0x04] v_mad_u16 v5, 0.5, v2, v3 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_mad_u16 v5, -4.0, v2, v3 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_mad_u16 v5, v1, v255, v3 // CHECK: [0x05,0x00,0xeb,0xd1,0x01,0xff,0x0f,0x04] @@ -51566,10 +51566,10 @@ // CHECK: [0x05,0x00,0xeb,0xd1,0x01,0x83,0x0d,0x04] v_mad_u16 v5, v1, 0.5, v3 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_mad_u16 v5, v1, -4.0, v3 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_mad_u16 v5, v1, v2, v255 // CHECK: [0x05,0x00,0xeb,0xd1,0x01,0x05,0xfe,0x07] @@ -51623,10 +51623,10 @@ // CHECK: [0x05,0x00,0xeb,0xd1,0x01,0x05,0x06,0x03] v_mad_u16 v5, v1, v2, 0.5 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_mad_u16 v5, v1, v2, -4.0 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_mad_u16 v5, v1, v2, v3 clamp // CHECK: [0x05,0x80,0xeb,0xd1,0x01,0x05,0x0e,0x04] @@ -51689,10 +51689,10 @@ // CHECK: [0x05,0x00,0xec,0xd1,0xc1,0x04,0x0e,0x04] v_mad_i16 v5, 0.5, v2, v3 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_mad_i16 v5, -4.0, v2, v3 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_mad_i16 v5, v1, v255, v3 // CHECK: [0x05,0x00,0xec,0xd1,0x01,0xff,0x0f,0x04] @@ -51746,10 +51746,10 @@ // CHECK: [0x05,0x00,0xec,0xd1,0x01,0x83,0x0d,0x04] v_mad_i16 v5, v1, 0.5, v3 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_mad_i16 v5, v1, -4.0, v3 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_mad_i16 v5, v1, v2, v255 // CHECK: [0x05,0x00,0xec,0xd1,0x01,0x05,0xfe,0x07] @@ -51803,10 +51803,10 @@ // CHECK: [0x05,0x00,0xec,0xd1,0x01,0x05,0x06,0x03] v_mad_i16 v5, v1, v2, 0.5 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_mad_i16 v5, v1, v2, -4.0 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_mad_i16 v5, v1, v2, v3 clamp // CHECK: [0x05,0x80,0xec,0xd1,0x01,0x05,0x0e,0x04] @@ -76394,10 +76394,10 @@ // CHECK: [0x0a,0x00,0xa0,0xd0,0xc1,0x04,0x02,0x00] v_cmp_f_i16_e64 s[10:11], 0.5, v2 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_cmp_f_i16_e64 s[10:11], -4.0, v2 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_cmp_f_i16_e64 s[10:11], v1, v255 // CHECK: [0x0a,0x00,0xa0,0xd0,0x01,0xff,0x03,0x00] @@ -76451,10 +76451,10 @@ // CHECK: [0x0a,0x00,0xa0,0xd0,0x01,0x83,0x01,0x00] v_cmp_f_i16_e64 s[10:11], v1, 0.5 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_cmp_f_i16_e64 s[10:11], v1, -4.0 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_cmp_lt_i16 vcc, v1, v2 // CHECK: [0x01,0x05,0x42,0x7d] @@ -76601,10 +76601,10 @@ // CHECK: [0x0a,0x00,0xa1,0xd0,0xc1,0x04,0x02,0x00] v_cmp_lt_i16_e64 s[10:11], 0.5, v2 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_cmp_lt_i16_e64 s[10:11], -4.0, v2 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_cmp_lt_i16_e64 s[10:11], v1, v255 // CHECK: [0x0a,0x00,0xa1,0xd0,0x01,0xff,0x03,0x00] @@ -76658,10 +76658,10 @@ // CHECK: [0x0a,0x00,0xa1,0xd0,0x01,0x83,0x01,0x00] v_cmp_lt_i16_e64 s[10:11], v1, 0.5 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_cmp_lt_i16_e64 s[10:11], v1, -4.0 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_cmp_eq_i16 vcc, v1, v2 // CHECK: [0x01,0x05,0x44,0x7d] @@ -76808,10 +76808,10 @@ // CHECK: [0x0a,0x00,0xa2,0xd0,0xc1,0x04,0x02,0x00] v_cmp_eq_i16_e64 s[10:11], 0.5, v2 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_cmp_eq_i16_e64 s[10:11], -4.0, v2 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_cmp_eq_i16_e64 s[10:11], v1, v255 // CHECK: [0x0a,0x00,0xa2,0xd0,0x01,0xff,0x03,0x00] @@ -76865,10 +76865,10 @@ // CHECK: [0x0a,0x00,0xa2,0xd0,0x01,0x83,0x01,0x00] v_cmp_eq_i16_e64 s[10:11], v1, 0.5 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_cmp_eq_i16_e64 s[10:11], v1, -4.0 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_cmp_le_i16 vcc, v1, v2 // CHECK: [0x01,0x05,0x46,0x7d] @@ -77015,10 +77015,10 @@ // CHECK: [0x0a,0x00,0xa3,0xd0,0xc1,0x04,0x02,0x00] v_cmp_le_i16_e64 s[10:11], 0.5, v2 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_cmp_le_i16_e64 s[10:11], -4.0, v2 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_cmp_le_i16_e64 s[10:11], v1, v255 // CHECK: [0x0a,0x00,0xa3,0xd0,0x01,0xff,0x03,0x00] @@ -77072,10 +77072,10 @@ // CHECK: [0x0a,0x00,0xa3,0xd0,0x01,0x83,0x01,0x00] v_cmp_le_i16_e64 s[10:11], v1, 0.5 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_cmp_le_i16_e64 s[10:11], v1, -4.0 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_cmp_gt_i16 vcc, v1, v2 // CHECK: [0x01,0x05,0x48,0x7d] @@ -77222,10 +77222,10 @@ // CHECK: [0x0a,0x00,0xa4,0xd0,0xc1,0x04,0x02,0x00] v_cmp_gt_i16_e64 s[10:11], 0.5, v2 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_cmp_gt_i16_e64 s[10:11], -4.0, v2 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_cmp_gt_i16_e64 s[10:11], v1, v255 // CHECK: [0x0a,0x00,0xa4,0xd0,0x01,0xff,0x03,0x00] @@ -77279,10 +77279,10 @@ // CHECK: [0x0a,0x00,0xa4,0xd0,0x01,0x83,0x01,0x00] v_cmp_gt_i16_e64 s[10:11], v1, 0.5 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_cmp_gt_i16_e64 s[10:11], v1, -4.0 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_cmp_ne_i16 vcc, v1, v2 // CHECK: [0x01,0x05,0x4a,0x7d] @@ -77429,10 +77429,10 @@ // CHECK: [0x0a,0x00,0xa5,0xd0,0xc1,0x04,0x02,0x00] v_cmp_ne_i16_e64 s[10:11], 0.5, v2 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_cmp_ne_i16_e64 s[10:11], -4.0, v2 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_cmp_ne_i16_e64 s[10:11], v1, v255 // CHECK: [0x0a,0x00,0xa5,0xd0,0x01,0xff,0x03,0x00] @@ -77486,10 +77486,10 @@ // CHECK: [0x0a,0x00,0xa5,0xd0,0x01,0x83,0x01,0x00] v_cmp_ne_i16_e64 s[10:11], v1, 0.5 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_cmp_ne_i16_e64 s[10:11], v1, -4.0 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_cmp_ge_i16 vcc, v1, v2 // CHECK: [0x01,0x05,0x4c,0x7d] @@ -77636,10 +77636,10 @@ // CHECK: [0x0a,0x00,0xa6,0xd0,0xc1,0x04,0x02,0x00] v_cmp_ge_i16_e64 s[10:11], 0.5, v2 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_cmp_ge_i16_e64 s[10:11], -4.0, v2 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_cmp_ge_i16_e64 s[10:11], v1, v255 // CHECK: [0x0a,0x00,0xa6,0xd0,0x01,0xff,0x03,0x00] @@ -77693,10 +77693,10 @@ // CHECK: [0x0a,0x00,0xa6,0xd0,0x01,0x83,0x01,0x00] v_cmp_ge_i16_e64 s[10:11], v1, 0.5 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_cmp_ge_i16_e64 s[10:11], v1, -4.0 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_cmp_t_i16 vcc, v1, v2 // CHECK: [0x01,0x05,0x4e,0x7d] @@ -77843,10 +77843,10 @@ // CHECK: [0x0a,0x00,0xa7,0xd0,0xc1,0x04,0x02,0x00] v_cmp_t_i16_e64 s[10:11], 0.5, v2 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_cmp_t_i16_e64 s[10:11], -4.0, v2 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_cmp_t_i16_e64 s[10:11], v1, v255 // CHECK: [0x0a,0x00,0xa7,0xd0,0x01,0xff,0x03,0x00] @@ -77900,10 +77900,10 @@ // CHECK: [0x0a,0x00,0xa7,0xd0,0x01,0x83,0x01,0x00] v_cmp_t_i16_e64 s[10:11], v1, 0.5 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_cmp_t_i16_e64 s[10:11], v1, -4.0 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_cmp_f_u16 vcc, v1, v2 // CHECK: [0x01,0x05,0x50,0x7d] @@ -78050,10 +78050,10 @@ // CHECK: [0x0a,0x00,0xa8,0xd0,0xc1,0x04,0x02,0x00] v_cmp_f_u16_e64 s[10:11], 0.5, v2 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_cmp_f_u16_e64 s[10:11], -4.0, v2 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_cmp_f_u16_e64 s[10:11], v1, v255 // CHECK: [0x0a,0x00,0xa8,0xd0,0x01,0xff,0x03,0x00] @@ -78107,10 +78107,10 @@ // CHECK: [0x0a,0x00,0xa8,0xd0,0x01,0x83,0x01,0x00] v_cmp_f_u16_e64 s[10:11], v1, 0.5 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_cmp_f_u16_e64 s[10:11], v1, -4.0 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_cmp_lt_u16 vcc, v1, v2 // CHECK: [0x01,0x05,0x52,0x7d] @@ -78257,10 +78257,10 @@ // CHECK: [0x0a,0x00,0xa9,0xd0,0xc1,0x04,0x02,0x00] v_cmp_lt_u16_e64 s[10:11], 0.5, v2 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_cmp_lt_u16_e64 s[10:11], -4.0, v2 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_cmp_lt_u16_e64 s[10:11], v1, v255 // CHECK: [0x0a,0x00,0xa9,0xd0,0x01,0xff,0x03,0x00] @@ -78314,10 +78314,10 @@ // CHECK: [0x0a,0x00,0xa9,0xd0,0x01,0x83,0x01,0x00] v_cmp_lt_u16_e64 s[10:11], v1, 0.5 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_cmp_lt_u16_e64 s[10:11], v1, -4.0 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_cmp_eq_u16 vcc, v1, v2 // CHECK: [0x01,0x05,0x54,0x7d] @@ -78464,10 +78464,10 @@ // CHECK: [0x0a,0x00,0xaa,0xd0,0xc1,0x04,0x02,0x00] v_cmp_eq_u16_e64 s[10:11], 0.5, v2 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_cmp_eq_u16_e64 s[10:11], -4.0, v2 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_cmp_eq_u16_e64 s[10:11], v1, v255 // CHECK: [0x0a,0x00,0xaa,0xd0,0x01,0xff,0x03,0x00] @@ -78521,10 +78521,10 @@ // CHECK: [0x0a,0x00,0xaa,0xd0,0x01,0x83,0x01,0x00] v_cmp_eq_u16_e64 s[10:11], v1, 0.5 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_cmp_eq_u16_e64 s[10:11], v1, -4.0 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_cmp_le_u16 vcc, v1, v2 // CHECK: [0x01,0x05,0x56,0x7d] @@ -78671,10 +78671,10 @@ // CHECK: [0x0a,0x00,0xab,0xd0,0xc1,0x04,0x02,0x00] v_cmp_le_u16_e64 s[10:11], 0.5, v2 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_cmp_le_u16_e64 s[10:11], -4.0, v2 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_cmp_le_u16_e64 s[10:11], v1, v255 // CHECK: [0x0a,0x00,0xab,0xd0,0x01,0xff,0x03,0x00] @@ -78728,10 +78728,10 @@ // CHECK: [0x0a,0x00,0xab,0xd0,0x01,0x83,0x01,0x00] v_cmp_le_u16_e64 s[10:11], v1, 0.5 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_cmp_le_u16_e64 s[10:11], v1, -4.0 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_cmp_gt_u16 vcc, v1, v2 // CHECK: [0x01,0x05,0x58,0x7d] @@ -78878,10 +78878,10 @@ // CHECK: [0x0a,0x00,0xac,0xd0,0xc1,0x04,0x02,0x00] v_cmp_gt_u16_e64 s[10:11], 0.5, v2 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_cmp_gt_u16_e64 s[10:11], -4.0, v2 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_cmp_gt_u16_e64 s[10:11], v1, v255 // CHECK: [0x0a,0x00,0xac,0xd0,0x01,0xff,0x03,0x00] @@ -78935,10 +78935,10 @@ // CHECK: [0x0a,0x00,0xac,0xd0,0x01,0x83,0x01,0x00] v_cmp_gt_u16_e64 s[10:11], v1, 0.5 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_cmp_gt_u16_e64 s[10:11], v1, -4.0 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_cmp_ne_u16 vcc, v1, v2 // CHECK: [0x01,0x05,0x5a,0x7d] @@ -79085,10 +79085,10 @@ // CHECK: [0x0a,0x00,0xad,0xd0,0xc1,0x04,0x02,0x00] v_cmp_ne_u16_e64 s[10:11], 0.5, v2 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_cmp_ne_u16_e64 s[10:11], -4.0, v2 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_cmp_ne_u16_e64 s[10:11], v1, v255 // CHECK: [0x0a,0x00,0xad,0xd0,0x01,0xff,0x03,0x00] @@ -79142,10 +79142,10 @@ // CHECK: [0x0a,0x00,0xad,0xd0,0x01,0x83,0x01,0x00] v_cmp_ne_u16_e64 s[10:11], v1, 0.5 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_cmp_ne_u16_e64 s[10:11], v1, -4.0 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_cmp_ge_u16 vcc, v1, v2 // CHECK: [0x01,0x05,0x5c,0x7d] @@ -79292,10 +79292,10 @@ // CHECK: [0x0a,0x00,0xae,0xd0,0xc1,0x04,0x02,0x00] v_cmp_ge_u16_e64 s[10:11], 0.5, v2 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_cmp_ge_u16_e64 s[10:11], -4.0, v2 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_cmp_ge_u16_e64 s[10:11], v1, v255 // CHECK: [0x0a,0x00,0xae,0xd0,0x01,0xff,0x03,0x00] @@ -79349,10 +79349,10 @@ // CHECK: [0x0a,0x00,0xae,0xd0,0x01,0x83,0x01,0x00] v_cmp_ge_u16_e64 s[10:11], v1, 0.5 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_cmp_ge_u16_e64 s[10:11], v1, -4.0 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_cmp_t_u16 vcc, v1, v2 // CHECK: [0x01,0x05,0x5e,0x7d] @@ -79499,10 +79499,10 @@ // CHECK: [0x0a,0x00,0xaf,0xd0,0xc1,0x04,0x02,0x00] v_cmp_t_u16_e64 s[10:11], 0.5, v2 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_cmp_t_u16_e64 s[10:11], -4.0, v2 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_cmp_t_u16_e64 s[10:11], v1, v255 // CHECK: [0x0a,0x00,0xaf,0xd0,0x01,0xff,0x03,0x00] @@ -79556,10 +79556,10 @@ // CHECK: [0x0a,0x00,0xaf,0xd0,0x01,0x83,0x01,0x00] v_cmp_t_u16_e64 s[10:11], v1, 0.5 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_cmp_t_u16_e64 s[10:11], v1, -4.0 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_cmpx_f_i16 vcc, v1, v2 // CHECK: [0x01,0x05,0x60,0x7d] @@ -79706,10 +79706,10 @@ // CHECK: [0x0a,0x00,0xb0,0xd0,0xc1,0x04,0x02,0x00] v_cmpx_f_i16_e64 s[10:11], 0.5, v2 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_cmpx_f_i16_e64 s[10:11], -4.0, v2 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_cmpx_f_i16_e64 s[10:11], v1, v255 // CHECK: [0x0a,0x00,0xb0,0xd0,0x01,0xff,0x03,0x00] @@ -79763,10 +79763,10 @@ // CHECK: [0x0a,0x00,0xb0,0xd0,0x01,0x83,0x01,0x00] v_cmpx_f_i16_e64 s[10:11], v1, 0.5 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_cmpx_f_i16_e64 s[10:11], v1, -4.0 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_cmpx_lt_i16 vcc, v1, v2 // CHECK: [0x01,0x05,0x62,0x7d] @@ -79913,10 +79913,10 @@ // CHECK: [0x0a,0x00,0xb1,0xd0,0xc1,0x04,0x02,0x00] v_cmpx_lt_i16_e64 s[10:11], 0.5, v2 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_cmpx_lt_i16_e64 s[10:11], -4.0, v2 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_cmpx_lt_i16_e64 s[10:11], v1, v255 // CHECK: [0x0a,0x00,0xb1,0xd0,0x01,0xff,0x03,0x00] @@ -79970,10 +79970,10 @@ // CHECK: [0x0a,0x00,0xb1,0xd0,0x01,0x83,0x01,0x00] v_cmpx_lt_i16_e64 s[10:11], v1, 0.5 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_cmpx_lt_i16_e64 s[10:11], v1, -4.0 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_cmpx_eq_i16 vcc, v1, v2 // CHECK: [0x01,0x05,0x64,0x7d] @@ -80120,10 +80120,10 @@ // CHECK: [0x0a,0x00,0xb2,0xd0,0xc1,0x04,0x02,0x00] v_cmpx_eq_i16_e64 s[10:11], 0.5, v2 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_cmpx_eq_i16_e64 s[10:11], -4.0, v2 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_cmpx_eq_i16_e64 s[10:11], v1, v255 // CHECK: [0x0a,0x00,0xb2,0xd0,0x01,0xff,0x03,0x00] @@ -80177,10 +80177,10 @@ // CHECK: [0x0a,0x00,0xb2,0xd0,0x01,0x83,0x01,0x00] v_cmpx_eq_i16_e64 s[10:11], v1, 0.5 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_cmpx_eq_i16_e64 s[10:11], v1, -4.0 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_cmpx_le_i16 vcc, v1, v2 // CHECK: [0x01,0x05,0x66,0x7d] @@ -80327,10 +80327,10 @@ // CHECK: [0x0a,0x00,0xb3,0xd0,0xc1,0x04,0x02,0x00] v_cmpx_le_i16_e64 s[10:11], 0.5, v2 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_cmpx_le_i16_e64 s[10:11], -4.0, v2 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_cmpx_le_i16_e64 s[10:11], v1, v255 // CHECK: [0x0a,0x00,0xb3,0xd0,0x01,0xff,0x03,0x00] @@ -80384,10 +80384,10 @@ // CHECK: [0x0a,0x00,0xb3,0xd0,0x01,0x83,0x01,0x00] v_cmpx_le_i16_e64 s[10:11], v1, 0.5 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_cmpx_le_i16_e64 s[10:11], v1, -4.0 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_cmpx_gt_i16 vcc, v1, v2 // CHECK: [0x01,0x05,0x68,0x7d] @@ -80534,10 +80534,10 @@ // CHECK: [0x0a,0x00,0xb4,0xd0,0xc1,0x04,0x02,0x00] v_cmpx_gt_i16_e64 s[10:11], 0.5, v2 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_cmpx_gt_i16_e64 s[10:11], -4.0, v2 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_cmpx_gt_i16_e64 s[10:11], v1, v255 // CHECK: [0x0a,0x00,0xb4,0xd0,0x01,0xff,0x03,0x00] @@ -80591,10 +80591,10 @@ // CHECK: [0x0a,0x00,0xb4,0xd0,0x01,0x83,0x01,0x00] v_cmpx_gt_i16_e64 s[10:11], v1, 0.5 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_cmpx_gt_i16_e64 s[10:11], v1, -4.0 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_cmpx_ne_i16 vcc, v1, v2 // CHECK: [0x01,0x05,0x6a,0x7d] @@ -80741,10 +80741,10 @@ // CHECK: [0x0a,0x00,0xb5,0xd0,0xc1,0x04,0x02,0x00] v_cmpx_ne_i16_e64 s[10:11], 0.5, v2 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_cmpx_ne_i16_e64 s[10:11], -4.0, v2 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_cmpx_ne_i16_e64 s[10:11], v1, v255 // CHECK: [0x0a,0x00,0xb5,0xd0,0x01,0xff,0x03,0x00] @@ -80798,10 +80798,10 @@ // CHECK: [0x0a,0x00,0xb5,0xd0,0x01,0x83,0x01,0x00] v_cmpx_ne_i16_e64 s[10:11], v1, 0.5 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_cmpx_ne_i16_e64 s[10:11], v1, -4.0 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_cmpx_ge_i16 vcc, v1, v2 // CHECK: [0x01,0x05,0x6c,0x7d] @@ -80948,10 +80948,10 @@ // CHECK: [0x0a,0x00,0xb6,0xd0,0xc1,0x04,0x02,0x00] v_cmpx_ge_i16_e64 s[10:11], 0.5, v2 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_cmpx_ge_i16_e64 s[10:11], -4.0, v2 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_cmpx_ge_i16_e64 s[10:11], v1, v255 // CHECK: [0x0a,0x00,0xb6,0xd0,0x01,0xff,0x03,0x00] @@ -81005,10 +81005,10 @@ // CHECK: [0x0a,0x00,0xb6,0xd0,0x01,0x83,0x01,0x00] v_cmpx_ge_i16_e64 s[10:11], v1, 0.5 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_cmpx_ge_i16_e64 s[10:11], v1, -4.0 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_cmpx_t_i16 vcc, v1, v2 // CHECK: [0x01,0x05,0x6e,0x7d] @@ -81155,10 +81155,10 @@ // CHECK: [0x0a,0x00,0xb7,0xd0,0xc1,0x04,0x02,0x00] v_cmpx_t_i16_e64 s[10:11], 0.5, v2 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_cmpx_t_i16_e64 s[10:11], -4.0, v2 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_cmpx_t_i16_e64 s[10:11], v1, v255 // CHECK: [0x0a,0x00,0xb7,0xd0,0x01,0xff,0x03,0x00] @@ -81212,10 +81212,10 @@ // CHECK: [0x0a,0x00,0xb7,0xd0,0x01,0x83,0x01,0x00] v_cmpx_t_i16_e64 s[10:11], v1, 0.5 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_cmpx_t_i16_e64 s[10:11], v1, -4.0 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_cmpx_f_u16 vcc, v1, v2 // CHECK: [0x01,0x05,0x70,0x7d] @@ -81362,10 +81362,10 @@ // CHECK: [0x0a,0x00,0xb8,0xd0,0xc1,0x04,0x02,0x00] v_cmpx_f_u16_e64 s[10:11], 0.5, v2 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_cmpx_f_u16_e64 s[10:11], -4.0, v2 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_cmpx_f_u16_e64 s[10:11], v1, v255 // CHECK: [0x0a,0x00,0xb8,0xd0,0x01,0xff,0x03,0x00] @@ -81419,10 +81419,10 @@ // CHECK: [0x0a,0x00,0xb8,0xd0,0x01,0x83,0x01,0x00] v_cmpx_f_u16_e64 s[10:11], v1, 0.5 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_cmpx_f_u16_e64 s[10:11], v1, -4.0 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_cmpx_lt_u16 vcc, v1, v2 // CHECK: [0x01,0x05,0x72,0x7d] @@ -81569,10 +81569,10 @@ // CHECK: [0x0a,0x00,0xb9,0xd0,0xc1,0x04,0x02,0x00] v_cmpx_lt_u16_e64 s[10:11], 0.5, v2 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_cmpx_lt_u16_e64 s[10:11], -4.0, v2 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_cmpx_lt_u16_e64 s[10:11], v1, v255 // CHECK: [0x0a,0x00,0xb9,0xd0,0x01,0xff,0x03,0x00] @@ -81626,10 +81626,10 @@ // CHECK: [0x0a,0x00,0xb9,0xd0,0x01,0x83,0x01,0x00] v_cmpx_lt_u16_e64 s[10:11], v1, 0.5 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_cmpx_lt_u16_e64 s[10:11], v1, -4.0 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_cmpx_eq_u16 vcc, v1, v2 // CHECK: [0x01,0x05,0x74,0x7d] @@ -81776,10 +81776,10 @@ // CHECK: [0x0a,0x00,0xba,0xd0,0xc1,0x04,0x02,0x00] v_cmpx_eq_u16_e64 s[10:11], 0.5, v2 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_cmpx_eq_u16_e64 s[10:11], -4.0, v2 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_cmpx_eq_u16_e64 s[10:11], v1, v255 // CHECK: [0x0a,0x00,0xba,0xd0,0x01,0xff,0x03,0x00] @@ -81833,10 +81833,10 @@ // CHECK: [0x0a,0x00,0xba,0xd0,0x01,0x83,0x01,0x00] v_cmpx_eq_u16_e64 s[10:11], v1, 0.5 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_cmpx_eq_u16_e64 s[10:11], v1, -4.0 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_cmpx_le_u16 vcc, v1, v2 // CHECK: [0x01,0x05,0x76,0x7d] @@ -81983,10 +81983,10 @@ // CHECK: [0x0a,0x00,0xbb,0xd0,0xc1,0x04,0x02,0x00] v_cmpx_le_u16_e64 s[10:11], 0.5, v2 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_cmpx_le_u16_e64 s[10:11], -4.0, v2 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_cmpx_le_u16_e64 s[10:11], v1, v255 // CHECK: [0x0a,0x00,0xbb,0xd0,0x01,0xff,0x03,0x00] @@ -82040,10 +82040,10 @@ // CHECK: [0x0a,0x00,0xbb,0xd0,0x01,0x83,0x01,0x00] v_cmpx_le_u16_e64 s[10:11], v1, 0.5 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_cmpx_le_u16_e64 s[10:11], v1, -4.0 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_cmpx_gt_u16 vcc, v1, v2 // CHECK: [0x01,0x05,0x78,0x7d] @@ -82190,10 +82190,10 @@ // CHECK: [0x0a,0x00,0xbc,0xd0,0xc1,0x04,0x02,0x00] v_cmpx_gt_u16_e64 s[10:11], 0.5, v2 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_cmpx_gt_u16_e64 s[10:11], -4.0, v2 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_cmpx_gt_u16_e64 s[10:11], v1, v255 // CHECK: [0x0a,0x00,0xbc,0xd0,0x01,0xff,0x03,0x00] @@ -82247,10 +82247,10 @@ // CHECK: [0x0a,0x00,0xbc,0xd0,0x01,0x83,0x01,0x00] v_cmpx_gt_u16_e64 s[10:11], v1, 0.5 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_cmpx_gt_u16_e64 s[10:11], v1, -4.0 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_cmpx_ne_u16 vcc, v1, v2 // CHECK: [0x01,0x05,0x7a,0x7d] @@ -82397,10 +82397,10 @@ // CHECK: [0x0a,0x00,0xbd,0xd0,0xc1,0x04,0x02,0x00] v_cmpx_ne_u16_e64 s[10:11], 0.5, v2 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_cmpx_ne_u16_e64 s[10:11], -4.0, v2 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_cmpx_ne_u16_e64 s[10:11], v1, v255 // CHECK: [0x0a,0x00,0xbd,0xd0,0x01,0xff,0x03,0x00] @@ -82454,10 +82454,10 @@ // CHECK: [0x0a,0x00,0xbd,0xd0,0x01,0x83,0x01,0x00] v_cmpx_ne_u16_e64 s[10:11], v1, 0.5 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_cmpx_ne_u16_e64 s[10:11], v1, -4.0 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_cmpx_ge_u16 vcc, v1, v2 // CHECK: [0x01,0x05,0x7c,0x7d] @@ -82604,10 +82604,10 @@ // CHECK: [0x0a,0x00,0xbe,0xd0,0xc1,0x04,0x02,0x00] v_cmpx_ge_u16_e64 s[10:11], 0.5, v2 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_cmpx_ge_u16_e64 s[10:11], -4.0, v2 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_cmpx_ge_u16_e64 s[10:11], v1, v255 // CHECK: [0x0a,0x00,0xbe,0xd0,0x01,0xff,0x03,0x00] @@ -82661,10 +82661,10 @@ // CHECK: [0x0a,0x00,0xbe,0xd0,0x01,0x83,0x01,0x00] v_cmpx_ge_u16_e64 s[10:11], v1, 0.5 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_cmpx_ge_u16_e64 s[10:11], v1, -4.0 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_cmpx_t_u16 vcc, v1, v2 // CHECK: [0x01,0x05,0x7e,0x7d] @@ -82811,10 +82811,10 @@ // CHECK: [0x0a,0x00,0xbf,0xd0,0xc1,0x04,0x02,0x00] v_cmpx_t_u16_e64 s[10:11], 0.5, v2 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_cmpx_t_u16_e64 s[10:11], -4.0, v2 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_cmpx_t_u16_e64 s[10:11], v1, v255 // CHECK: [0x0a,0x00,0xbf,0xd0,0x01,0xff,0x03,0x00] @@ -82868,10 +82868,10 @@ // CHECK: [0x0a,0x00,0xbf,0xd0,0x01,0x83,0x01,0x00] v_cmpx_t_u16_e64 s[10:11], v1, 0.5 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_cmpx_t_u16_e64 s[10:11], v1, -4.0 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_cmp_f_i32 vcc, v1, v2 // CHECK: [0x01,0x05,0x80,0x7d] diff --git a/llvm/test/MC/AMDGPU/gfx9-asm-err.s b/llvm/test/MC/AMDGPU/gfx9-asm-err.s --- a/llvm/test/MC/AMDGPU/gfx9-asm-err.s +++ b/llvm/test/MC/AMDGPU/gfx9-asm-err.s @@ -1,31 +1,31 @@ // RUN: not llvm-mc -arch=amdgcn -mcpu=gfx900 %s 2>&1 | FileCheck -check-prefix=GFX9ERR --implicit-check-not=error: %s v_cvt_f16_u16_e64 v5, 0.5 -// GFX9ERR: error: invalid literal operand +// GFX9ERR: error: literal operands are not supported v_cvt_f16_u16_e64 v5, -4.0 -// GFX9ERR: error: invalid literal operand +// GFX9ERR: error: literal operands are not supported v_add_u16_e64 v5, v1, 0.5 -// GFX9ERR: error: invalid literal operand +// GFX9ERR: error: literal operands are not supported v_add_u16_e64 v5, v1, -4.0 -// GFX9ERR: error: invalid literal operand +// GFX9ERR: error: literal operands are not supported v_cvt_f16_i16_e64 v5, 0.5 -// GFX9ERR: error: invalid literal operand +// GFX9ERR: error: literal operands are not supported v_cvt_f16_i16_e64 v5, -4.0 -// GFX9ERR: error: invalid literal operand +// GFX9ERR: error: literal operands are not supported v_add_u16_e64 v5, 0.5, v2 -// GFX9ERR: error: invalid literal operand +// GFX9ERR: error: literal operands are not supported v_add_u16_e64 v5, -4.0, v2 -// GFX9ERR: error: invalid literal operand +// GFX9ERR: error: literal operands are not supported v_subrev_u16_e64 v5, v1, 0.5 -// GFX9ERR: error: invalid literal operand +// GFX9ERR: error: literal operands are not supported v_subrev_u16_e64 v5, v1, -4.0 -// GFX9ERR: error: invalid literal operand +// GFX9ERR: error: literal operands are not supported diff --git a/llvm/test/MC/AMDGPU/gfx908_err_pos.s b/llvm/test/MC/AMDGPU/gfx908_err_pos.s --- a/llvm/test/MC/AMDGPU/gfx908_err_pos.s +++ b/llvm/test/MC/AMDGPU/gfx908_err_pos.s @@ -1,5 +1,13 @@ // RUN: not llvm-mc -arch=amdgcn -mcpu=gfx908 %s 2>&1 | FileCheck %s --implicit-check-not=error: --strict-whitespace +//============================================================================== +// inline constants are not allowed for this operand + +v_mfma_f32_32x32x1f32 a[0:31], v0, v1, 0 +// CHECK: error: inline constants are not allowed for this operand +// CHECK-NEXT:{{^}}v_mfma_f32_32x32x1f32 a[0:31], v0, v1, 0 +// CHECK-NEXT:{{^}} ^ + //============================================================================== // invalid neg_hi value diff --git a/llvm/test/MC/AMDGPU/gfx9_asm_all.s b/llvm/test/MC/AMDGPU/gfx9_asm_all.s --- a/llvm/test/MC/AMDGPU/gfx9_asm_all.s +++ b/llvm/test/MC/AMDGPU/gfx9_asm_all.s @@ -34899,10 +34899,10 @@ // CHECK: [0x05,0x00,0x27,0xd1,0xc1,0x04,0x02,0x00] v_sub_u16_e64 v5, 0.5, v2 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_sub_u16_e64 v5, -4.0, v2 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_sub_u16_e64 v5, v1, v255 // CHECK: [0x05,0x00,0x27,0xd1,0x01,0xff,0x03,0x00] @@ -34941,10 +34941,10 @@ // CHECK: [0x05,0x00,0x27,0xd1,0x01,0x83,0x01,0x00] v_sub_u16_e64 v5, v1, 0.5 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_sub_u16_e64 v5, v1, -4.0 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_subrev_u16 v5, v1, v2 // CHECK: [0x01,0x05,0x0a,0x50] @@ -35046,10 +35046,10 @@ // CHECK: [0x05,0x00,0x28,0xd1,0xc1,0x04,0x02,0x00] v_subrev_u16_e64 v5, 0.5, v2 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_subrev_u16_e64 v5, -4.0, v2 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_subrev_u16_e64 v5, v1, v255 // CHECK: [0x05,0x00,0x28,0xd1,0x01,0xff,0x03,0x00] @@ -35187,10 +35187,10 @@ // CHECK: [0x05,0x00,0x29,0xd1,0xc1,0x04,0x02,0x00] v_mul_lo_u16_e64 v5, 0.5, v2 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_mul_lo_u16_e64 v5, -4.0, v2 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_mul_lo_u16_e64 v5, v1, v255 // CHECK: [0x05,0x00,0x29,0xd1,0x01,0xff,0x03,0x00] @@ -35229,10 +35229,10 @@ // CHECK: [0x05,0x00,0x29,0xd1,0x01,0x83,0x01,0x00] v_mul_lo_u16_e64 v5, v1, 0.5 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_mul_lo_u16_e64 v5, v1, -4.0 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_lshlrev_b16 v5, v1, v2 // CHECK: [0x01,0x05,0x0a,0x54] @@ -35334,10 +35334,10 @@ // CHECK: [0x05,0x00,0x2a,0xd1,0xc1,0x04,0x02,0x00] v_lshlrev_b16_e64 v5, 0.5, v2 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_lshlrev_b16_e64 v5, -4.0, v2 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_lshlrev_b16_e64 v5, v1, v255 // CHECK: [0x05,0x00,0x2a,0xd1,0x01,0xff,0x03,0x00] @@ -35376,10 +35376,10 @@ // CHECK: [0x05,0x00,0x2a,0xd1,0x01,0x83,0x01,0x00] v_lshlrev_b16_e64 v5, v1, 0.5 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_lshlrev_b16_e64 v5, v1, -4.0 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_lshrrev_b16 v5, v1, v2 // CHECK: [0x01,0x05,0x0a,0x56] @@ -35481,10 +35481,10 @@ // CHECK: [0x05,0x00,0x2b,0xd1,0xc1,0x04,0x02,0x00] v_lshrrev_b16_e64 v5, 0.5, v2 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_lshrrev_b16_e64 v5, -4.0, v2 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_lshrrev_b16_e64 v5, v1, v255 // CHECK: [0x05,0x00,0x2b,0xd1,0x01,0xff,0x03,0x00] @@ -35523,10 +35523,10 @@ // CHECK: [0x05,0x00,0x2b,0xd1,0x01,0x83,0x01,0x00] v_lshrrev_b16_e64 v5, v1, 0.5 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_lshrrev_b16_e64 v5, v1, -4.0 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_ashrrev_i16 v5, v1, v2 // CHECK: [0x01,0x05,0x0a,0x58] @@ -35628,10 +35628,10 @@ // CHECK: [0x05,0x00,0x2c,0xd1,0xc1,0x04,0x02,0x00] v_ashrrev_i16_e64 v5, 0.5, v2 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_ashrrev_i16_e64 v5, -4.0, v2 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_ashrrev_i16_e64 v5, v1, v255 // CHECK: [0x05,0x00,0x2c,0xd1,0x01,0xff,0x03,0x00] @@ -35670,10 +35670,10 @@ // CHECK: [0x05,0x00,0x2c,0xd1,0x01,0x83,0x01,0x00] v_ashrrev_i16_e64 v5, v1, 0.5 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_ashrrev_i16_e64 v5, v1, -4.0 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_max_f16 v5, v1, v2 // CHECK: [0x01,0x05,0x0a,0x5a] @@ -36111,10 +36111,10 @@ // CHECK: [0x05,0x00,0x2f,0xd1,0xc1,0x04,0x02,0x00] v_max_u16_e64 v5, 0.5, v2 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_max_u16_e64 v5, -4.0, v2 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_max_u16_e64 v5, v1, v255 // CHECK: [0x05,0x00,0x2f,0xd1,0x01,0xff,0x03,0x00] @@ -36153,10 +36153,10 @@ // CHECK: [0x05,0x00,0x2f,0xd1,0x01,0x83,0x01,0x00] v_max_u16_e64 v5, v1, 0.5 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_max_u16_e64 v5, v1, -4.0 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_max_i16 v5, v1, v2 // CHECK: [0x01,0x05,0x0a,0x60] @@ -36258,10 +36258,10 @@ // CHECK: [0x05,0x00,0x30,0xd1,0xc1,0x04,0x02,0x00] v_max_i16_e64 v5, 0.5, v2 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_max_i16_e64 v5, -4.0, v2 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_max_i16_e64 v5, v1, v255 // CHECK: [0x05,0x00,0x30,0xd1,0x01,0xff,0x03,0x00] @@ -36300,10 +36300,10 @@ // CHECK: [0x05,0x00,0x30,0xd1,0x01,0x83,0x01,0x00] v_max_i16_e64 v5, v1, 0.5 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_max_i16_e64 v5, v1, -4.0 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_min_u16 v5, v1, v2 // CHECK: [0x01,0x05,0x0a,0x62] @@ -36405,10 +36405,10 @@ // CHECK: [0x05,0x00,0x31,0xd1,0xc1,0x04,0x02,0x00] v_min_u16_e64 v5, 0.5, v2 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_min_u16_e64 v5, -4.0, v2 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_min_u16_e64 v5, v1, v255 // CHECK: [0x05,0x00,0x31,0xd1,0x01,0xff,0x03,0x00] @@ -36447,10 +36447,10 @@ // CHECK: [0x05,0x00,0x31,0xd1,0x01,0x83,0x01,0x00] v_min_u16_e64 v5, v1, 0.5 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_min_u16_e64 v5, v1, -4.0 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_min_i16 v5, v1, v2 // CHECK: [0x01,0x05,0x0a,0x64] @@ -36552,10 +36552,10 @@ // CHECK: [0x05,0x00,0x32,0xd1,0xc1,0x04,0x02,0x00] v_min_i16_e64 v5, 0.5, v2 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_min_i16_e64 v5, -4.0, v2 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_min_i16_e64 v5, v1, v255 // CHECK: [0x05,0x00,0x32,0xd1,0x01,0xff,0x03,0x00] @@ -36594,10 +36594,10 @@ // CHECK: [0x05,0x00,0x32,0xd1,0x01,0x83,0x01,0x00] v_min_i16_e64 v5, v1, 0.5 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_min_i16_e64 v5, v1, -4.0 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_ldexp_f16 v5, v1, v2 // CHECK: [0x01,0x05,0x0a,0x66] @@ -43764,10 +43764,10 @@ // CHECK: [0x05,0x00,0xf1,0xd1,0xc1,0x04,0x0e,0x04] v_mad_u32_u16 v5, 0.5, v2, v3 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_mad_u32_u16 v5, -4.0, v2, v3 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_mad_u32_u16 v5, v1, v255, v3 // CHECK: [0x05,0x00,0xf1,0xd1,0x01,0xff,0x0f,0x04] @@ -43806,10 +43806,10 @@ // CHECK: [0x05,0x00,0xf1,0xd1,0x01,0x83,0x0d,0x04] v_mad_u32_u16 v5, v1, 0.5, v3 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_mad_u32_u16 v5, v1, -4.0, v3 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_mad_u32_u16 v5, v1, v2, v255 // CHECK: [0x05,0x00,0xf1,0xd1,0x01,0x05,0xfe,0x07] @@ -43917,10 +43917,10 @@ // CHECK: [0x05,0x00,0xf2,0xd1,0xc1,0x04,0x0e,0x04] v_mad_i32_i16 v5, 0.5, v2, v3 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_mad_i32_i16 v5, -4.0, v2, v3 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_mad_i32_i16 v5, v1, v255, v3 // CHECK: [0x05,0x00,0xf2,0xd1,0x01,0xff,0x0f,0x04] @@ -43959,10 +43959,10 @@ // CHECK: [0x05,0x00,0xf2,0xd1,0x01,0x83,0x0d,0x04] v_mad_i32_i16 v5, v1, 0.5, v3 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_mad_i32_i16 v5, v1, -4.0, v3 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_mad_i32_i16 v5, v1, v2, v255 // CHECK: [0x05,0x00,0xf2,0xd1,0x01,0x05,0xfe,0x07] @@ -44379,10 +44379,10 @@ // CHECK: [0x05,0x00,0xf5,0xd1,0xc1,0x04,0x0e,0x04] v_min3_i16 v5, 0.5, v2, v3 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_min3_i16 v5, -4.0, v2, v3 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_min3_i16 v5, v1, v255, v3 // CHECK: [0x05,0x00,0xf5,0xd1,0x01,0xff,0x0f,0x04] @@ -44421,10 +44421,10 @@ // CHECK: [0x05,0x00,0xf5,0xd1,0x01,0x83,0x0d,0x04] v_min3_i16 v5, v1, 0.5, v3 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_min3_i16 v5, v1, -4.0, v3 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_min3_i16 v5, v1, v2, v255 // CHECK: [0x05,0x00,0xf5,0xd1,0x01,0x05,0xfe,0x07] @@ -44463,10 +44463,10 @@ // CHECK: [0x05,0x00,0xf5,0xd1,0x01,0x05,0x06,0x03] v_min3_i16 v5, v1, v2, 0.5 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_min3_i16 v5, v1, v2, -4.0 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_min3_i16 v5, v1, v2, v3 op_sel:[0,0,0,0] // CHECK: [0x05,0x00,0xf5,0xd1,0x01,0x05,0x0e,0x04] @@ -44529,10 +44529,10 @@ // CHECK: [0x05,0x00,0xf6,0xd1,0xc1,0x04,0x0e,0x04] v_min3_u16 v5, 0.5, v2, v3 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_min3_u16 v5, -4.0, v2, v3 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_min3_u16 v5, v1, v255, v3 // CHECK: [0x05,0x00,0xf6,0xd1,0x01,0xff,0x0f,0x04] @@ -44571,10 +44571,10 @@ // CHECK: [0x05,0x00,0xf6,0xd1,0x01,0x83,0x0d,0x04] v_min3_u16 v5, v1, 0.5, v3 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_min3_u16 v5, v1, -4.0, v3 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_min3_u16 v5, v1, v2, v255 // CHECK: [0x05,0x00,0xf6,0xd1,0x01,0x05,0xfe,0x07] @@ -44613,10 +44613,10 @@ // CHECK: [0x05,0x00,0xf6,0xd1,0x01,0x05,0x06,0x03] v_min3_u16 v5, v1, v2, 0.5 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_min3_u16 v5, v1, v2, -4.0 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_min3_u16 v5, v1, v2, v3 op_sel:[0,0,0,0] // CHECK: [0x05,0x00,0xf6,0xd1,0x01,0x05,0x0e,0x04] @@ -44856,10 +44856,10 @@ // CHECK: [0x05,0x00,0xf8,0xd1,0xc1,0x04,0x0e,0x04] v_max3_i16 v5, 0.5, v2, v3 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_max3_i16 v5, -4.0, v2, v3 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_max3_i16 v5, v1, v255, v3 // CHECK: [0x05,0x00,0xf8,0xd1,0x01,0xff,0x0f,0x04] @@ -44898,10 +44898,10 @@ // CHECK: [0x05,0x00,0xf8,0xd1,0x01,0x83,0x0d,0x04] v_max3_i16 v5, v1, 0.5, v3 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_max3_i16 v5, v1, -4.0, v3 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_max3_i16 v5, v1, v2, v255 // CHECK: [0x05,0x00,0xf8,0xd1,0x01,0x05,0xfe,0x07] @@ -44940,10 +44940,10 @@ // CHECK: [0x05,0x00,0xf8,0xd1,0x01,0x05,0x06,0x03] v_max3_i16 v5, v1, v2, 0.5 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_max3_i16 v5, v1, v2, -4.0 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_max3_i16 v5, v1, v2, v3 op_sel:[0,0,0,0] // CHECK: [0x05,0x00,0xf8,0xd1,0x01,0x05,0x0e,0x04] @@ -45006,10 +45006,10 @@ // CHECK: [0x05,0x00,0xf9,0xd1,0xc1,0x04,0x0e,0x04] v_max3_u16 v5, 0.5, v2, v3 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_max3_u16 v5, -4.0, v2, v3 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_max3_u16 v5, v1, v255, v3 // CHECK: [0x05,0x00,0xf9,0xd1,0x01,0xff,0x0f,0x04] @@ -45048,10 +45048,10 @@ // CHECK: [0x05,0x00,0xf9,0xd1,0x01,0x83,0x0d,0x04] v_max3_u16 v5, v1, 0.5, v3 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_max3_u16 v5, v1, -4.0, v3 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_max3_u16 v5, v1, v2, v255 // CHECK: [0x05,0x00,0xf9,0xd1,0x01,0x05,0xfe,0x07] @@ -45090,10 +45090,10 @@ // CHECK: [0x05,0x00,0xf9,0xd1,0x01,0x05,0x06,0x03] v_max3_u16 v5, v1, v2, 0.5 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_max3_u16 v5, v1, v2, -4.0 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_max3_u16 v5, v1, v2, v3 op_sel:[0,0,0,0] // CHECK: [0x05,0x00,0xf9,0xd1,0x01,0x05,0x0e,0x04] @@ -45333,10 +45333,10 @@ // CHECK: [0x05,0x00,0xfb,0xd1,0xc1,0x04,0x0e,0x04] v_med3_i16 v5, 0.5, v2, v3 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_med3_i16 v5, -4.0, v2, v3 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_med3_i16 v5, v1, v255, v3 // CHECK: [0x05,0x00,0xfb,0xd1,0x01,0xff,0x0f,0x04] @@ -45375,10 +45375,10 @@ // CHECK: [0x05,0x00,0xfb,0xd1,0x01,0x83,0x0d,0x04] v_med3_i16 v5, v1, 0.5, v3 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_med3_i16 v5, v1, -4.0, v3 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_med3_i16 v5, v1, v2, v255 // CHECK: [0x05,0x00,0xfb,0xd1,0x01,0x05,0xfe,0x07] @@ -45417,10 +45417,10 @@ // CHECK: [0x05,0x00,0xfb,0xd1,0x01,0x05,0x06,0x03] v_med3_i16 v5, v1, v2, 0.5 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_med3_i16 v5, v1, v2, -4.0 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_med3_i16 v5, v1, v2, v3 op_sel:[0,0,0,0] // CHECK: [0x05,0x00,0xfb,0xd1,0x01,0x05,0x0e,0x04] @@ -45483,10 +45483,10 @@ // CHECK: [0x05,0x00,0xfc,0xd1,0xc1,0x04,0x0e,0x04] v_med3_u16 v5, 0.5, v2, v3 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_med3_u16 v5, -4.0, v2, v3 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_med3_u16 v5, v1, v255, v3 // CHECK: [0x05,0x00,0xfc,0xd1,0x01,0xff,0x0f,0x04] @@ -45525,10 +45525,10 @@ // CHECK: [0x05,0x00,0xfc,0xd1,0x01,0x83,0x0d,0x04] v_med3_u16 v5, v1, 0.5, v3 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_med3_u16 v5, v1, -4.0, v3 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_med3_u16 v5, v1, v2, v255 // CHECK: [0x05,0x00,0xfc,0xd1,0x01,0x05,0xfe,0x07] @@ -45567,10 +45567,10 @@ // CHECK: [0x05,0x00,0xfc,0xd1,0x01,0x05,0x06,0x03] v_med3_u16 v5, v1, v2, 0.5 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_med3_u16 v5, v1, v2, -4.0 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_med3_u16 v5, v1, v2, v3 op_sel:[0,0,0,0] // CHECK: [0x05,0x00,0xfc,0xd1,0x01,0x05,0x0e,0x04] @@ -46602,10 +46602,10 @@ // CHECK: [0x05,0x00,0x04,0xd2,0xc1,0x04,0x0e,0x04] v_mad_u16 v5, 0.5, v2, v3 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_mad_u16 v5, -4.0, v2, v3 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_mad_u16 v5, v1, v255, v3 // CHECK: [0x05,0x00,0x04,0xd2,0x01,0xff,0x0f,0x04] @@ -46644,10 +46644,10 @@ // CHECK: [0x05,0x00,0x04,0xd2,0x01,0x83,0x0d,0x04] v_mad_u16 v5, v1, 0.5, v3 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_mad_u16 v5, v1, -4.0, v3 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_mad_u16 v5, v1, v2, v255 // CHECK: [0x05,0x00,0x04,0xd2,0x01,0x05,0xfe,0x07] @@ -46686,10 +46686,10 @@ // CHECK: [0x05,0x00,0x04,0xd2,0x01,0x05,0x06,0x03] v_mad_u16 v5, v1, v2, 0.5 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_mad_u16 v5, v1, v2, -4.0 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_mad_u16 v5, v1, v2, v3 op_sel:[0,0,0,0] // CHECK: [0x05,0x00,0x04,0xd2,0x01,0x05,0x0e,0x04] @@ -46755,10 +46755,10 @@ // CHECK: [0x05,0x00,0x05,0xd2,0xc1,0x04,0x0e,0x04] v_mad_i16 v5, 0.5, v2, v3 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_mad_i16 v5, -4.0, v2, v3 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_mad_i16 v5, v1, v255, v3 // CHECK: [0x05,0x00,0x05,0xd2,0x01,0xff,0x0f,0x04] @@ -46797,10 +46797,10 @@ // CHECK: [0x05,0x00,0x05,0xd2,0x01,0x83,0x0d,0x04] v_mad_i16 v5, v1, 0.5, v3 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_mad_i16 v5, v1, -4.0, v3 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_mad_i16 v5, v1, v2, v255 // CHECK: [0x05,0x00,0x05,0xd2,0x01,0x05,0xfe,0x07] @@ -46839,10 +46839,10 @@ // CHECK: [0x05,0x00,0x05,0xd2,0x01,0x05,0x06,0x03] v_mad_i16 v5, v1, v2, 0.5 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_mad_i16 v5, v1, v2, -4.0 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_mad_i16 v5, v1, v2, v3 op_sel:[0,0,0,0] // CHECK: [0x05,0x00,0x05,0xd2,0x01,0x05,0x0e,0x04] @@ -50052,10 +50052,10 @@ // CHECK: [0x05,0x00,0x9e,0xd2,0xc1,0x04,0x02,0x00] v_add_i16 v5, 0.5, v2 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_add_i16 v5, -4.0, v2 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_add_i16 v5, v1, v255 // CHECK: [0x05,0x00,0x9e,0xd2,0x01,0xff,0x03,0x00] @@ -50094,10 +50094,10 @@ // CHECK: [0x05,0x00,0x9e,0xd2,0x01,0x83,0x01,0x00] v_add_i16 v5, v1, 0.5 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_add_i16 v5, v1, -4.0 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_add_i16 v5, v1, v2 op_sel:[0,0,0] // CHECK: [0x05,0x00,0x9e,0xd2,0x01,0x05,0x02,0x00] @@ -50160,10 +50160,10 @@ // CHECK: [0x05,0x00,0x9f,0xd2,0xc1,0x04,0x02,0x00] v_sub_i16 v5, 0.5, v2 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_sub_i16 v5, -4.0, v2 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_sub_i16 v5, v1, v255 // CHECK: [0x05,0x00,0x9f,0xd2,0x01,0xff,0x03,0x00] @@ -50202,10 +50202,10 @@ // CHECK: [0x05,0x00,0x9f,0xd2,0x01,0x83,0x01,0x00] v_sub_i16 v5, v1, 0.5 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_sub_i16 v5, v1, -4.0 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_sub_i16 v5, v1, v2 op_sel:[0,0,0] // CHECK: [0x05,0x00,0x9f,0xd2,0x01,0x05,0x02,0x00] @@ -50391,10 +50391,10 @@ // CHECK: [0x05,0x40,0x80,0xd3,0xc1,0x04,0x0e,0x1c] v_pk_mad_i16 v5, 0.5, v2, v3 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_pk_mad_i16 v5, -4.0, v2, v3 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_pk_mad_i16 v5, v1, v255, v3 // CHECK: [0x05,0x40,0x80,0xd3,0x01,0xff,0x0f,0x1c] @@ -50433,10 +50433,10 @@ // CHECK: [0x05,0x40,0x80,0xd3,0x01,0x83,0x0d,0x1c] v_pk_mad_i16 v5, v1, 0.5, v3 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_pk_mad_i16 v5, v1, -4.0, v3 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_pk_mad_i16 v5, v1, v2, v255 // CHECK: [0x05,0x40,0x80,0xd3,0x01,0x05,0xfe,0x1f] @@ -50475,10 +50475,10 @@ // CHECK: [0x05,0x40,0x80,0xd3,0x01,0x05,0x06,0x1b] v_pk_mad_i16 v5, v1, v2, 0.5 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_pk_mad_i16 v5, v1, v2, -4.0 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_pk_mad_i16 v5, v1, v2, v3 op_sel:[1,0,0] // CHECK: [0x05,0x48,0x80,0xd3,0x01,0x05,0x0e,0x1c] @@ -50550,10 +50550,10 @@ // CHECK: [0x05,0x00,0x81,0xd3,0xc1,0x04,0x02,0x18] v_pk_mul_lo_u16 v5, 0.5, v2 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_pk_mul_lo_u16 v5, -4.0, v2 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_pk_mul_lo_u16 v5, v1, v255 // CHECK: [0x05,0x00,0x81,0xd3,0x01,0xff,0x03,0x18] @@ -50595,10 +50595,10 @@ // CHECK: [0x05,0x00,0x81,0xd3,0x01,0x83,0x01,0x18] v_pk_mul_lo_u16 v5, v1, 0.5 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_pk_mul_lo_u16 v5, v1, -4.0 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_pk_mul_lo_u16 v5, v1, v2 op_sel:[1,0] // CHECK: [0x05,0x08,0x81,0xd3,0x01,0x05,0x02,0x18] @@ -50664,10 +50664,10 @@ // CHECK: [0x05,0x00,0x82,0xd3,0xc1,0x04,0x02,0x18] v_pk_add_i16 v5, 0.5, v2 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_pk_add_i16 v5, -4.0, v2 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_pk_add_i16 v5, v1, v255 // CHECK: [0x05,0x00,0x82,0xd3,0x01,0xff,0x03,0x18] @@ -50709,10 +50709,10 @@ // CHECK: [0x05,0x00,0x82,0xd3,0x01,0x83,0x01,0x18] v_pk_add_i16 v5, v1, 0.5 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_pk_add_i16 v5, v1, -4.0 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_pk_add_i16 v5, v1, v2 op_sel:[1,0] // CHECK: [0x05,0x08,0x82,0xd3,0x01,0x05,0x02,0x18] @@ -50781,10 +50781,10 @@ // CHECK: [0x05,0x00,0x83,0xd3,0xc1,0x04,0x02,0x18] v_pk_sub_i16 v5, 0.5, v2 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_pk_sub_i16 v5, -4.0, v2 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_pk_sub_i16 v5, v1, v255 // CHECK: [0x05,0x00,0x83,0xd3,0x01,0xff,0x03,0x18] @@ -50826,10 +50826,10 @@ // CHECK: [0x05,0x00,0x83,0xd3,0x01,0x83,0x01,0x18] v_pk_sub_i16 v5, v1, 0.5 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_pk_sub_i16 v5, v1, -4.0 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_pk_sub_i16 v5, v1, v2 op_sel:[1,0] // CHECK: [0x05,0x08,0x83,0xd3,0x01,0x05,0x02,0x18] @@ -50898,10 +50898,10 @@ // CHECK: [0x05,0x00,0x84,0xd3,0xc1,0x04,0x02,0x18] v_pk_lshlrev_b16 v5, 0.5, v2 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_pk_lshlrev_b16 v5, -4.0, v2 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_pk_lshlrev_b16 v5, v1, v255 // CHECK: [0x05,0x00,0x84,0xd3,0x01,0xff,0x03,0x18] @@ -50943,10 +50943,10 @@ // CHECK: [0x05,0x00,0x84,0xd3,0x01,0x83,0x01,0x18] v_pk_lshlrev_b16 v5, v1, 0.5 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_pk_lshlrev_b16 v5, v1, -4.0 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_pk_lshlrev_b16 v5, v1, v2 op_sel:[1,0] // CHECK: [0x05,0x08,0x84,0xd3,0x01,0x05,0x02,0x18] @@ -51012,10 +51012,10 @@ // CHECK: [0x05,0x00,0x85,0xd3,0xc1,0x04,0x02,0x18] v_pk_lshrrev_b16 v5, 0.5, v2 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_pk_lshrrev_b16 v5, -4.0, v2 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_pk_lshrrev_b16 v5, v1, v255 // CHECK: [0x05,0x00,0x85,0xd3,0x01,0xff,0x03,0x18] @@ -51057,10 +51057,10 @@ // CHECK: [0x05,0x00,0x85,0xd3,0x01,0x83,0x01,0x18] v_pk_lshrrev_b16 v5, v1, 0.5 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_pk_lshrrev_b16 v5, v1, -4.0 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_pk_lshrrev_b16 v5, v1, v2 op_sel:[1,0] // CHECK: [0x05,0x08,0x85,0xd3,0x01,0x05,0x02,0x18] @@ -51126,10 +51126,10 @@ // CHECK: [0x05,0x00,0x86,0xd3,0xc1,0x04,0x02,0x18] v_pk_ashrrev_i16 v5, 0.5, v2 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_pk_ashrrev_i16 v5, -4.0, v2 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_pk_ashrrev_i16 v5, v1, v255 // CHECK: [0x05,0x00,0x86,0xd3,0x01,0xff,0x03,0x18] @@ -51171,10 +51171,10 @@ // CHECK: [0x05,0x00,0x86,0xd3,0x01,0x83,0x01,0x18] v_pk_ashrrev_i16 v5, v1, 0.5 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_pk_ashrrev_i16 v5, v1, -4.0 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_pk_ashrrev_i16 v5, v1, v2 op_sel:[1,0] // CHECK: [0x05,0x08,0x86,0xd3,0x01,0x05,0x02,0x18] @@ -51240,10 +51240,10 @@ // CHECK: [0x05,0x00,0x87,0xd3,0xc1,0x04,0x02,0x18] v_pk_max_i16 v5, 0.5, v2 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_pk_max_i16 v5, -4.0, v2 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_pk_max_i16 v5, v1, v255 // CHECK: [0x05,0x00,0x87,0xd3,0x01,0xff,0x03,0x18] @@ -51285,10 +51285,10 @@ // CHECK: [0x05,0x00,0x87,0xd3,0x01,0x83,0x01,0x18] v_pk_max_i16 v5, v1, 0.5 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_pk_max_i16 v5, v1, -4.0 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_pk_max_i16 v5, v1, v2 op_sel:[1,0] // CHECK: [0x05,0x08,0x87,0xd3,0x01,0x05,0x02,0x18] @@ -51354,10 +51354,10 @@ // CHECK: [0x05,0x00,0x88,0xd3,0xc1,0x04,0x02,0x18] v_pk_min_i16 v5, 0.5, v2 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_pk_min_i16 v5, -4.0, v2 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_pk_min_i16 v5, v1, v255 // CHECK: [0x05,0x00,0x88,0xd3,0x01,0xff,0x03,0x18] @@ -51399,10 +51399,10 @@ // CHECK: [0x05,0x00,0x88,0xd3,0x01,0x83,0x01,0x18] v_pk_min_i16 v5, v1, 0.5 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_pk_min_i16 v5, v1, -4.0 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_pk_min_i16 v5, v1, v2 op_sel:[1,0] // CHECK: [0x05,0x08,0x88,0xd3,0x01,0x05,0x02,0x18] @@ -51468,10 +51468,10 @@ // CHECK: [0x05,0x40,0x89,0xd3,0xc1,0x04,0x0e,0x1c] v_pk_mad_u16 v5, 0.5, v2, v3 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_pk_mad_u16 v5, -4.0, v2, v3 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_pk_mad_u16 v5, v1, v255, v3 // CHECK: [0x05,0x40,0x89,0xd3,0x01,0xff,0x0f,0x1c] @@ -51510,10 +51510,10 @@ // CHECK: [0x05,0x40,0x89,0xd3,0x01,0x83,0x0d,0x1c] v_pk_mad_u16 v5, v1, 0.5, v3 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_pk_mad_u16 v5, v1, -4.0, v3 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_pk_mad_u16 v5, v1, v2, v255 // CHECK: [0x05,0x40,0x89,0xd3,0x01,0x05,0xfe,0x1f] @@ -51552,10 +51552,10 @@ // CHECK: [0x05,0x40,0x89,0xd3,0x01,0x05,0x06,0x1b] v_pk_mad_u16 v5, v1, v2, 0.5 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_pk_mad_u16 v5, v1, v2, -4.0 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_pk_mad_u16 v5, v1, v2, v3 op_sel:[1,0,0] // CHECK: [0x05,0x48,0x89,0xd3,0x01,0x05,0x0e,0x1c] @@ -51627,10 +51627,10 @@ // CHECK: [0x05,0x00,0x8a,0xd3,0xc1,0x04,0x02,0x18] v_pk_add_u16 v5, 0.5, v2 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_pk_add_u16 v5, -4.0, v2 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_pk_add_u16 v5, v1, v255 // CHECK: [0x05,0x00,0x8a,0xd3,0x01,0xff,0x03,0x18] @@ -51672,10 +51672,10 @@ // CHECK: [0x05,0x00,0x8a,0xd3,0x01,0x83,0x01,0x18] v_pk_add_u16 v5, v1, 0.5 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_pk_add_u16 v5, v1, -4.0 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_pk_add_u16 v5, v1, v2 op_sel:[1,0] // CHECK: [0x05,0x08,0x8a,0xd3,0x01,0x05,0x02,0x18] @@ -51744,10 +51744,10 @@ // CHECK: [0x05,0x00,0x8b,0xd3,0xc1,0x04,0x02,0x18] v_pk_sub_u16 v5, 0.5, v2 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_pk_sub_u16 v5, -4.0, v2 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_pk_sub_u16 v5, v1, v255 // CHECK: [0x05,0x00,0x8b,0xd3,0x01,0xff,0x03,0x18] @@ -51789,10 +51789,10 @@ // CHECK: [0x05,0x00,0x8b,0xd3,0x01,0x83,0x01,0x18] v_pk_sub_u16 v5, v1, 0.5 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_pk_sub_u16 v5, v1, -4.0 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_pk_sub_u16 v5, v1, v2 op_sel:[1,0] // CHECK: [0x05,0x08,0x8b,0xd3,0x01,0x05,0x02,0x18] @@ -51861,10 +51861,10 @@ // CHECK: [0x05,0x00,0x8c,0xd3,0xc1,0x04,0x02,0x18] v_pk_max_u16 v5, 0.5, v2 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_pk_max_u16 v5, -4.0, v2 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_pk_max_u16 v5, v1, v255 // CHECK: [0x05,0x00,0x8c,0xd3,0x01,0xff,0x03,0x18] @@ -51906,10 +51906,10 @@ // CHECK: [0x05,0x00,0x8c,0xd3,0x01,0x83,0x01,0x18] v_pk_max_u16 v5, v1, 0.5 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_pk_max_u16 v5, v1, -4.0 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_pk_max_u16 v5, v1, v2 op_sel:[1,0] // CHECK: [0x05,0x08,0x8c,0xd3,0x01,0x05,0x02,0x18] @@ -51975,10 +51975,10 @@ // CHECK: [0x05,0x00,0x8d,0xd3,0xc1,0x04,0x02,0x18] v_pk_min_u16 v5, 0.5, v2 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_pk_min_u16 v5, -4.0, v2 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_pk_min_u16 v5, v1, v255 // CHECK: [0x05,0x00,0x8d,0xd3,0x01,0xff,0x03,0x18] @@ -52020,10 +52020,10 @@ // CHECK: [0x05,0x00,0x8d,0xd3,0x01,0x83,0x01,0x18] v_pk_min_u16 v5, v1, 0.5 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_pk_min_u16 v5, v1, -4.0 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_pk_min_u16 v5, v1, v2 op_sel:[1,0] // CHECK: [0x05,0x08,0x8d,0xd3,0x01,0x05,0x02,0x18] @@ -69471,10 +69471,10 @@ // CHECK: [0x0a,0x00,0xa0,0xd0,0xc1,0x04,0x02,0x00] v_cmp_f_i16_e64 s[10:11], 0.5, v2 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_cmp_f_i16_e64 s[10:11], -4.0, v2 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_cmp_f_i16_e64 s[10:11], v1, v255 // CHECK: [0x0a,0x00,0xa0,0xd0,0x01,0xff,0x03,0x00] @@ -69513,10 +69513,10 @@ // CHECK: [0x0a,0x00,0xa0,0xd0,0x01,0x83,0x01,0x00] v_cmp_f_i16_e64 s[10:11], v1, 0.5 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_cmp_f_i16_e64 s[10:11], v1, -4.0 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_cmp_lt_i16 vcc, v1, v2 // CHECK: [0x01,0x05,0x42,0x7d] @@ -69624,10 +69624,10 @@ // CHECK: [0x0a,0x00,0xa1,0xd0,0xc1,0x04,0x02,0x00] v_cmp_lt_i16_e64 s[10:11], 0.5, v2 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_cmp_lt_i16_e64 s[10:11], -4.0, v2 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_cmp_lt_i16_e64 s[10:11], v1, v255 // CHECK: [0x0a,0x00,0xa1,0xd0,0x01,0xff,0x03,0x00] @@ -69666,10 +69666,10 @@ // CHECK: [0x0a,0x00,0xa1,0xd0,0x01,0x83,0x01,0x00] v_cmp_lt_i16_e64 s[10:11], v1, 0.5 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_cmp_lt_i16_e64 s[10:11], v1, -4.0 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_cmp_eq_i16 vcc, v1, v2 // CHECK: [0x01,0x05,0x44,0x7d] @@ -69777,10 +69777,10 @@ // CHECK: [0x0a,0x00,0xa2,0xd0,0xc1,0x04,0x02,0x00] v_cmp_eq_i16_e64 s[10:11], 0.5, v2 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_cmp_eq_i16_e64 s[10:11], -4.0, v2 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_cmp_eq_i16_e64 s[10:11], v1, v255 // CHECK: [0x0a,0x00,0xa2,0xd0,0x01,0xff,0x03,0x00] @@ -69819,10 +69819,10 @@ // CHECK: [0x0a,0x00,0xa2,0xd0,0x01,0x83,0x01,0x00] v_cmp_eq_i16_e64 s[10:11], v1, 0.5 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_cmp_eq_i16_e64 s[10:11], v1, -4.0 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_cmp_le_i16 vcc, v1, v2 // CHECK: [0x01,0x05,0x46,0x7d] @@ -69930,10 +69930,10 @@ // CHECK: [0x0a,0x00,0xa3,0xd0,0xc1,0x04,0x02,0x00] v_cmp_le_i16_e64 s[10:11], 0.5, v2 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_cmp_le_i16_e64 s[10:11], -4.0, v2 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_cmp_le_i16_e64 s[10:11], v1, v255 // CHECK: [0x0a,0x00,0xa3,0xd0,0x01,0xff,0x03,0x00] @@ -69972,10 +69972,10 @@ // CHECK: [0x0a,0x00,0xa3,0xd0,0x01,0x83,0x01,0x00] v_cmp_le_i16_e64 s[10:11], v1, 0.5 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_cmp_le_i16_e64 s[10:11], v1, -4.0 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_cmp_gt_i16 vcc, v1, v2 // CHECK: [0x01,0x05,0x48,0x7d] @@ -70083,10 +70083,10 @@ // CHECK: [0x0a,0x00,0xa4,0xd0,0xc1,0x04,0x02,0x00] v_cmp_gt_i16_e64 s[10:11], 0.5, v2 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_cmp_gt_i16_e64 s[10:11], -4.0, v2 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_cmp_gt_i16_e64 s[10:11], v1, v255 // CHECK: [0x0a,0x00,0xa4,0xd0,0x01,0xff,0x03,0x00] @@ -70125,10 +70125,10 @@ // CHECK: [0x0a,0x00,0xa4,0xd0,0x01,0x83,0x01,0x00] v_cmp_gt_i16_e64 s[10:11], v1, 0.5 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_cmp_gt_i16_e64 s[10:11], v1, -4.0 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_cmp_ne_i16 vcc, v1, v2 // CHECK: [0x01,0x05,0x4a,0x7d] @@ -70236,10 +70236,10 @@ // CHECK: [0x0a,0x00,0xa5,0xd0,0xc1,0x04,0x02,0x00] v_cmp_ne_i16_e64 s[10:11], 0.5, v2 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_cmp_ne_i16_e64 s[10:11], -4.0, v2 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_cmp_ne_i16_e64 s[10:11], v1, v255 // CHECK: [0x0a,0x00,0xa5,0xd0,0x01,0xff,0x03,0x00] @@ -70278,10 +70278,10 @@ // CHECK: [0x0a,0x00,0xa5,0xd0,0x01,0x83,0x01,0x00] v_cmp_ne_i16_e64 s[10:11], v1, 0.5 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_cmp_ne_i16_e64 s[10:11], v1, -4.0 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_cmp_ge_i16 vcc, v1, v2 // CHECK: [0x01,0x05,0x4c,0x7d] @@ -70389,10 +70389,10 @@ // CHECK: [0x0a,0x00,0xa6,0xd0,0xc1,0x04,0x02,0x00] v_cmp_ge_i16_e64 s[10:11], 0.5, v2 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_cmp_ge_i16_e64 s[10:11], -4.0, v2 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_cmp_ge_i16_e64 s[10:11], v1, v255 // CHECK: [0x0a,0x00,0xa6,0xd0,0x01,0xff,0x03,0x00] @@ -70431,10 +70431,10 @@ // CHECK: [0x0a,0x00,0xa6,0xd0,0x01,0x83,0x01,0x00] v_cmp_ge_i16_e64 s[10:11], v1, 0.5 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_cmp_ge_i16_e64 s[10:11], v1, -4.0 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_cmp_t_i16 vcc, v1, v2 // CHECK: [0x01,0x05,0x4e,0x7d] @@ -70542,10 +70542,10 @@ // CHECK: [0x0a,0x00,0xa7,0xd0,0xc1,0x04,0x02,0x00] v_cmp_t_i16_e64 s[10:11], 0.5, v2 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_cmp_t_i16_e64 s[10:11], -4.0, v2 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_cmp_t_i16_e64 s[10:11], v1, v255 // CHECK: [0x0a,0x00,0xa7,0xd0,0x01,0xff,0x03,0x00] @@ -70584,10 +70584,10 @@ // CHECK: [0x0a,0x00,0xa7,0xd0,0x01,0x83,0x01,0x00] v_cmp_t_i16_e64 s[10:11], v1, 0.5 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_cmp_t_i16_e64 s[10:11], v1, -4.0 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_cmp_f_u16 vcc, v1, v2 // CHECK: [0x01,0x05,0x50,0x7d] @@ -70695,10 +70695,10 @@ // CHECK: [0x0a,0x00,0xa8,0xd0,0xc1,0x04,0x02,0x00] v_cmp_f_u16_e64 s[10:11], 0.5, v2 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_cmp_f_u16_e64 s[10:11], -4.0, v2 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_cmp_f_u16_e64 s[10:11], v1, v255 // CHECK: [0x0a,0x00,0xa8,0xd0,0x01,0xff,0x03,0x00] @@ -70737,10 +70737,10 @@ // CHECK: [0x0a,0x00,0xa8,0xd0,0x01,0x83,0x01,0x00] v_cmp_f_u16_e64 s[10:11], v1, 0.5 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_cmp_f_u16_e64 s[10:11], v1, -4.0 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_cmp_lt_u16 vcc, v1, v2 // CHECK: [0x01,0x05,0x52,0x7d] @@ -70848,10 +70848,10 @@ // CHECK: [0x0a,0x00,0xa9,0xd0,0xc1,0x04,0x02,0x00] v_cmp_lt_u16_e64 s[10:11], 0.5, v2 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_cmp_lt_u16_e64 s[10:11], -4.0, v2 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_cmp_lt_u16_e64 s[10:11], v1, v255 // CHECK: [0x0a,0x00,0xa9,0xd0,0x01,0xff,0x03,0x00] @@ -70890,10 +70890,10 @@ // CHECK: [0x0a,0x00,0xa9,0xd0,0x01,0x83,0x01,0x00] v_cmp_lt_u16_e64 s[10:11], v1, 0.5 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_cmp_lt_u16_e64 s[10:11], v1, -4.0 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_cmp_eq_u16 vcc, v1, v2 // CHECK: [0x01,0x05,0x54,0x7d] @@ -71001,10 +71001,10 @@ // CHECK: [0x0a,0x00,0xaa,0xd0,0xc1,0x04,0x02,0x00] v_cmp_eq_u16_e64 s[10:11], 0.5, v2 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_cmp_eq_u16_e64 s[10:11], -4.0, v2 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_cmp_eq_u16_e64 s[10:11], v1, v255 // CHECK: [0x0a,0x00,0xaa,0xd0,0x01,0xff,0x03,0x00] @@ -71043,10 +71043,10 @@ // CHECK: [0x0a,0x00,0xaa,0xd0,0x01,0x83,0x01,0x00] v_cmp_eq_u16_e64 s[10:11], v1, 0.5 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_cmp_eq_u16_e64 s[10:11], v1, -4.0 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_cmp_le_u16 vcc, v1, v2 // CHECK: [0x01,0x05,0x56,0x7d] @@ -71154,10 +71154,10 @@ // CHECK: [0x0a,0x00,0xab,0xd0,0xc1,0x04,0x02,0x00] v_cmp_le_u16_e64 s[10:11], 0.5, v2 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_cmp_le_u16_e64 s[10:11], -4.0, v2 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_cmp_le_u16_e64 s[10:11], v1, v255 // CHECK: [0x0a,0x00,0xab,0xd0,0x01,0xff,0x03,0x00] @@ -71196,10 +71196,10 @@ // CHECK: [0x0a,0x00,0xab,0xd0,0x01,0x83,0x01,0x00] v_cmp_le_u16_e64 s[10:11], v1, 0.5 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_cmp_le_u16_e64 s[10:11], v1, -4.0 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_cmp_gt_u16 vcc, v1, v2 // CHECK: [0x01,0x05,0x58,0x7d] @@ -71307,10 +71307,10 @@ // CHECK: [0x0a,0x00,0xac,0xd0,0xc1,0x04,0x02,0x00] v_cmp_gt_u16_e64 s[10:11], 0.5, v2 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_cmp_gt_u16_e64 s[10:11], -4.0, v2 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_cmp_gt_u16_e64 s[10:11], v1, v255 // CHECK: [0x0a,0x00,0xac,0xd0,0x01,0xff,0x03,0x00] @@ -71349,10 +71349,10 @@ // CHECK: [0x0a,0x00,0xac,0xd0,0x01,0x83,0x01,0x00] v_cmp_gt_u16_e64 s[10:11], v1, 0.5 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_cmp_gt_u16_e64 s[10:11], v1, -4.0 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_cmp_ne_u16 vcc, v1, v2 // CHECK: [0x01,0x05,0x5a,0x7d] @@ -71460,10 +71460,10 @@ // CHECK: [0x0a,0x00,0xad,0xd0,0xc1,0x04,0x02,0x00] v_cmp_ne_u16_e64 s[10:11], 0.5, v2 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_cmp_ne_u16_e64 s[10:11], -4.0, v2 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_cmp_ne_u16_e64 s[10:11], v1, v255 // CHECK: [0x0a,0x00,0xad,0xd0,0x01,0xff,0x03,0x00] @@ -71502,10 +71502,10 @@ // CHECK: [0x0a,0x00,0xad,0xd0,0x01,0x83,0x01,0x00] v_cmp_ne_u16_e64 s[10:11], v1, 0.5 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_cmp_ne_u16_e64 s[10:11], v1, -4.0 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_cmp_ge_u16 vcc, v1, v2 // CHECK: [0x01,0x05,0x5c,0x7d] @@ -71613,10 +71613,10 @@ // CHECK: [0x0a,0x00,0xae,0xd0,0xc1,0x04,0x02,0x00] v_cmp_ge_u16_e64 s[10:11], 0.5, v2 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_cmp_ge_u16_e64 s[10:11], -4.0, v2 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_cmp_ge_u16_e64 s[10:11], v1, v255 // CHECK: [0x0a,0x00,0xae,0xd0,0x01,0xff,0x03,0x00] @@ -71655,10 +71655,10 @@ // CHECK: [0x0a,0x00,0xae,0xd0,0x01,0x83,0x01,0x00] v_cmp_ge_u16_e64 s[10:11], v1, 0.5 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_cmp_ge_u16_e64 s[10:11], v1, -4.0 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_cmp_t_u16 vcc, v1, v2 // CHECK: [0x01,0x05,0x5e,0x7d] @@ -71766,10 +71766,10 @@ // CHECK: [0x0a,0x00,0xaf,0xd0,0xc1,0x04,0x02,0x00] v_cmp_t_u16_e64 s[10:11], 0.5, v2 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_cmp_t_u16_e64 s[10:11], -4.0, v2 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_cmp_t_u16_e64 s[10:11], v1, v255 // CHECK: [0x0a,0x00,0xaf,0xd0,0x01,0xff,0x03,0x00] @@ -71808,10 +71808,10 @@ // CHECK: [0x0a,0x00,0xaf,0xd0,0x01,0x83,0x01,0x00] v_cmp_t_u16_e64 s[10:11], v1, 0.5 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_cmp_t_u16_e64 s[10:11], v1, -4.0 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_cmpx_f_i16 vcc, v1, v2 // CHECK: [0x01,0x05,0x60,0x7d] @@ -71922,10 +71922,10 @@ // CHECK: [0x0a,0x00,0xb0,0xd0,0xc1,0x04,0x02,0x00] v_cmpx_f_i16_e64 s[10:11], 0.5, v2 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_cmpx_f_i16_e64 s[10:11], -4.0, v2 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_cmpx_f_i16_e64 s[10:11], v1, v255 // CHECK: [0x0a,0x00,0xb0,0xd0,0x01,0xff,0x03,0x00] @@ -71964,10 +71964,10 @@ // CHECK: [0x0a,0x00,0xb0,0xd0,0x01,0x83,0x01,0x00] v_cmpx_f_i16_e64 s[10:11], v1, 0.5 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_cmpx_f_i16_e64 s[10:11], v1, -4.0 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_cmpx_lt_i16 vcc, v1, v2 // CHECK: [0x01,0x05,0x62,0x7d] @@ -72078,10 +72078,10 @@ // CHECK: [0x0a,0x00,0xb1,0xd0,0xc1,0x04,0x02,0x00] v_cmpx_lt_i16_e64 s[10:11], 0.5, v2 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_cmpx_lt_i16_e64 s[10:11], -4.0, v2 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_cmpx_lt_i16_e64 s[10:11], v1, v255 // CHECK: [0x0a,0x00,0xb1,0xd0,0x01,0xff,0x03,0x00] @@ -72120,10 +72120,10 @@ // CHECK: [0x0a,0x00,0xb1,0xd0,0x01,0x83,0x01,0x00] v_cmpx_lt_i16_e64 s[10:11], v1, 0.5 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_cmpx_lt_i16_e64 s[10:11], v1, -4.0 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_cmpx_eq_i16 vcc, v1, v2 // CHECK: [0x01,0x05,0x64,0x7d] @@ -72234,10 +72234,10 @@ // CHECK: [0x0a,0x00,0xb2,0xd0,0xc1,0x04,0x02,0x00] v_cmpx_eq_i16_e64 s[10:11], 0.5, v2 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_cmpx_eq_i16_e64 s[10:11], -4.0, v2 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_cmpx_eq_i16_e64 s[10:11], v1, v255 // CHECK: [0x0a,0x00,0xb2,0xd0,0x01,0xff,0x03,0x00] @@ -72276,10 +72276,10 @@ // CHECK: [0x0a,0x00,0xb2,0xd0,0x01,0x83,0x01,0x00] v_cmpx_eq_i16_e64 s[10:11], v1, 0.5 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_cmpx_eq_i16_e64 s[10:11], v1, -4.0 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_cmpx_le_i16 vcc, v1, v2 // CHECK: [0x01,0x05,0x66,0x7d] @@ -72390,10 +72390,10 @@ // CHECK: [0x0a,0x00,0xb3,0xd0,0xc1,0x04,0x02,0x00] v_cmpx_le_i16_e64 s[10:11], 0.5, v2 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_cmpx_le_i16_e64 s[10:11], -4.0, v2 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_cmpx_le_i16_e64 s[10:11], v1, v255 // CHECK: [0x0a,0x00,0xb3,0xd0,0x01,0xff,0x03,0x00] @@ -72432,10 +72432,10 @@ // CHECK: [0x0a,0x00,0xb3,0xd0,0x01,0x83,0x01,0x00] v_cmpx_le_i16_e64 s[10:11], v1, 0.5 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_cmpx_le_i16_e64 s[10:11], v1, -4.0 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_cmpx_gt_i16 vcc, v1, v2 // CHECK: [0x01,0x05,0x68,0x7d] @@ -72546,10 +72546,10 @@ // CHECK: [0x0a,0x00,0xb4,0xd0,0xc1,0x04,0x02,0x00] v_cmpx_gt_i16_e64 s[10:11], 0.5, v2 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_cmpx_gt_i16_e64 s[10:11], -4.0, v2 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_cmpx_gt_i16_e64 s[10:11], v1, v255 // CHECK: [0x0a,0x00,0xb4,0xd0,0x01,0xff,0x03,0x00] @@ -72588,10 +72588,10 @@ // CHECK: [0x0a,0x00,0xb4,0xd0,0x01,0x83,0x01,0x00] v_cmpx_gt_i16_e64 s[10:11], v1, 0.5 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_cmpx_gt_i16_e64 s[10:11], v1, -4.0 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_cmpx_ne_i16 vcc, v1, v2 // CHECK: [0x01,0x05,0x6a,0x7d] @@ -72702,10 +72702,10 @@ // CHECK: [0x0a,0x00,0xb5,0xd0,0xc1,0x04,0x02,0x00] v_cmpx_ne_i16_e64 s[10:11], 0.5, v2 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_cmpx_ne_i16_e64 s[10:11], -4.0, v2 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_cmpx_ne_i16_e64 s[10:11], v1, v255 // CHECK: [0x0a,0x00,0xb5,0xd0,0x01,0xff,0x03,0x00] @@ -72744,10 +72744,10 @@ // CHECK: [0x0a,0x00,0xb5,0xd0,0x01,0x83,0x01,0x00] v_cmpx_ne_i16_e64 s[10:11], v1, 0.5 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_cmpx_ne_i16_e64 s[10:11], v1, -4.0 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_cmpx_ge_i16 vcc, v1, v2 // CHECK: [0x01,0x05,0x6c,0x7d] @@ -72858,10 +72858,10 @@ // CHECK: [0x0a,0x00,0xb6,0xd0,0xc1,0x04,0x02,0x00] v_cmpx_ge_i16_e64 s[10:11], 0.5, v2 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_cmpx_ge_i16_e64 s[10:11], -4.0, v2 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_cmpx_ge_i16_e64 s[10:11], v1, v255 // CHECK: [0x0a,0x00,0xb6,0xd0,0x01,0xff,0x03,0x00] @@ -72900,10 +72900,10 @@ // CHECK: [0x0a,0x00,0xb6,0xd0,0x01,0x83,0x01,0x00] v_cmpx_ge_i16_e64 s[10:11], v1, 0.5 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_cmpx_ge_i16_e64 s[10:11], v1, -4.0 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_cmpx_t_i16 vcc, v1, v2 // CHECK: [0x01,0x05,0x6e,0x7d] @@ -73014,10 +73014,10 @@ // CHECK: [0x0a,0x00,0xb7,0xd0,0xc1,0x04,0x02,0x00] v_cmpx_t_i16_e64 s[10:11], 0.5, v2 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_cmpx_t_i16_e64 s[10:11], -4.0, v2 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_cmpx_t_i16_e64 s[10:11], v1, v255 // CHECK: [0x0a,0x00,0xb7,0xd0,0x01,0xff,0x03,0x00] @@ -73056,10 +73056,10 @@ // CHECK: [0x0a,0x00,0xb7,0xd0,0x01,0x83,0x01,0x00] v_cmpx_t_i16_e64 s[10:11], v1, 0.5 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_cmpx_t_i16_e64 s[10:11], v1, -4.0 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_cmpx_f_u16 vcc, v1, v2 // CHECK: [0x01,0x05,0x70,0x7d] @@ -73170,10 +73170,10 @@ // CHECK: [0x0a,0x00,0xb8,0xd0,0xc1,0x04,0x02,0x00] v_cmpx_f_u16_e64 s[10:11], 0.5, v2 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_cmpx_f_u16_e64 s[10:11], -4.0, v2 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_cmpx_f_u16_e64 s[10:11], v1, v255 // CHECK: [0x0a,0x00,0xb8,0xd0,0x01,0xff,0x03,0x00] @@ -73212,10 +73212,10 @@ // CHECK: [0x0a,0x00,0xb8,0xd0,0x01,0x83,0x01,0x00] v_cmpx_f_u16_e64 s[10:11], v1, 0.5 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_cmpx_f_u16_e64 s[10:11], v1, -4.0 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_cmpx_lt_u16 vcc, v1, v2 // CHECK: [0x01,0x05,0x72,0x7d] @@ -73326,10 +73326,10 @@ // CHECK: [0x0a,0x00,0xb9,0xd0,0xc1,0x04,0x02,0x00] v_cmpx_lt_u16_e64 s[10:11], 0.5, v2 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_cmpx_lt_u16_e64 s[10:11], -4.0, v2 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_cmpx_lt_u16_e64 s[10:11], v1, v255 // CHECK: [0x0a,0x00,0xb9,0xd0,0x01,0xff,0x03,0x00] @@ -73368,10 +73368,10 @@ // CHECK: [0x0a,0x00,0xb9,0xd0,0x01,0x83,0x01,0x00] v_cmpx_lt_u16_e64 s[10:11], v1, 0.5 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_cmpx_lt_u16_e64 s[10:11], v1, -4.0 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_cmpx_eq_u16 vcc, v1, v2 // CHECK: [0x01,0x05,0x74,0x7d] @@ -73482,10 +73482,10 @@ // CHECK: [0x0a,0x00,0xba,0xd0,0xc1,0x04,0x02,0x00] v_cmpx_eq_u16_e64 s[10:11], 0.5, v2 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_cmpx_eq_u16_e64 s[10:11], -4.0, v2 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_cmpx_eq_u16_e64 s[10:11], v1, v255 // CHECK: [0x0a,0x00,0xba,0xd0,0x01,0xff,0x03,0x00] @@ -73524,10 +73524,10 @@ // CHECK: [0x0a,0x00,0xba,0xd0,0x01,0x83,0x01,0x00] v_cmpx_eq_u16_e64 s[10:11], v1, 0.5 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_cmpx_eq_u16_e64 s[10:11], v1, -4.0 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_cmpx_le_u16 vcc, v1, v2 // CHECK: [0x01,0x05,0x76,0x7d] @@ -73638,10 +73638,10 @@ // CHECK: [0x0a,0x00,0xbb,0xd0,0xc1,0x04,0x02,0x00] v_cmpx_le_u16_e64 s[10:11], 0.5, v2 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_cmpx_le_u16_e64 s[10:11], -4.0, v2 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_cmpx_le_u16_e64 s[10:11], v1, v255 // CHECK: [0x0a,0x00,0xbb,0xd0,0x01,0xff,0x03,0x00] @@ -73680,10 +73680,10 @@ // CHECK: [0x0a,0x00,0xbb,0xd0,0x01,0x83,0x01,0x00] v_cmpx_le_u16_e64 s[10:11], v1, 0.5 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_cmpx_le_u16_e64 s[10:11], v1, -4.0 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_cmpx_gt_u16 vcc, v1, v2 // CHECK: [0x01,0x05,0x78,0x7d] @@ -73794,10 +73794,10 @@ // CHECK: [0x0a,0x00,0xbc,0xd0,0xc1,0x04,0x02,0x00] v_cmpx_gt_u16_e64 s[10:11], 0.5, v2 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_cmpx_gt_u16_e64 s[10:11], -4.0, v2 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_cmpx_gt_u16_e64 s[10:11], v1, v255 // CHECK: [0x0a,0x00,0xbc,0xd0,0x01,0xff,0x03,0x00] @@ -73836,10 +73836,10 @@ // CHECK: [0x0a,0x00,0xbc,0xd0,0x01,0x83,0x01,0x00] v_cmpx_gt_u16_e64 s[10:11], v1, 0.5 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_cmpx_gt_u16_e64 s[10:11], v1, -4.0 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_cmpx_ne_u16 vcc, v1, v2 // CHECK: [0x01,0x05,0x7a,0x7d] @@ -73950,10 +73950,10 @@ // CHECK: [0x0a,0x00,0xbd,0xd0,0xc1,0x04,0x02,0x00] v_cmpx_ne_u16_e64 s[10:11], 0.5, v2 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_cmpx_ne_u16_e64 s[10:11], -4.0, v2 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_cmpx_ne_u16_e64 s[10:11], v1, v255 // CHECK: [0x0a,0x00,0xbd,0xd0,0x01,0xff,0x03,0x00] @@ -73992,10 +73992,10 @@ // CHECK: [0x0a,0x00,0xbd,0xd0,0x01,0x83,0x01,0x00] v_cmpx_ne_u16_e64 s[10:11], v1, 0.5 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_cmpx_ne_u16_e64 s[10:11], v1, -4.0 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_cmpx_ge_u16 vcc, v1, v2 // CHECK: [0x01,0x05,0x7c,0x7d] @@ -74106,10 +74106,10 @@ // CHECK: [0x0a,0x00,0xbe,0xd0,0xc1,0x04,0x02,0x00] v_cmpx_ge_u16_e64 s[10:11], 0.5, v2 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_cmpx_ge_u16_e64 s[10:11], -4.0, v2 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_cmpx_ge_u16_e64 s[10:11], v1, v255 // CHECK: [0x0a,0x00,0xbe,0xd0,0x01,0xff,0x03,0x00] @@ -74148,10 +74148,10 @@ // CHECK: [0x0a,0x00,0xbe,0xd0,0x01,0x83,0x01,0x00] v_cmpx_ge_u16_e64 s[10:11], v1, 0.5 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_cmpx_ge_u16_e64 s[10:11], v1, -4.0 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_cmpx_t_u16 vcc, v1, v2 // CHECK: [0x01,0x05,0x7e,0x7d] @@ -74262,10 +74262,10 @@ // CHECK: [0x0a,0x00,0xbf,0xd0,0xc1,0x04,0x02,0x00] v_cmpx_t_u16_e64 s[10:11], 0.5, v2 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_cmpx_t_u16_e64 s[10:11], -4.0, v2 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_cmpx_t_u16_e64 s[10:11], v1, v255 // CHECK: [0x0a,0x00,0xbf,0xd0,0x01,0xff,0x03,0x00] @@ -74304,10 +74304,10 @@ // CHECK: [0x0a,0x00,0xbf,0xd0,0x01,0x83,0x01,0x00] v_cmpx_t_u16_e64 s[10:11], v1, 0.5 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_cmpx_t_u16_e64 s[10:11], v1, -4.0 -// CHECK-ERR: error: invalid literal operand +// CHECK-ERR: error: literal operands are not supported v_cmp_f_i32 vcc, v1, v2 // CHECK: [0x01,0x05,0x80,0x7d] diff --git a/llvm/test/MC/AMDGPU/gfx9_err_pos.s b/llvm/test/MC/AMDGPU/gfx9_err_pos.s --- a/llvm/test/MC/AMDGPU/gfx9_err_pos.s +++ b/llvm/test/MC/AMDGPU/gfx9_err_pos.s @@ -124,6 +124,29 @@ // CHECK-NEXT:{{^}}v_add_f32_e64 v0, flat_scratch_hi, m0 // CHECK-NEXT:{{^}}^ +//============================================================================== +// literal operands are not supported + +v_bfe_u32 v0, v2, v3, undef +// CHECK: error: literal operands are not supported +// CHECK-NEXT:{{^}}v_bfe_u32 v0, v2, v3, undef +// CHECK-NEXT:{{^}} ^ + +v_bfe_u32 v0, v2, undef, v3 +// CHECK: error: literal operands are not supported +// CHECK-NEXT:{{^}}v_bfe_u32 v0, v2, undef, v3 +// CHECK-NEXT:{{^}} ^ + +v_add_i16 v5, v1, 0.5 +// CHECK: error: literal operands are not supported +// CHECK-NEXT:{{^}}v_add_i16 v5, v1, 0.5 +// CHECK-NEXT:{{^}} ^ + +v_add_i16 v5, 0.5, v2 +// CHECK: error: literal operands are not supported +// CHECK-NEXT:{{^}}v_add_i16 v5, 0.5, v2 +// CHECK-NEXT:{{^}} ^ + //============================================================================== // r128 modifier is not supported on this GPU diff --git a/llvm/test/MC/AMDGPU/literals.s b/llvm/test/MC/AMDGPU/literals.s --- a/llvm/test/MC/AMDGPU/literals.s +++ b/llvm/test/MC/AMDGPU/literals.s @@ -282,12 +282,12 @@ // GFX89: v_fract_f64_e32 v[0:1], 0x4d2 ; encoding: [0xff,0x64,0x00,0x7e,0xd2,0x04,0x00,0x00] v_fract_f64_e32 v[0:1], 1234 -// NOSICI: error: invalid literal operand -// NOGFX89: error: invalid literal operand +// NOSICI: error: literal operands are not supported +// NOGFX89: error: literal operands are not supported v_trunc_f32_e64 v0, 1234 -// NOSICI: error: invalid literal operand -// NOGFX89: error: invalid literal operand +// NOSICI: error: literal operands are not supported +// NOGFX89: error: literal operands are not supported v_fract_f64_e64 v[0:1], 1234 // SICI: v_trunc_f32_e32 v0, 0xffff2bcf ; encoding: [0xff,0x42,0x00,0x7e,0xcf,0x2b,0xff,0xff] @@ -378,8 +378,8 @@ // GFX89: v_and_b32_e32 v0, 0x4d2, v1 ; encoding: [0xff,0x02,0x00,0x26,0xd2,0x04,0x00,0x00] v_and_b32_e32 v0, 1234, v1 -// NOSICI: error: invalid literal operand -// NOGFX89: error: invalid literal operand +// NOSICI: error: literal operands are not supported +// NOGFX89: error: literal operands are not supported v_and_b32_e64 v0, 1234, v1 // SICI: s_mov_b64 s[0:1], 0xffff2bcf ; encoding: [0xff,0x04,0x80,0xbe,0xcf,0x2b,0xff,0xff] @@ -450,12 +450,12 @@ // GFX89: v_fract_f64_e64 v[0:1], 0.15915494309189532 ; encoding: [0x00,0x00,0x72,0xd1,0xf8,0x00,0x00,0x00] v_fract_f64_e64 v[0:1], 0x3fc45f306dc9c882 -// NOSICI: error: invalid literal operand +// NOSICI: error: literal operands are not supported // GFX89: v_trunc_f32_e64 v0, 0.15915494 ; encoding: [0x00,0x00,0x5c,0xd1,0xf8,0x00,0x00,0x00] v_trunc_f32_e64 v0, 0x3e22f983 -// NOSICI: error: invalid literal operand -// NOGFX89: error: invalid literal operand +// NOSICI: error: literal operands are not supported +// NOGFX89: error: literal operands are not supported v_fract_f64_e64 v[0:1], 0x3e22f983 // NOSICI: error: invalid operand for instruction @@ -466,7 +466,7 @@ // GFX89: v_and_b32_e32 v0, 0.15915494, v1 ; encoding: [0xf8,0x02,0x00,0x26] v_and_b32_e32 v0, 0.159154943091895317852646485335, v1 -// NOSICI: error: invalid literal operand +// NOSICI: error: literal operands are not supported // GFX89: v_and_b32_e64 v0, 0.15915494, v1 ; encoding: [0x00,0x00,0x13,0xd1,0xf8,0x02,0x02,0x00] v_and_b32_e64 v0, 0.159154943091895317852646485335, v1 diff --git a/llvm/test/MC/AMDGPU/literalv216-err.s b/llvm/test/MC/AMDGPU/literalv216-err.s --- a/llvm/test/MC/AMDGPU/literalv216-err.s +++ b/llvm/test/MC/AMDGPU/literalv216-err.s @@ -2,45 +2,45 @@ // RUN: not llvm-mc -arch=amdgcn -mcpu=gfx1010 %s 2>&1 | FileCheck -check-prefix=GFX10 --implicit-check-not=error: %s v_pk_add_f16 v1, -17, v2 -// GFX9: error: invalid literal operand +// GFX9: error: literal operands are not supported v_pk_add_f16 v1, 65, v2 -// GFX9: error: invalid literal operand +// GFX9: error: literal operands are not supported v_pk_add_f16 v1, 64.0, v2 -// GFX9: error: invalid literal operand +// GFX9: error: literal operands are not supported v_pk_add_f16 v1, -0.15915494, v2 -// GFX9: error: invalid literal operand +// GFX9: error: literal operands are not supported v_pk_add_f16 v1, -0.0, v2 -// GFX9: error: invalid literal operand +// GFX9: error: literal operands are not supported v_pk_add_f16 v1, -32768, v2 -// GFX9: error: invalid literal operand +// GFX9: error: literal operands are not supported v_pk_add_f16 v1, 32767, v2 -// GFX9: error: invalid literal operand +// GFX9: error: literal operands are not supported v_pk_add_f16 v1, 0xffffffffffff000f, v2 -// GFX9: error: invalid literal operand +// GFX9: error: literal operands are not supported v_pk_add_f16 v1, 0x1000ffff, v2 -// GFX9: error: invalid literal operand +// GFX9: error: literal operands are not supported v_pk_mad_i16 v5, 0x3c00, 0x4000, 0x4400 -// GFX9: error: invalid literal operand -// GFX10: error: invalid literal operand +// GFX9: error: literal operands are not supported +// GFX10: error: only one literal operand is allowed v_pk_mad_i16 v5, 0x3c00, 0x4000, 2 -// GFX9: error: invalid literal operand -// GFX10: error: invalid literal operand +// GFX9: error: literal operands are not supported +// GFX10: error: only one literal operand is allowed v_pk_mad_i16 v5, 0x3c00, 3, 2 -// GFX9: error: invalid literal operand +// GFX9: error: literal operands are not supported v_pk_mad_i16 v5, 3, 0x3c00, 2 -// GFX9: error: invalid literal operand +// GFX9: error: literal operands are not supported v_pk_mad_i16 v5, 3, 2, 0x3c00 -// GFX9: error: invalid literal operand +// GFX9: error: literal operands are not supported diff --git a/llvm/test/MC/AMDGPU/literalv216.s b/llvm/test/MC/AMDGPU/literalv216.s --- a/llvm/test/MC/AMDGPU/literalv216.s +++ b/llvm/test/MC/AMDGPU/literalv216.s @@ -161,35 +161,35 @@ //===----------------------------------------------------------------------===// v_pk_add_f16 v5, v1, 0x12345678 -// NOGFX9: error: invalid literal operand +// NOGFX9: error: literal operands are not supported // GFX10: v_pk_add_f16 v5, v1, 0x12345678 ; encoding: [0x05,0x00,0x0f,0xcc,0x01,0xff,0x01,0x18,0x78,0x56,0x34,0x12] v_pk_add_f16 v5, 0x12345678, v2 -// NOGFX9: error: invalid literal operand +// NOGFX9: error: literal operands are not supported // GFX10: v_pk_add_f16 v5, 0x12345678, v2 ; encoding: [0x05,0x00,0x0f,0xcc,0xff,0x04,0x02,0x18,0x78,0x56,0x34,0x12] v_pk_add_f16 v5, -256, v2 -// NOGFX9: error: invalid literal operand +// NOGFX9: error: literal operands are not supported // GFX10: v_pk_add_f16 v5, 0xffffff00, v2 ; encoding: [0x05,0x00,0x0f,0xcc,0xff,0x04,0x02,0x18,0x00,0xff,0xff,0xff] v_pk_add_f16 v5, v1, 256 -// NOGFX9: error: invalid literal operand +// NOGFX9: error: literal operands are not supported // GFX10: v_pk_add_f16 v5, v1, 0x100 ; encoding: [0x05,0x00,0x0f,0xcc,0x01,0xff,0x01,0x18,0x00,0x01,0x00,0x00] v_pk_add_u16 v5, v1, 0x12345678 -// NOGFX9: error: invalid literal operand +// NOGFX9: error: literal operands are not supported // GFX10: v_pk_add_u16 v5, v1, 0x12345678 ; encoding: [0x05,0x00,0x0a,0xcc,0x01,0xff,0x01,0x18,0x78,0x56,0x34,0x12] v_pk_add_u16 v5, 0x12345678, v2 -// NOGFX9: error: invalid literal operand +// NOGFX9: error: literal operands are not supported // GFX10: v_pk_add_u16 v5, 0x12345678, v2 ; encoding: [0x05,0x00,0x0a,0xcc,0xff,0x04,0x02,0x18,0x78,0x56,0x34,0x12] v_pk_add_u16 v5, -256, v2 -// NOGFX9: error: invalid literal operand +// NOGFX9: error: literal operands are not supported // GFX10: v_pk_add_u16 v5, 0xffffff00, v2 ; encoding: [0x05,0x00,0x0a,0xcc,0xff,0x04,0x02,0x18,0x00,0xff,0xff,0xff] v_pk_add_u16 v5, v1, 256 -// NOGFX9: error: invalid literal operand +// NOGFX9: error: literal operands are not supported // GFX10: v_pk_add_u16 v5, v1, 0x100 ; encoding: [0x05,0x00,0x0a,0xcc,0x01,0xff,0x01,0x18,0x00,0x01,0x00,0x00] v_pk_add_f16 v5, v1, 0x123456780 @@ -201,35 +201,35 @@ // NOGFX10: error: invalid operand for instruction v_pk_fma_f16 v5, 0xaf123456, v2, v3 -// NOGFX9: error: invalid literal operand +// NOGFX9: error: literal operands are not supported // GFX10: v_pk_fma_f16 v5, 0xaf123456, v2, v3 ; encoding: [0x05,0x40,0x0e,0xcc,0xff,0x04,0x0e,0x1c,0x56,0x34,0x12,0xaf] v_pk_fma_f16 v5, v1, 0xaf123456, v3 -// NOGFX9: error: invalid literal operand +// NOGFX9: error: literal operands are not supported // GFX10: v_pk_fma_f16 v5, v1, 0xaf123456, v3 ; encoding: [0x05,0x40,0x0e,0xcc,0x01,0xff,0x0d,0x1c,0x56,0x34,0x12,0xaf] v_pk_fma_f16 v5, v1, v2, 0xaf123456 -// NOGFX9: error: invalid literal operand +// NOGFX9: error: literal operands are not supported // GFX10: v_pk_fma_f16 v5, v1, v2, 0xaf123456 ; encoding: [0x05,0x40,0x0e,0xcc,0x01,0x05,0xfe,0x1b,0x56,0x34,0x12,0xaf] v_pk_mad_i16 v5, 0xaf123456, v2, v3 -// NOGFX9: error: invalid literal operand +// NOGFX9: error: literal operands are not supported // GFX10: v_pk_mad_i16 v5, 0xaf123456, v2, v3 ; encoding: [0x05,0x40,0x00,0xcc,0xff,0x04,0x0e,0x1c,0x56,0x34,0x12,0xaf] v_pk_mad_i16 v5, v1, 0xaf123456, v3 -// NOGFX9: error: invalid literal operand +// NOGFX9: error: literal operands are not supported // GFX10: v_pk_mad_i16 v5, v1, 0xaf123456, v3 ; encoding: [0x05,0x40,0x00,0xcc,0x01,0xff,0x0d,0x1c,0x56,0x34,0x12,0xaf] v_pk_mad_i16 v5, v1, v2, 0xaf123456 -// NOGFX9: error: invalid literal operand +// NOGFX9: error: literal operands are not supported // GFX10: v_pk_mad_i16 v5, v1, v2, 0xaf123456 ; encoding: [0x05,0x40,0x00,0xcc,0x01,0x05,0xfe,0x1b,0x56,0x34,0x12,0xaf] v_pk_ashrrev_i16 v5, 0x12345678, v2 -// NOGFX9: error: invalid literal operand +// NOGFX9: error: literal operands are not supported // GFX10: v_pk_ashrrev_i16 v5, 0x12345678, v2 ; encoding: [0x05,0x00,0x06,0xcc,0xff,0x04,0x02,0x18,0x78,0x56,0x34,0x12] v_pk_ashrrev_i16 v5, v1, 0x12345678 -// NOGFX9: error: invalid literal operand +// NOGFX9: error: literal operands are not supported // GFX10: v_pk_ashrrev_i16 v5, v1, 0x12345678 ; encoding: [0x05,0x00,0x06,0xcc,0x01,0xff,0x01,0x18,0x78,0x56,0x34,0x12] //===----------------------------------------------------------------------===// @@ -237,35 +237,35 @@ //===----------------------------------------------------------------------===// v_pk_add_f16 v5, v1, 0.1234 -// NOGFX9: error: invalid literal operand +// NOGFX9: error: literal operands are not supported // GFX10: v_pk_add_f16 v5, v1, 0x2fe6 ; encoding: [0x05,0x00,0x0f,0xcc,0x01,0xff,0x01,0x18,0xe6,0x2f,0x00,0x00] v_pk_add_u16 v5, v1, 0.1234 -// NOGFX9: error: invalid literal operand +// NOGFX9: error: literal operands are not supported // GFX10: v_pk_add_u16 v5, v1, 0x2fe6 ; encoding: [0x05,0x00,0x0a,0xcc,0x01,0xff,0x01,0x18,0xe6,0x2f,0x00,0x00] v_pk_fma_f16 v5, 0.1234, v2, v3 -// NOGFX9: error: invalid literal operand +// NOGFX9: error: literal operands are not supported // GFX10: v_pk_fma_f16 v5, 0x2fe6, v2, v3 ; encoding: [0x05,0x40,0x0e,0xcc,0xff,0x04,0x0e,0x1c,0xe6,0x2f,0x00,0x00] v_pk_fma_f16 v5, v1, 0.1234, v3 -// NOGFX9: error: invalid literal operand +// NOGFX9: error: literal operands are not supported // GFX10: v_pk_fma_f16 v5, v1, 0x2fe6, v3 ; encoding: [0x05,0x40,0x0e,0xcc,0x01,0xff,0x0d,0x1c,0xe6,0x2f,0x00,0x00] v_pk_fma_f16 v5, v1, v2, 0.1234 -// NOGFX9: error: invalid literal operand +// NOGFX9: error: literal operands are not supported // GFX10: v_pk_fma_f16 v5, v1, v2, 0x2fe6 ; encoding: [0x05,0x40,0x0e,0xcc,0x01,0x05,0xfe,0x1b,0xe6,0x2f,0x00,0x00] v_pk_mad_i16 v5, 0.1234, v2, v3 -// NOGFX9: error: invalid literal operand +// NOGFX9: error: literal operands are not supported // GFX10: v_pk_mad_i16 v5, 0x2fe6, v2, v3 ; encoding: [0x05,0x40,0x00,0xcc,0xff,0x04,0x0e,0x1c,0xe6,0x2f,0x00,0x00] v_pk_mad_i16 v5, v1, 0.1234, v3 -// NOGFX9: error: invalid literal operand +// NOGFX9: error: literal operands are not supported // GFX10: v_pk_mad_i16 v5, v1, 0x2fe6, v3 ; encoding: [0x05,0x40,0x00,0xcc,0x01,0xff,0x0d,0x1c,0xe6,0x2f,0x00,0x00] v_pk_mad_i16 v5, v1, v2, 0.1234 -// NOGFX9: error: invalid literal operand +// NOGFX9: error: literal operands are not supported // GFX10: v_pk_mad_i16 v5, v1, v2, 0x2fe6 ; encoding: [0x05,0x40,0x00,0xcc,0x01,0x05,0xfe,0x1b,0xe6,0x2f,0x00,0x00] v_pk_add_f16 v5, v1, 123456.0 diff --git a/llvm/test/MC/AMDGPU/mai-err.s b/llvm/test/MC/AMDGPU/mai-err.s --- a/llvm/test/MC/AMDGPU/mai-err.s +++ b/llvm/test/MC/AMDGPU/mai-err.s @@ -61,640 +61,640 @@ v_mfma_f32_32x32x1f32 a[0:31], v0, v1, 0 // GFX900: error: instruction not supported on this GPU -// GFX908: error: invalid literal operand +// GFX908: error: inline constants are not allowed for this operand v_mfma_f32_32x32x1f32 a[0:31], v0, v1, -2.0 -// GFX908: error: invalid literal operand +// GFX908: error: inline constants are not allowed for this operand // GFX900: error: instruction not supported on this GPU v_mfma_f32_32x32x1f32 a[0:31], v0, v1, -2.0 cbsz:3 abid:2 blgp:7 -// GFX908: error: invalid literal operand +// GFX908: error: inline constants are not allowed for this operand // GFX900: error: instruction not supported on this GPU v_mfma_f32_32x32x1f32 a[0:31], v0, a1, -2.0 -// GFX908: error: invalid literal operand +// GFX908: error: inline constants are not allowed for this operand // GFX900: error: instruction not supported on this GPU v_mfma_f32_32x32x1f32 a[0:31], v0, a1, -2.0 cbsz:3 abid:2 blgp:7 -// GFX908: error: invalid literal operand +// GFX908: error: inline constants are not allowed for this operand // GFX900: error: instruction not supported on this GPU v_mfma_f32_32x32x1f32 a[0:31], a0, v1, -2.0 -// GFX908: error: invalid literal operand +// GFX908: error: inline constants are not allowed for this operand // GFX900: error: instruction not supported on this GPU v_mfma_f32_32x32x1f32 a[0:31], a0, v1, -2.0 cbsz:3 abid:2 blgp:7 -// GFX908: error: invalid literal operand +// GFX908: error: inline constants are not allowed for this operand // GFX900: error: instruction not supported on this GPU v_mfma_f32_32x32x1f32 a[0:31], a0, a1, -2.0 -// GFX908: error: invalid literal operand +// GFX908: error: inline constants are not allowed for this operand // GFX900: error: instruction not supported on this GPU v_mfma_f32_32x32x1f32 a[0:31], a0, a1, -2.0 cbsz:3 abid:2 blgp:7 -// GFX908: error: invalid literal operand +// GFX908: error: inline constants are not allowed for this operand // GFX900: error: instruction not supported on this GPU v_mfma_f32_16x16x1f32 a[0:15], v0, v1, -2.0 -// GFX908: error: invalid literal operand +// GFX908: error: inline constants are not allowed for this operand // GFX900: error: instruction not supported on this GPU v_mfma_f32_16x16x1f32 a[0:15], v0, v1, -2.0 cbsz:3 abid:2 blgp:7 -// GFX908: error: invalid literal operand +// GFX908: error: inline constants are not allowed for this operand // GFX900: error: instruction not supported on this GPU v_mfma_f32_16x16x1f32 a[0:15], v0, a1, -2.0 -// GFX908: error: invalid literal operand +// GFX908: error: inline constants are not allowed for this operand // GFX900: error: instruction not supported on this GPU v_mfma_f32_16x16x1f32 a[0:15], v0, a1, -2.0 cbsz:3 abid:2 blgp:7 -// GFX908: error: invalid literal operand +// GFX908: error: inline constants are not allowed for this operand // GFX900: error: instruction not supported on this GPU v_mfma_f32_16x16x1f32 a[0:15], a0, v1, -2.0 -// GFX908: error: invalid literal operand +// GFX908: error: inline constants are not allowed for this operand // GFX900: error: instruction not supported on this GPU v_mfma_f32_16x16x1f32 a[0:15], a0, v1, -2.0 cbsz:3 abid:2 blgp:7 -// GFX908: error: invalid literal operand +// GFX908: error: inline constants are not allowed for this operand // GFX900: error: instruction not supported on this GPU v_mfma_f32_16x16x1f32 a[0:15], a0, a1, -2.0 -// GFX908: error: invalid literal operand +// GFX908: error: inline constants are not allowed for this operand // GFX900: error: instruction not supported on this GPU v_mfma_f32_16x16x1f32 a[0:15], a0, a1, -2.0 cbsz:3 abid:2 blgp:7 -// GFX908: error: invalid literal operand +// GFX908: error: inline constants are not allowed for this operand // GFX900: error: instruction not supported on this GPU v_mfma_f32_4x4x1f32 a[0:3], v0, v1, -2.0 -// GFX908: error: invalid literal operand +// GFX908: error: inline constants are not allowed for this operand // GFX900: error: instruction not supported on this GPU v_mfma_f32_4x4x1f32 a[0:3], v0, v1, -2.0 cbsz:3 abid:2 blgp:7 -// GFX908: error: invalid literal operand +// GFX908: error: inline constants are not allowed for this operand // GFX900: error: instruction not supported on this GPU v_mfma_f32_4x4x1f32 a[0:3], v0, a1, -2.0 -// GFX908: error: invalid literal operand +// GFX908: error: inline constants are not allowed for this operand // GFX900: error: instruction not supported on this GPU v_mfma_f32_4x4x1f32 a[0:3], v0, a1, -2.0 cbsz:3 abid:2 blgp:7 -// GFX908: error: invalid literal operand +// GFX908: error: inline constants are not allowed for this operand // GFX900: error: instruction not supported on this GPU v_mfma_f32_4x4x1f32 a[0:3], a0, v1, -2.0 -// GFX908: error: invalid literal operand +// GFX908: error: inline constants are not allowed for this operand // GFX900: error: instruction not supported on this GPU v_mfma_f32_4x4x1f32 a[0:3], a0, v1, -2.0 cbsz:3 abid:2 blgp:7 -// GFX908: error: invalid literal operand +// GFX908: error: inline constants are not allowed for this operand // GFX900: error: instruction not supported on this GPU v_mfma_f32_4x4x1f32 a[0:3], a0, a1, -2.0 -// GFX908: error: invalid literal operand +// GFX908: error: inline constants are not allowed for this operand // GFX900: error: instruction not supported on this GPU v_mfma_f32_4x4x1f32 a[0:3], a0, a1, -2.0 cbsz:3 abid:2 blgp:7 -// GFX908: error: invalid literal operand +// GFX908: error: inline constants are not allowed for this operand // GFX900: error: instruction not supported on this GPU v_mfma_f32_32x32x2f32 a[0:15], v0, v1, -2.0 -// GFX908: error: invalid literal operand +// GFX908: error: inline constants are not allowed for this operand // GFX900: error: instruction not supported on this GPU v_mfma_f32_32x32x2f32 a[0:15], v0, v1, -2.0 cbsz:3 abid:2 blgp:7 -// GFX908: error: invalid literal operand +// GFX908: error: inline constants are not allowed for this operand // GFX900: error: instruction not supported on this GPU v_mfma_f32_32x32x2f32 a[0:15], v0, a1, -2.0 -// GFX908: error: invalid literal operand +// GFX908: error: inline constants are not allowed for this operand // GFX900: error: instruction not supported on this GPU v_mfma_f32_32x32x2f32 a[0:15], v0, a1, -2.0 cbsz:3 abid:2 blgp:7 -// GFX908: error: invalid literal operand +// GFX908: error: inline constants are not allowed for this operand // GFX900: error: instruction not supported on this GPU v_mfma_f32_32x32x2f32 a[0:15], a0, v1, -2.0 -// GFX908: error: invalid literal operand +// GFX908: error: inline constants are not allowed for this operand // GFX900: error: instruction not supported on this GPU v_mfma_f32_32x32x2f32 a[0:15], a0, v1, -2.0 cbsz:3 abid:2 blgp:7 -// GFX908: error: invalid literal operand +// GFX908: error: inline constants are not allowed for this operand // GFX900: error: instruction not supported on this GPU v_mfma_f32_32x32x2f32 a[0:15], a0, a1, -2.0 -// GFX908: error: invalid literal operand +// GFX908: error: inline constants are not allowed for this operand // GFX900: error: instruction not supported on this GPU v_mfma_f32_32x32x2f32 a[0:15], a0, a1, -2.0 cbsz:3 abid:2 blgp:7 -// GFX908: error: invalid literal operand +// GFX908: error: inline constants are not allowed for this operand // GFX900: error: instruction not supported on this GPU v_mfma_f32_16x16x4f32 a[0:3], v0, v1, -2.0 -// GFX908: error: invalid literal operand +// GFX908: error: inline constants are not allowed for this operand // GFX900: error: instruction not supported on this GPU v_mfma_f32_16x16x4f32 a[0:3], v0, v1, -2.0 cbsz:3 abid:2 blgp:7 -// GFX908: error: invalid literal operand +// GFX908: error: inline constants are not allowed for this operand // GFX900: error: instruction not supported on this GPU v_mfma_f32_16x16x4f32 a[0:3], v0, a1, -2.0 -// GFX908: error: invalid literal operand +// GFX908: error: inline constants are not allowed for this operand // GFX900: error: instruction not supported on this GPU v_mfma_f32_16x16x4f32 a[0:3], v0, a1, -2.0 cbsz:3 abid:2 blgp:7 -// GFX908: error: invalid literal operand +// GFX908: error: inline constants are not allowed for this operand // GFX900: error: instruction not supported on this GPU v_mfma_f32_16x16x4f32 a[0:3], a0, v1, -2.0 -// GFX908: error: invalid literal operand +// GFX908: error: inline constants are not allowed for this operand // GFX900: error: instruction not supported on this GPU v_mfma_f32_16x16x4f32 a[0:3], a0, v1, -2.0 cbsz:3 abid:2 blgp:7 -// GFX908: error: invalid literal operand +// GFX908: error: inline constants are not allowed for this operand // GFX900: error: instruction not supported on this GPU v_mfma_f32_16x16x4f32 a[0:3], a0, a1, -2.0 -// GFX908: error: invalid literal operand +// GFX908: error: inline constants are not allowed for this operand // GFX900: error: instruction not supported on this GPU v_mfma_f32_16x16x4f32 a[0:3], a0, a1, -2.0 cbsz:3 abid:2 blgp:7 -// GFX908: error: invalid literal operand +// GFX908: error: inline constants are not allowed for this operand // GFX900: error: instruction not supported on this GPU v_mfma_f32_32x32x4f16 a[0:31], v[0:1], v[1:2], -2.0 -// GFX908: error: invalid literal operand +// GFX908: error: inline constants are not allowed for this operand // GFX900: error: instruction not supported on this GPU v_mfma_f32_32x32x4f16 a[0:31], v[0:1], v[1:2], -2.0 cbsz:3 abid:2 blgp:7 -// GFX908: error: invalid literal operand +// GFX908: error: inline constants are not allowed for this operand // GFX900: error: instruction not supported on this GPU v_mfma_f32_32x32x4f16 a[0:31], v[0:1], a[1:2], -2.0 -// GFX908: error: invalid literal operand +// GFX908: error: inline constants are not allowed for this operand // GFX900: error: instruction not supported on this GPU v_mfma_f32_32x32x4f16 a[0:31], v[0:1], a[1:2], -2.0 cbsz:3 abid:2 blgp:7 -// GFX908: error: invalid literal operand +// GFX908: error: inline constants are not allowed for this operand // GFX900: error: instruction not supported on this GPU v_mfma_f32_32x32x4f16 a[0:31], a[0:1], v[1:2], -2.0 -// GFX908: error: invalid literal operand +// GFX908: error: inline constants are not allowed for this operand // GFX900: error: instruction not supported on this GPU v_mfma_f32_32x32x4f16 a[0:31], a[0:1], v[1:2], -2.0 cbsz:3 abid:2 blgp:7 -// GFX908: error: invalid literal operand +// GFX908: error: inline constants are not allowed for this operand // GFX900: error: instruction not supported on this GPU v_mfma_f32_32x32x4f16 a[0:31], a[0:1], a[1:2], -2.0 -// GFX908: error: invalid literal operand +// GFX908: error: inline constants are not allowed for this operand // GFX900: error: instruction not supported on this GPU v_mfma_f32_32x32x4f16 a[0:31], a[0:1], a[1:2], -2.0 cbsz:3 abid:2 blgp:7 -// GFX908: error: invalid literal operand +// GFX908: error: inline constants are not allowed for this operand // GFX900: error: instruction not supported on this GPU v_mfma_f32_16x16x4f16 a[0:15], v[0:1], v[1:2], -2.0 -// GFX908: error: invalid literal operand +// GFX908: error: inline constants are not allowed for this operand // GFX900: error: instruction not supported on this GPU v_mfma_f32_16x16x4f16 a[0:15], v[0:1], v[1:2], -2.0 cbsz:3 abid:2 blgp:7 -// GFX908: error: invalid literal operand +// GFX908: error: inline constants are not allowed for this operand // GFX900: error: instruction not supported on this GPU v_mfma_f32_16x16x4f16 a[0:15], v[0:1], a[1:2], -2.0 -// GFX908: error: invalid literal operand +// GFX908: error: inline constants are not allowed for this operand // GFX900: error: instruction not supported on this GPU v_mfma_f32_16x16x4f16 a[0:15], v[0:1], a[1:2], -2.0 cbsz:3 abid:2 blgp:7 -// GFX908: error: invalid literal operand +// GFX908: error: inline constants are not allowed for this operand // GFX900: error: instruction not supported on this GPU v_mfma_f32_16x16x4f16 a[0:15], a[0:1], v[1:2], -2.0 -// GFX908: error: invalid literal operand +// GFX908: error: inline constants are not allowed for this operand // GFX900: error: instruction not supported on this GPU v_mfma_f32_16x16x4f16 a[0:15], a[0:1], v[1:2], -2.0 cbsz:3 abid:2 blgp:7 -// GFX908: error: invalid literal operand +// GFX908: error: inline constants are not allowed for this operand // GFX900: error: instruction not supported on this GPU v_mfma_f32_16x16x4f16 a[0:15], a[0:1], a[1:2], -2.0 -// GFX908: error: invalid literal operand +// GFX908: error: inline constants are not allowed for this operand // GFX900: error: instruction not supported on this GPU v_mfma_f32_16x16x4f16 a[0:15], a[0:1], a[1:2], -2.0 cbsz:3 abid:2 blgp:7 -// GFX908: error: invalid literal operand +// GFX908: error: inline constants are not allowed for this operand // GFX900: error: instruction not supported on this GPU v_mfma_f32_4x4x4f16 a[0:3], v[0:1], v[1:2], -2.0 -// GFX908: error: invalid literal operand +// GFX908: error: inline constants are not allowed for this operand // GFX900: error: instruction not supported on this GPU v_mfma_f32_4x4x4f16 a[0:3], v[0:1], v[1:2], -2.0 cbsz:3 abid:2 blgp:7 -// GFX908: error: invalid literal operand +// GFX908: error: inline constants are not allowed for this operand // GFX900: error: instruction not supported on this GPU v_mfma_f32_4x4x4f16 a[0:3], v[0:1], a[1:2], -2.0 -// GFX908: error: invalid literal operand +// GFX908: error: inline constants are not allowed for this operand // GFX900: error: instruction not supported on this GPU v_mfma_f32_4x4x4f16 a[0:3], v[0:1], a[1:2], -2.0 cbsz:3 abid:2 blgp:7 -// GFX908: error: invalid literal operand +// GFX908: error: inline constants are not allowed for this operand // GFX900: error: instruction not supported on this GPU v_mfma_f32_4x4x4f16 a[0:3], a[0:1], v[1:2], -2.0 -// GFX908: error: invalid literal operand +// GFX908: error: inline constants are not allowed for this operand // GFX900: error: instruction not supported on this GPU v_mfma_f32_4x4x4f16 a[0:3], a[0:1], v[1:2], -2.0 cbsz:3 abid:2 blgp:7 -// GFX908: error: invalid literal operand +// GFX908: error: inline constants are not allowed for this operand // GFX900: error: instruction not supported on this GPU v_mfma_f32_4x4x4f16 a[0:3], a[0:1], a[1:2], -2.0 -// GFX908: error: invalid literal operand +// GFX908: error: inline constants are not allowed for this operand // GFX900: error: instruction not supported on this GPU v_mfma_f32_4x4x4f16 a[0:3], a[0:1], a[1:2], -2.0 cbsz:3 abid:2 blgp:7 -// GFX908: error: invalid literal operand +// GFX908: error: inline constants are not allowed for this operand // GFX900: error: instruction not supported on this GPU v_mfma_f32_32x32x8f16 a[0:15], v[0:1], v[1:2], -2.0 -// GFX908: error: invalid literal operand +// GFX908: error: inline constants are not allowed for this operand // GFX900: error: instruction not supported on this GPU v_mfma_f32_32x32x8f16 a[0:15], v[0:1], v[1:2], -2.0 cbsz:3 abid:2 blgp:7 -// GFX908: error: invalid literal operand +// GFX908: error: inline constants are not allowed for this operand // GFX900: error: instruction not supported on this GPU v_mfma_f32_32x32x8f16 a[0:15], v[0:1], a[1:2], -2.0 -// GFX908: error: invalid literal operand +// GFX908: error: inline constants are not allowed for this operand // GFX900: error: instruction not supported on this GPU v_mfma_f32_32x32x8f16 a[0:15], v[0:1], a[1:2], -2.0 cbsz:3 abid:2 blgp:7 -// GFX908: error: invalid literal operand +// GFX908: error: inline constants are not allowed for this operand // GFX900: error: instruction not supported on this GPU v_mfma_f32_32x32x8f16 a[0:15], a[0:1], v[1:2], -2.0 -// GFX908: error: invalid literal operand +// GFX908: error: inline constants are not allowed for this operand // GFX900: error: instruction not supported on this GPU v_mfma_f32_32x32x8f16 a[0:15], a[0:1], v[1:2], -2.0 cbsz:3 abid:2 blgp:7 -// GFX908: error: invalid literal operand +// GFX908: error: inline constants are not allowed for this operand // GFX900: error: instruction not supported on this GPU v_mfma_f32_32x32x8f16 a[0:15], a[0:1], a[1:2], -2.0 -// GFX908: error: invalid literal operand +// GFX908: error: inline constants are not allowed for this operand // GFX900: error: instruction not supported on this GPU v_mfma_f32_32x32x8f16 a[0:15], a[0:1], a[1:2], -2.0 cbsz:3 abid:2 blgp:7 -// GFX908: error: invalid literal operand +// GFX908: error: inline constants are not allowed for this operand // GFX900: error: instruction not supported on this GPU v_mfma_f32_16x16x16f16 a[0:3], v[0:1], v[1:2], -2.0 -// GFX908: error: invalid literal operand +// GFX908: error: inline constants are not allowed for this operand // GFX900: error: instruction not supported on this GPU v_mfma_f32_16x16x16f16 a[0:3], v[0:1], v[1:2], -2.0 cbsz:3 abid:2 blgp:7 -// GFX908: error: invalid literal operand +// GFX908: error: inline constants are not allowed for this operand // GFX900: error: instruction not supported on this GPU v_mfma_f32_16x16x16f16 a[0:3], v[0:1], a[1:2], -2.0 -// GFX908: error: invalid literal operand +// GFX908: error: inline constants are not allowed for this operand // GFX900: error: instruction not supported on this GPU v_mfma_f32_16x16x16f16 a[0:3], v[0:1], a[1:2], -2.0 cbsz:3 abid:2 blgp:7 -// GFX908: error: invalid literal operand +// GFX908: error: inline constants are not allowed for this operand // GFX900: error: instruction not supported on this GPU v_mfma_f32_16x16x16f16 a[0:3], a[0:1], v[1:2], -2.0 -// GFX908: error: invalid literal operand +// GFX908: error: inline constants are not allowed for this operand // GFX900: error: instruction not supported on this GPU v_mfma_f32_16x16x16f16 a[0:3], a[0:1], v[1:2], -2.0 cbsz:3 abid:2 blgp:7 -// GFX908: error: invalid literal operand +// GFX908: error: inline constants are not allowed for this operand // GFX900: error: instruction not supported on this GPU v_mfma_f32_16x16x16f16 a[0:3], a[0:1], a[1:2], -2.0 -// GFX908: error: invalid literal operand +// GFX908: error: inline constants are not allowed for this operand // GFX900: error: instruction not supported on this GPU v_mfma_f32_16x16x16f16 a[0:3], a[0:1], a[1:2], -2.0 cbsz:3 abid:2 blgp:7 -// GFX908: error: invalid literal operand +// GFX908: error: inline constants are not allowed for this operand // GFX900: error: instruction not supported on this GPU v_mfma_i32_32x32x4i8 a[0:31], v0, v1, 2 -// GFX908: error: invalid literal operand +// GFX908: error: inline constants are not allowed for this operand // GFX900: error: instruction not supported on this GPU v_mfma_i32_32x32x4i8 a[0:31], v0, v1, 2 cbsz:3 abid:2 blgp:7 -// GFX908: error: invalid literal operand +// GFX908: error: inline constants are not allowed for this operand // GFX900: error: instruction not supported on this GPU v_mfma_i32_32x32x4i8 a[0:31], v0, a1, 2 -// GFX908: error: invalid literal operand +// GFX908: error: inline constants are not allowed for this operand // GFX900: error: instruction not supported on this GPU v_mfma_i32_32x32x4i8 a[0:31], v0, a1, 2 cbsz:3 abid:2 blgp:7 -// GFX908: error: invalid literal operand +// GFX908: error: inline constants are not allowed for this operand // GFX900: error: instruction not supported on this GPU v_mfma_i32_32x32x4i8 a[0:31], a0, v1, 2 -// GFX908: error: invalid literal operand +// GFX908: error: inline constants are not allowed for this operand // GFX900: error: instruction not supported on this GPU v_mfma_i32_32x32x4i8 a[0:31], a0, v1, 2 cbsz:3 abid:2 blgp:7 -// GFX908: error: invalid literal operand +// GFX908: error: inline constants are not allowed for this operand // GFX900: error: instruction not supported on this GPU v_mfma_i32_32x32x4i8 a[0:31], a0, a1, 2 -// GFX908: error: invalid literal operand +// GFX908: error: inline constants are not allowed for this operand // GFX900: error: instruction not supported on this GPU v_mfma_i32_32x32x4i8 a[0:31], a0, a1, 2 cbsz:3 abid:2 blgp:7 -// GFX908: error: invalid literal operand +// GFX908: error: inline constants are not allowed for this operand // GFX900: error: instruction not supported on this GPU v_mfma_i32_16x16x4i8 a[0:15], v0, v1, 2 -// GFX908: error: invalid literal operand +// GFX908: error: inline constants are not allowed for this operand // GFX900: error: instruction not supported on this GPU v_mfma_i32_16x16x4i8 a[0:15], v0, v1, 2 cbsz:3 abid:2 blgp:7 -// GFX908: error: invalid literal operand +// GFX908: error: inline constants are not allowed for this operand // GFX900: error: instruction not supported on this GPU v_mfma_i32_16x16x4i8 a[0:15], v0, a1, 2 -// GFX908: error: invalid literal operand +// GFX908: error: inline constants are not allowed for this operand // GFX900: error: instruction not supported on this GPU v_mfma_i32_16x16x4i8 a[0:15], v0, a1, 2 cbsz:3 abid:2 blgp:7 -// GFX908: error: invalid literal operand +// GFX908: error: inline constants are not allowed for this operand // GFX900: error: instruction not supported on this GPU v_mfma_i32_16x16x4i8 a[0:15], a0, v1, 2 -// GFX908: error: invalid literal operand +// GFX908: error: inline constants are not allowed for this operand // GFX900: error: instruction not supported on this GPU v_mfma_i32_16x16x4i8 a[0:15], a0, v1, 2 cbsz:3 abid:2 blgp:7 -// GFX908: error: invalid literal operand +// GFX908: error: inline constants are not allowed for this operand // GFX900: error: instruction not supported on this GPU v_mfma_i32_16x16x4i8 a[0:15], a0, a1, 2 -// GFX908: error: invalid literal operand +// GFX908: error: inline constants are not allowed for this operand // GFX900: error: instruction not supported on this GPU v_mfma_i32_16x16x4i8 a[0:15], a0, a1, 2 cbsz:3 abid:2 blgp:7 -// GFX908: error: invalid literal operand +// GFX908: error: inline constants are not allowed for this operand // GFX900: error: instruction not supported on this GPU v_mfma_i32_4x4x4i8 a[0:3], v0, v1, 2 -// GFX908: error: invalid literal operand +// GFX908: error: inline constants are not allowed for this operand // GFX900: error: instruction not supported on this GPU v_mfma_i32_4x4x4i8 a[0:3], v0, v1, 2 cbsz:3 abid:2 blgp:7 -// GFX908: error: invalid literal operand +// GFX908: error: inline constants are not allowed for this operand // GFX900: error: instruction not supported on this GPU v_mfma_i32_4x4x4i8 a[0:3], v0, a1, 2 -// GFX908: error: invalid literal operand +// GFX908: error: inline constants are not allowed for this operand // GFX900: error: instruction not supported on this GPU v_mfma_i32_4x4x4i8 a[0:3], v0, a1, 2 cbsz:3 abid:2 blgp:7 -// GFX908: error: invalid literal operand +// GFX908: error: inline constants are not allowed for this operand // GFX900: error: instruction not supported on this GPU v_mfma_i32_4x4x4i8 a[0:3], a0, v1, 2 -// GFX908: error: invalid literal operand +// GFX908: error: inline constants are not allowed for this operand // GFX900: error: instruction not supported on this GPU v_mfma_i32_4x4x4i8 a[0:3], a0, v1, 2 cbsz:3 abid:2 blgp:7 -// GFX908: error: invalid literal operand +// GFX908: error: inline constants are not allowed for this operand // GFX900: error: instruction not supported on this GPU v_mfma_i32_4x4x4i8 a[0:3], a0, a1, 2 -// GFX908: error: invalid literal operand +// GFX908: error: inline constants are not allowed for this operand // GFX900: error: instruction not supported on this GPU v_mfma_i32_4x4x4i8 a[0:3], a0, a1, 2 cbsz:3 abid:2 blgp:7 -// GFX908: error: invalid literal operand +// GFX908: error: inline constants are not allowed for this operand // GFX900: error: instruction not supported on this GPU v_mfma_i32_32x32x8i8 a[0:15], v0, v1, 2 -// GFX908: error: invalid literal operand +// GFX908: error: inline constants are not allowed for this operand // GFX900: error: instruction not supported on this GPU v_mfma_i32_32x32x8i8 a[0:15], v0, v1, 2 cbsz:3 abid:2 blgp:7 -// GFX908: error: invalid literal operand +// GFX908: error: inline constants are not allowed for this operand // GFX900: error: instruction not supported on this GPU v_mfma_i32_32x32x8i8 a[0:15], v0, a1, 2 -// GFX908: error: invalid literal operand +// GFX908: error: inline constants are not allowed for this operand // GFX900: error: instruction not supported on this GPU v_mfma_i32_32x32x8i8 a[0:15], v0, a1, 2 cbsz:3 abid:2 blgp:7 -// GFX908: error: invalid literal operand +// GFX908: error: inline constants are not allowed for this operand // GFX900: error: instruction not supported on this GPU v_mfma_i32_32x32x8i8 a[0:15], a0, v1, 2 -// GFX908: error: invalid literal operand +// GFX908: error: inline constants are not allowed for this operand // GFX900: error: instruction not supported on this GPU v_mfma_i32_32x32x8i8 a[0:15], a0, v1, 2 cbsz:3 abid:2 blgp:7 -// GFX908: error: invalid literal operand +// GFX908: error: inline constants are not allowed for this operand // GFX900: error: instruction not supported on this GPU v_mfma_i32_32x32x8i8 a[0:15], a0, a1, 2 -// GFX908: error: invalid literal operand +// GFX908: error: inline constants are not allowed for this operand // GFX900: error: instruction not supported on this GPU v_mfma_i32_32x32x8i8 a[0:15], a0, a1, 2 cbsz:3 abid:2 blgp:7 -// GFX908: error: invalid literal operand +// GFX908: error: inline constants are not allowed for this operand // GFX900: error: instruction not supported on this GPU v_mfma_i32_16x16x16i8 a[0:3], v0, v1, 2 -// GFX908: error: invalid literal operand +// GFX908: error: inline constants are not allowed for this operand // GFX900: error: instruction not supported on this GPU v_mfma_i32_16x16x16i8 a[0:3], v0, v1, 2 cbsz:3 abid:2 blgp:7 -// GFX908: error: invalid literal operand +// GFX908: error: inline constants are not allowed for this operand // GFX900: error: instruction not supported on this GPU v_mfma_i32_16x16x16i8 a[0:3], v0, a1, 2 -// GFX908: error: invalid literal operand +// GFX908: error: inline constants are not allowed for this operand // GFX900: error: instruction not supported on this GPU v_mfma_i32_16x16x16i8 a[0:3], v0, a1, 2 cbsz:3 abid:2 blgp:7 -// GFX908: error: invalid literal operand +// GFX908: error: inline constants are not allowed for this operand // GFX900: error: instruction not supported on this GPU v_mfma_i32_16x16x16i8 a[0:3], a0, v1, 2 -// GFX908: error: invalid literal operand +// GFX908: error: inline constants are not allowed for this operand // GFX900: error: instruction not supported on this GPU v_mfma_i32_16x16x16i8 a[0:3], a0, v1, 2 cbsz:3 abid:2 blgp:7 -// GFX908: error: invalid literal operand +// GFX908: error: inline constants are not allowed for this operand // GFX900: error: instruction not supported on this GPU v_mfma_i32_16x16x16i8 a[0:3], a0, a1, 2 -// GFX908: error: invalid literal operand +// GFX908: error: inline constants are not allowed for this operand // GFX900: error: instruction not supported on this GPU v_mfma_f32_32x32x2bf16 a[0:31], v0, v1, -2.0 -// GFX908: error: invalid literal operand +// GFX908: error: inline constants are not allowed for this operand // GFX900: error: instruction not supported on this GPU v_mfma_f32_32x32x2bf16 a[0:31], v0, v1, -2.0 cbsz:3 abid:2 blgp:7 -// GFX908: error: invalid literal operand +// GFX908: error: inline constants are not allowed for this operand // GFX900: error: instruction not supported on this GPU v_mfma_f32_32x32x2bf16 a[0:31], v0, a1, -2.0 -// GFX908: error: invalid literal operand +// GFX908: error: inline constants are not allowed for this operand // GFX900: error: instruction not supported on this GPU v_mfma_f32_32x32x2bf16 a[0:31], v0, a1, -2.0 cbsz:3 abid:2 blgp:7 -// GFX908: error: invalid literal operand +// GFX908: error: inline constants are not allowed for this operand // GFX900: error: instruction not supported on this GPU v_mfma_f32_32x32x2bf16 a[0:31], a0, v1, -2.0 -// GFX908: error: invalid literal operand +// GFX908: error: inline constants are not allowed for this operand // GFX900: error: instruction not supported on this GPU v_mfma_f32_32x32x2bf16 a[0:31], a0, v1, -2.0 cbsz:3 abid:2 blgp:7 -// GFX908: error: invalid literal operand +// GFX908: error: inline constants are not allowed for this operand // GFX900: error: instruction not supported on this GPU v_mfma_f32_32x32x2bf16 a[0:31], a0, a1, -2.0 -// GFX908: error: invalid literal operand +// GFX908: error: inline constants are not allowed for this operand // GFX900: error: instruction not supported on this GPU v_mfma_f32_32x32x2bf16 a[0:31], a0, a1, -2.0 cbsz:3 abid:2 blgp:7 -// GFX908: error: invalid literal operand +// GFX908: error: inline constants are not allowed for this operand // GFX900: error: instruction not supported on this GPU v_mfma_f32_16x16x2bf16 a[0:15], v0, v1, -2.0 -// GFX908: error: invalid literal operand +// GFX908: error: inline constants are not allowed for this operand // GFX900: error: instruction not supported on this GPU v_mfma_f32_16x16x2bf16 a[0:15], v0, v1, -2.0 cbsz:3 abid:2 blgp:7 -// GFX908: error: invalid literal operand +// GFX908: error: inline constants are not allowed for this operand // GFX900: error: instruction not supported on this GPU v_mfma_f32_16x16x2bf16 a[0:15], v0, a1, -2.0 -// GFX908: error: invalid literal operand +// GFX908: error: inline constants are not allowed for this operand // GFX900: error: instruction not supported on this GPU v_mfma_f32_16x16x2bf16 a[0:15], v0, a1, -2.0 cbsz:3 abid:2 blgp:7 -// GFX908: error: invalid literal operand +// GFX908: error: inline constants are not allowed for this operand // GFX900: error: instruction not supported on this GPU v_mfma_f32_16x16x2bf16 a[0:15], a0, v1, -2.0 -// GFX908: error: invalid literal operand +// GFX908: error: inline constants are not allowed for this operand // GFX900: error: instruction not supported on this GPU v_mfma_f32_16x16x2bf16 a[0:15], a0, v1, -2.0 cbsz:3 abid:2 blgp:7 -// GFX908: error: invalid literal operand +// GFX908: error: inline constants are not allowed for this operand // GFX900: error: instruction not supported on this GPU v_mfma_f32_16x16x2bf16 a[0:15], a0, a1, -2.0 -// GFX908: error: invalid literal operand +// GFX908: error: inline constants are not allowed for this operand // GFX900: error: instruction not supported on this GPU v_mfma_f32_16x16x2bf16 a[0:15], a0, a1, -2.0 cbsz:3 abid:2 blgp:7 -// GFX908: error: invalid literal operand +// GFX908: error: inline constants are not allowed for this operand // GFX900: error: instruction not supported on this GPU v_mfma_f32_4x4x2bf16 a[0:3], v0, v1, -2.0 -// GFX908: error: invalid literal operand +// GFX908: error: inline constants are not allowed for this operand // GFX900: error: instruction not supported on this GPU v_mfma_f32_4x4x2bf16 a[0:3], v0, v1, -2.0 cbsz:3 abid:2 blgp:7 -// GFX908: error: invalid literal operand +// GFX908: error: inline constants are not allowed for this operand // GFX900: error: instruction not supported on this GPU v_mfma_f32_4x4x2bf16 a[0:3], v0, a1, -2.0 -// GFX908: error: invalid literal operand +// GFX908: error: inline constants are not allowed for this operand // GFX900: error: instruction not supported on this GPU v_mfma_f32_4x4x2bf16 a[0:3], v0, a1, -2.0 cbsz:3 abid:2 blgp:7 -// GFX908: error: invalid literal operand +// GFX908: error: inline constants are not allowed for this operand // GFX900: error: instruction not supported on this GPU v_mfma_f32_4x4x2bf16 a[0:3], a0, v1, -2.0 -// GFX908: error: invalid literal operand +// GFX908: error: inline constants are not allowed for this operand // GFX900: error: instruction not supported on this GPU v_mfma_f32_4x4x2bf16 a[0:3], a0, v1, -2.0 cbsz:3 abid:2 blgp:7 -// GFX908: error: invalid literal operand +// GFX908: error: inline constants are not allowed for this operand // GFX900: error: instruction not supported on this GPU v_mfma_f32_4x4x2bf16 a[0:3], a0, a1, -2.0 -// GFX908: error: invalid literal operand +// GFX908: error: inline constants are not allowed for this operand // GFX900: error: instruction not supported on this GPU v_mfma_f32_4x4x2bf16 a[0:3], a0, a1, -2.0 cbsz:3 abid:2 blgp:7 -// GFX908: error: invalid literal operand +// GFX908: error: inline constants are not allowed for this operand // GFX900: error: instruction not supported on this GPU v_mfma_f32_32x32x4bf16 a[0:15], v0, v1, -2.0 -// GFX908: error: invalid literal operand +// GFX908: error: inline constants are not allowed for this operand // GFX900: error: instruction not supported on this GPU v_mfma_f32_32x32x4bf16 a[0:15], v0, v1, -2.0 cbsz:3 abid:2 blgp:7 -// GFX908: error: invalid literal operand +// GFX908: error: inline constants are not allowed for this operand // GFX900: error: instruction not supported on this GPU v_mfma_f32_32x32x4bf16 a[0:15], v0, a1, -2.0 -// GFX908: error: invalid literal operand +// GFX908: error: inline constants are not allowed for this operand // GFX900: error: instruction not supported on this GPU v_mfma_f32_32x32x4bf16 a[0:15], v0, a1, -2.0 cbsz:3 abid:2 blgp:7 -// GFX908: error: invalid literal operand +// GFX908: error: inline constants are not allowed for this operand // GFX900: error: instruction not supported on this GPU v_mfma_f32_32x32x4bf16 a[0:15], a0, v1, -2.0 -// GFX908: error: invalid literal operand +// GFX908: error: inline constants are not allowed for this operand // GFX900: error: instruction not supported on this GPU v_mfma_f32_32x32x4bf16 a[0:15], a0, v1, -2.0 cbsz:3 abid:2 blgp:7 -// GFX908: error: invalid literal operand +// GFX908: error: inline constants are not allowed for this operand // GFX900: error: instruction not supported on this GPU v_mfma_f32_32x32x4bf16 a[0:15], a0, a1, -2.0 -// GFX908: error: invalid literal operand +// GFX908: error: inline constants are not allowed for this operand // GFX900: error: instruction not supported on this GPU v_mfma_f32_32x32x4bf16 a[0:15], a0, a1, -2.0 cbsz:3 abid:2 blgp:7 -// GFX908: error: invalid literal operand +// GFX908: error: inline constants are not allowed for this operand // GFX900: error: instruction not supported on this GPU v_mfma_f32_16x16x8bf16 a[0:3], v0, v1, -2.0 -// GFX908: error: invalid literal operand +// GFX908: error: inline constants are not allowed for this operand // GFX900: error: instruction not supported on this GPU v_mfma_f32_16x16x8bf16 a[0:3], v0, v1, -2.0 cbsz:3 abid:2 blgp:7 -// GFX908: error: invalid literal operand +// GFX908: error: inline constants are not allowed for this operand // GFX900: error: instruction not supported on this GPU v_mfma_f32_16x16x8bf16 a[0:3], v0, a1, -2.0 -// GFX908: error: invalid literal operand +// GFX908: error: inline constants are not allowed for this operand // GFX900: error: instruction not supported on this GPU v_mfma_f32_16x16x8bf16 a[0:3], v0, a1, -2.0 cbsz:3 abid:2 blgp:7 -// GFX908: error: invalid literal operand +// GFX908: error: inline constants are not allowed for this operand // GFX900: error: instruction not supported on this GPU v_mfma_f32_16x16x8bf16 a[0:3], a0, v1, -2.0 -// GFX908: error: invalid literal operand +// GFX908: error: inline constants are not allowed for this operand // GFX900: error: instruction not supported on this GPU v_mfma_f32_16x16x8bf16 a[0:3], a0, v1, -2.0 cbsz:3 abid:2 blgp:7 -// GFX908: error: invalid literal operand +// GFX908: error: inline constants are not allowed for this operand // GFX900: error: instruction not supported on this GPU v_mfma_f32_16x16x8bf16 a[0:3], a0, a1, -2.0 -// GFX908: error: invalid literal operand +// GFX908: error: inline constants are not allowed for this operand // GFX900: error: instruction not supported on this GPU v_mfma_f32_16x16x8bf16 a[0:3], a0, a1, -2.0 cbsz:3 abid:2 blgp:7 -// GFX908: error: invalid literal operand +// GFX908: error: inline constants are not allowed for this operand // GFX900: error: instruction not supported on this GPU diff --git a/llvm/test/MC/AMDGPU/vop2-err.s b/llvm/test/MC/AMDGPU/vop2-err.s --- a/llvm/test/MC/AMDGPU/vop2-err.s +++ b/llvm/test/MC/AMDGPU/vop2-err.s @@ -6,7 +6,7 @@ //===----------------------------------------------------------------------===// v_mul_i32_i24 v1, v2, 100 -// CHECK: error: invalid literal operand +// CHECK: error: literal operands are not supported //===----------------------------------------------------------------------===// // _e32 checks @@ -29,11 +29,11 @@ // Immediate src0 v_mul_i32_i24_e64 v1, 100, v3 -// CHECK: error: invalid literal operand +// CHECK: error: literal operands are not supported // Immediate src1 v_mul_i32_i24_e64 v1, v2, 100 -// CHECK: error: invalid literal operand +// CHECK: error: literal operands are not supported v_add_i32_e32 v1, s[0:1], v2, v3 // CHECK: error: invalid operand for instruction diff --git a/llvm/test/MC/AMDGPU/vop3-errs.s b/llvm/test/MC/AMDGPU/vop3-errs.s --- a/llvm/test/MC/AMDGPU/vop3-errs.s +++ b/llvm/test/MC/AMDGPU/vop3-errs.s @@ -7,7 +7,7 @@ // GCN: error: too few operands for instruction v_div_scale_f32 v24, vcc, v22, 1.1, v22 -// GCN: error: invalid literal operand +// GCN: error: literal operands are not supported v_mqsad_u32_u8 v[0:3], s[2:3], v4, v[0:3] // GFX67: error: instruction not supported on this GPU diff --git a/llvm/test/MC/AMDGPU/vop3-gfx9.s b/llvm/test/MC/AMDGPU/vop3-gfx9.s --- a/llvm/test/MC/AMDGPU/vop3-gfx9.s +++ b/llvm/test/MC/AMDGPU/vop3-gfx9.s @@ -426,9 +426,9 @@ // VI: v_mad_i16 v5, v1, -1, v3 ; encoding: [0x05,0x00,0xec,0xd1,0x01,0x83,0x0d,0x04] v_mad_i16 v5, v1, v2, -4.0 -// NOGFX9: error: invalid literal operand +// NOGFX9: error: literal operands are not supported // NOSICI: error: instruction not supported on this GPU -// NOVI: error: invalid literal operand +// NOVI: error: literal operands are not supported v_mad_i16 v5, v1, v2, v3 clamp // GFX9: v_mad_i16 v5, v1, v2, v3 clamp ; encoding: [0x05,0x80,0x05,0xd2,0x01,0x05,0x0e,0x04] @@ -478,11 +478,11 @@ // NOGCN: error: instruction not supported on this GPU v_mad_legacy_i16 v5, v1, v2, -4.0 -// NOGFX9: error: invalid literal operand +// NOGFX9: error: literal operands are not supported // NOGCN: error: instruction not supported on this GPU v_mad_legacy_i16 v5, v1, v2, -4.0 clamp -// NOGFX9: error: invalid literal operand +// NOGFX9: error: literal operands are not supported // NOGCN: error: instruction not supported on this GPU v_mad_legacy_u16_e64 v5, 0, v2, v3 @@ -494,11 +494,11 @@ // NOGCN: error: instruction not supported on this GPU v_mad_legacy_u16 v5, v1, v2, -4.0 -// NOGFX9: error: invalid literal operand +// NOGFX9: error: literal operands are not supported // NOGCN: error: instruction not supported on this GPU v_mad_legacy_u16 v5, v1, v2, -4.0 clamp -// NOGFX9: error: invalid literal operand +// NOGFX9: error: literal operands are not supported // NOGCN: error: instruction not supported on this GPU v_mad_u16_e64 v5, 0, v2, v3 @@ -512,9 +512,9 @@ // VI: v_mad_u16 v5, v1, -1, v3 ; encoding: [0x05,0x00,0xeb,0xd1,0x01,0x83,0x0d,0x04] v_mad_u16 v5, v1, v2, -4.0 -// NOGFX9: error: invalid literal operand +// NOGFX9: error: literal operands are not supported // NOSICI: error: instruction not supported on this GPU -// NOVI: error: invalid literal operand +// NOVI: error: literal operands are not supported v_mad_u16 v5, v1, v2, v3 clamp // GFX9: v_mad_u16 v5, v1, v2, v3 clamp ; encoding: [0x05,0x80,0x04,0xd2,0x01,0x05,0x0e,0x04] diff --git a/llvm/test/MC/AMDGPU/vop3-literal.s b/llvm/test/MC/AMDGPU/vop3-literal.s --- a/llvm/test/MC/AMDGPU/vop3-literal.s +++ b/llvm/test/MC/AMDGPU/vop3-literal.s @@ -5,91 +5,91 @@ v_bfe_u32 v0, 0x3039, v1, s1 // GFX10: v_bfe_u32 v0, 0x3039, v1, s1 ; encoding: [0x00,0x00,0x48,0xd5,0xff,0x02,0x06,0x00,0x39,0x30,0x00,0x00] -// GFX9-ERR: error: invalid literal operand +// GFX9-ERR: error: literal operands are not supported v_bfe_u32 v0, v1, 0x3039, s1 // GFX10: v_bfe_u32 v0, v1, 0x3039, s1 ; encoding: [0x00,0x00,0x48,0xd5,0x01,0xff,0x05,0x00,0x39,0x30,0x00,0x00] -// GFX9-ERR: error: invalid literal operand +// GFX9-ERR: error: literal operands are not supported v_bfe_u32 v0, v1, s1, 0x3039 // GFX10: v_bfe_u32 v0, v1, s1, 0x3039 ; encoding: [0x00,0x00,0x48,0xd5,0x01,0x03,0xfc,0x03,0x39,0x30,0x00,0x00] -// GFX9-ERR: error: invalid literal operand +// GFX9-ERR: error: literal operands are not supported 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] -// GFX9-ERR: error: invalid literal operand +// GFX9-ERR: 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] -// GFX9-ERR: error: invalid literal operand +// GFX9-ERR: 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] -// GFX9-ERR: error: invalid literal operand +// GFX9-ERR: 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] -// GFX9-ERR: error: invalid literal operand +// GFX9-ERR: error: literal operands are not supported v_bfe_u32 v0, 0x3039, s1, 0x3038 -// GFX10-ERR: error: invalid literal operand -// GFX9-ERR: error: invalid literal operand +// GFX10-ERR: error: only one literal operand is allowed +// GFX9-ERR: error: literal operands are not supported v_bfe_u32 v0, 0x3039, v1, v2 // GFX10: v_bfe_u32 v0, 0x3039, v1, v2 ; encoding: [0x00,0x00,0x48,0xd5,0xff,0x02,0x0a,0x04,0x39,0x30,0x00,0x00] -// GFX9-ERR: error: invalid literal operand +// GFX9-ERR: error: literal operands are not supported v_bfe_u32 v0, 0x3039, 0x12345, v2 -// GFX10-ERR: error: invalid literal operand -// GFX9-ERR: error: invalid literal operand +// GFX10-ERR: error: only one literal operand is allowed +// GFX9-ERR: error: literal operands are not supported v_bfe_u32 v0, s1, 0x3039, s1 -// GFX9-ERR: error: invalid literal operand +// GFX9-ERR: error: literal operands are not supported // GFX10: v_bfe_u32 v0, s1, 0x3039, s1 ; encoding: [0x00,0x00,0x48,0xd5,0x01,0xfe,0x05,0x00,0x39,0x30,0x00,0x00] v_bfe_u32 v0, s1, 0x3039, s2 -// GFX9-ERR: error: invalid literal operand +// GFX9-ERR: error: literal operands are not supported // GFX10-ERR: error: invalid operand (violates constant bus restrictions) v_bfm_b32_e64 v0, 0x3039, s1 // GFX10: v_bfm_b32_e64 v0, 0x3039, s1 ; encoding: [0x00,0x00,0x63,0xd7,0xff,0x02,0x00,0x00,0x39,0x30,0x00,0x00] -// GFX9-ERR: error: invalid literal operand +// GFX9-ERR: error: literal operands are not supported v_bfm_b32_e64 v0, 0x3039, v1 // GFX10: v_bfm_b32_e64 v0, 0x3039, v1 ; encoding: [0x00,0x00,0x63,0xd7,0xff,0x02,0x02,0x00,0x39,0x30,0x00,0x00] -// GFX9-ERR: error: invalid literal operand +// GFX9-ERR: error: literal operands are not supported v_bfm_b32_e64 v0, 0x3039, 0x3039 // GFX10: v_bfm_b32_e64 v0, 0x3039, 0x3039 ; encoding: [0x00,0x00,0x63,0xd7,0xff,0xfe,0x01,0x00,0x39,0x30,0x00,0x00] -// GFX9-ERR: error: invalid literal operand +// GFX9-ERR: error: literal operands are not supported v_bfm_b32_e64 v0, 0x3039, 0x3038 -// GFX10-ERR: error: invalid literal operand -// GFX9-ERR: error: invalid literal operand +// GFX10-ERR: error: only one literal operand is allowed +// GFX9-ERR: error: literal operands are not supported v_pk_add_f16 v1, 25.0, v2 // GFX10: v_pk_add_f16 v1, 0x4e40, v2 ; encoding: [0x01,0x00,0x0f,0xcc,0xff,0x04,0x02,0x18,0x40,0x4e,0x00,0x00] -// GFX9-ERR: error: invalid literal operand +// GFX9-ERR: error: literal operands are not supported v_pk_add_f16 v1, 123456, v2 // GFX10: v_pk_add_f16 v1, 0x1e240, v2 ; encoding: [0x01,0x00,0x0f,0xcc,0xff,0x04,0x02,0x18,0x40,0xe2,0x01,0x00] -// GFX9-ERR: error: invalid literal operand +// GFX9-ERR: error: literal operands are not supported v_pk_add_f16 v1, -200, v2 // GFX10: v_pk_add_f16 v1, 0xffffff38, v2 ; encoding: [0x01,0x00,0x0f,0xcc,0xff,0x04,0x02,0x18,0x38,0xff,0xff,0xff] -// GFX9-ERR: error: invalid literal operand +// GFX9-ERR: error: literal operands are not supported v_pk_add_f16 v1, 25.0, 25.0 // GFX10: v_pk_add_f16 v1, 0x4e40, 0x4e40 ; encoding: [0x01,0x00,0x0f,0xcc,0xff,0xfe,0x01,0x18,0x40,0x4e,0x00,0x00] -// GFX9-ERR: error: invalid literal operand +// GFX9-ERR: error: literal operands are not supported v_pk_add_f16 v1, 25.0, 25.1 -// GFX10-ERR: error: invalid literal operand -// GFX9-ERR: error: invalid literal operand +// GFX10-ERR: error: only one literal operand is allowed +// GFX9-ERR: error: literal operands are not supported v_pk_add_u16 v1, -200, v2 // GFX10: v_pk_add_u16 v1, 0xffffff38, v2 ; encoding: [0x01,0x00,0x0a,0xcc,0xff,0x04,0x02,0x18,0x38,0xff,0xff,0xff] -// GFX9-ERR: error: invalid literal operand +// GFX9-ERR: error: literal operands are not supported v_pk_add_u16 v1, 64, v2 // GFX10: v_pk_add_u16 v1, 64, v2 ; encoding: [0x01,0x00,0x0a,0xcc,0xc0,0x04,0x02,0x18] @@ -97,7 +97,7 @@ v_pk_add_u16 v1, 65, v2 // GFX10: v_pk_add_u16 v1, 0x41, v2 ; encoding: [0x01,0x00,0x0a,0xcc,0xff,0x04,0x02,0x18,0x41,0x00,0x00,0x00] -// GFX9-ERR: error: invalid literal operand +// GFX9-ERR: error: literal operands are not supported v_pk_add_u16 v1, -1, v2 // GFX10: v_pk_add_u16 v1, -1, v2 ; encoding: [0x01,0x00,0x0a,0xcc,0xc1,0x04,0x02,0x18] @@ -109,91 +109,91 @@ v_pk_add_u16 v1, -100, v2 // GFX10: v_pk_add_u16 v1, 0xffffff9c, v2 ; encoding: [0x01,0x00,0x0a,0xcc,0xff,0x04,0x02,0x18,0x9c,0xff,0xff,0xff] -// GFX9-ERR: error: invalid literal operand +// GFX9-ERR: error: literal operands are not supported v_pk_add_u16 v1, -100, -100 // GFX10: v_pk_add_u16 v1, 0xffffff9c, 0xffffff9c ; encoding: [0x01,0x00,0x0a,0xcc,0xff,0xfe,0x01,0x18,0x9c,0xff,0xff,0xff] -// GFX9-ERR: error: invalid literal operand +// GFX9-ERR: 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] -// GFX9-ERR: error: invalid literal operand +// GFX9-ERR: error: literal operands are not supported v_add_f32_e64 v1, v3, neg(0x123) // GFX10: v_add_f32_e64 v1, v3, neg(0x123) ; encoding: [0x01,0x00,0x03,0xd5,0x03,0xff,0x01,0x40,0x23,0x01,0x00,0x00] -// GFX9-ERR: error: invalid literal operand +// GFX9-ERR: error: literal operands are not supported 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] -// GFX9-ERR: error: invalid literal operand +// GFX9-ERR: 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] -// GFX9-ERR: error: invalid literal operand +// GFX9-ERR: error: literal operands are not supported v_add_f16_e64 v0, v0, neg(0xfe0b) // GFX10: v_add_f16_e64 v0, v0, neg(0xfe0b) ; encoding: [0x00,0x00,0x32,0xd5,0x00,0xff,0x01,0x40,0x0b,0xfe,0x00,0x00] -// GFX9-ERR: error: invalid literal operand +// GFX9-ERR: error: literal operands are not supported v_add_f16_e64 v0, 0x3456, v0 // GFX10: v_add_f16_e64 v0, 0x3456, v0 ; encoding: [0x00,0x00,0x32,0xd5,0xff,0x00,0x02,0x00,0x56,0x34,0x00,0x00] -// GFX9-ERR: error: invalid literal operand +// GFX9-ERR: error: literal operands are not supported 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] -// GFX9-ERR: error: invalid literal operand +// GFX9-ERR: 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] -// GFX9-ERR: error: invalid literal operand +// GFX9-ERR: error: literal operands are not supported v_add_f64 v[0:1], v[0:1], -abs(1.23456) // GFX10: v_add_f64 v[0:1], v[0:1], -|0x3ff3c0c1| ; encoding: [0x00,0x02,0x64,0xd5,0x00,0xff,0x01,0x40,0xc1,0xc0,0xf3,0x3f] -// GFX9-ERR: error: invalid literal operand +// GFX9-ERR: error: literal operands are not supported 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] -// GFX9-ERR: error: invalid literal operand +// GFX9-ERR: error: literal operands are not supported v_add_f64 v[0:1], 1.23456, -abs(1.2345) -// GFX10-ERR: error: invalid literal operand -// GFX9-ERR: error: invalid literal operand +// GFX10-ERR: error: only one literal operand is allowed +// GFX9-ERR: error: literal operands are not supported v_max_i16_e64 v5, 0xfe0b, v2 // GFX10: v_max_i16_e64 v5, 0xfe0b, v2 ; encoding: [0x05,0x00,0x0a,0xd7,0xff,0x04,0x02,0x00,0x0b,0xfe,0x00,0x00] -// GFX9-ERR: error: invalid literal operand +// GFX9-ERR: error: literal operands are not supported v_max_i16_e64 v5, v1, 0x123 // GFX10: v_max_i16_e64 v5, v1, 0x123 ; encoding: [0x05,0x00,0x0a,0xd7,0x01,0xff,0x01,0x00,0x23,0x01,0x00,0x00] -// GFX9-ERR: error: invalid literal operand +// GFX9-ERR: error: literal operands are not supported v_max_i16_e64 v5, 0x1234, 0x1234 // GFX10: v_max_i16_e64 v5, 0x1234, 0x1234 ; encoding: [0x05,0x00,0x0a,0xd7,0xff,0xfe,0x01,0x00,0x34,0x12,0x00,0x00] -// GFX9-ERR: error: invalid literal operand +// GFX9-ERR: 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] -// GFX9-ERR: error: invalid literal operand +// GFX9-ERR: error: literal operands are not supported v_min3_i16 v5, v1, 0x1234, v3 // GFX10: v_min3_i16 v5, v1, 0x1234, v3 ; encoding: [0x05,0x00,0x52,0xd7,0x01,0xff,0x0d,0x04,0x34,0x12,0x00,0x00] -// GFX9-ERR: error: invalid literal operand +// GFX9-ERR: error: literal operands are not supported v_min3_i16 v5, v1, v2, 0x5678 // GFX10: v_min3_i16 v5, v1, v2, 0x5678 ; encoding: [0x05,0x00,0x52,0xd7,0x01,0x05,0xfe,0x03,0x78,0x56,0x00,0x00] -// GFX9-ERR: error: invalid literal operand +// GFX9-ERR: error: literal operands are not supported 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] -// GFX9-ERR: error: invalid literal operand +// GFX9-ERR: error: literal operands are not supported v_min3_i16 v5, 0x5678, 0x5679, 0x5678 -// GFX10-ERR: error: invalid literal operand -// GFX9-ERR: error: invalid literal operand +// GFX10-ERR: error: only one literal operand is allowed +// GFX9-ERR: error: literal operands are not supported v_min3_i16 v5, 0x5678, 0x5678, 0x5679 -// GFX10-ERR: error: invalid literal operand -// GFX9-ERR: error: invalid literal operand +// GFX10-ERR: error: only one literal operand is allowed +// GFX9-ERR: error: literal operands are not supported v_add_nc_u16 v5, 0xfe0b, v2 // GFX10: v_add_nc_u16_e64 v5, 0xfe0b, v2 ; encoding: [0x05,0x00,0x03,0xd7,0xff,0x04,0x02,0x00,0x0b,0xfe,0x00,0x00] @@ -209,67 +209,67 @@ v_ashrrev_i16_e64 v5, 0x3456, v2 // GFX10: v_ashrrev_i16_e64 v5, 0x3456, v2 ; encoding: [0x05,0x00,0x08,0xd7,0xff,0x04,0x02,0x00,0x56,0x34,0x00,0x00] -// GFX9-ERR: error: invalid literal operand +// GFX9-ERR: error: literal operands are not supported v_mad_u16 v5, 0xfe0b, v2, v3 // GFX10: v_mad_u16 v5, 0xfe0b, v2, v3 ; encoding: [0x05,0x00,0x40,0xd7,0xff,0x04,0x0e,0x04,0x0b,0xfe,0x00,0x00] -// GFX9-ERR: error: invalid literal operand +// GFX9-ERR: error: literal operands are not supported v_mad_u16 v5, v1, 0x1234, v3 // GFX10: v_mad_u16 v5, v1, 0x1234, v3 ; encoding: [0x05,0x00,0x40,0xd7,0x01,0xff,0x0d,0x04,0x34,0x12,0x00,0x00] -// GFX9-ERR: error: invalid literal operand +// GFX9-ERR: error: literal operands are not supported v_mad_u16 v5, v1, v2, 0x5678 // GFX10: v_mad_u16 v5, v1, v2, 0x5678 ; encoding: [0x05,0x00,0x40,0xd7,0x01,0x05,0xfe,0x03,0x78,0x56,0x00,0x00] -// GFX9-ERR: error: invalid literal operand +// GFX9-ERR: error: literal operands are not supported 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] -// GFX9-ERR: error: invalid literal operand +// GFX9-ERR: 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] -// GFX9-ERR: error: invalid literal operand +// GFX9-ERR: error: literal operands are not supported v_mad_legacy_f32 v5, v1, 0xaf123456, v3 // GFX10: v_mad_legacy_f32 v5, v1, 0xaf123456, v3 ; encoding: [0x05,0x00,0x40,0xd5,0x01,0xff,0x0d,0x04,0x56,0x34,0x12,0xaf] -// GFX9-ERR: error: invalid literal operand +// GFX9-ERR: error: literal operands are not supported v_mad_legacy_f32 v5, v1, v2, 0xaf123456 // GFX10: v_mad_legacy_f32 v5, v1, v2, 0xaf123456 ; encoding: [0x05,0x00,0x40,0xd5,0x01,0x05,0xfe,0x03,0x56,0x34,0x12,0xaf] -// GFX9-ERR: error: invalid literal operand +// GFX9-ERR: error: literal operands are not supported 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: error: invalid literal operand +// GFX9-ERR: error: literal operands are not supported v_cmp_f_i32_e64 s[10:11], 0xaf123456, v2 // GFX10: v_cmp_f_i32_e64 s[10:11], 0xaf123456, v2 ; encoding: [0x0a,0x00,0x80,0xd4,0xff,0x04,0x02,0x00,0x56,0x34,0x12,0xaf] -// GFX9-ERR: error: invalid literal operand +// GFX9-ERR: error: literal operands are not supported v_cmp_f_i32_e64 s[10:11], v1, 0xaf123456 // GFX10: v_cmp_f_i32_e64 s[10:11], v1, 0xaf123456 ; encoding: [0x0a,0x00,0x80,0xd4,0x01,0xff,0x01,0x00,0x56,0x34,0x12,0xaf] -// GFX9-ERR: error: invalid literal operand +// GFX9-ERR: error: literal operands are not supported v_cmp_f_i32_e64 s[10:11], 0xaf123456, 0xaf123456 // GFX10: v_cmp_f_i32_e64 s[10:11], 0xaf123456, 0xaf123456 ; encoding: [0x0a,0x00,0x80,0xd4,0xff,0xfe,0x01,0x00,0x56,0x34,0x12,0xaf] -// GFX9-ERR: error: invalid literal operand +// GFX9-ERR: error: literal operands are not supported v_cmp_f_i32_e64 s[10:11], 0xaf123456, 0xaf123455 -// GFX10-ERR: error: invalid literal operand -// GFX9-ERR: error: invalid literal operand +// GFX10-ERR: error: only one literal operand is allowed +// GFX9-ERR: error: literal operands are not supported v_cmp_f_u64_e64 s[10:11], 0xaf123456, v[2:3] // GFX10: v_cmp_f_u64_e64 s[10:11], 0xaf123456, v[2:3] ; encoding: [0x0a,0x00,0xe0,0xd4,0xff,0x04,0x02,0x00,0x56,0x34,0x12,0xaf] -// GFX9-ERR: error: invalid literal operand +// GFX9-ERR: error: literal operands are not supported v_cmp_f_u64_e64 s[10:11], v[1:2], 0x3f717273 // GFX10: v_cmp_f_u64_e64 s[10:11], v[1:2], 0x3f717273 ; encoding: [0x0a,0x00,0xe0,0xd4,0x01,0xff,0x01,0x00,0x73,0x72,0x71,0x3f] -// GFX9-ERR: error: invalid literal operand +// GFX9-ERR: error: literal operands are not supported v_cmp_f_u64_e64 s[10:11], 0x3f717273, 0x3f717273 // GFX10: v_cmp_f_u64_e64 s[10:11], 0x3f717273, 0x3f717273 ; encoding: [0x0a,0x00,0xe0,0xd4,0xff,0xfe,0x01,0x00,0x73,0x72,0x71,0x3f] -// GFX9-ERR: error: invalid literal operand +// GFX9-ERR: error: literal operands are not supported v_cmpx_class_f32_e64 0xaf123456, v2 // GFX10: v_cmpx_class_f32_e64 0xaf123456, v2 ; encoding: [0x00,0x00,0x98,0xd4,0xff,0x04,0x02,0x00,0x56,0x34,0x12,0xaf] @@ -284,7 +284,7 @@ // GFX9-ERR: error: operands are not valid for this GPU or mode v_cmpx_class_f32_e64 0xaf123456, 0xaf123455 -// GFX10-ERR: error: invalid literal operand +// GFX10-ERR: error: only one literal operand is allowed // GFX9-ERR: error: operands are not valid for this GPU or mode v_cmpx_lt_i16_e64 v1, 0x3456 @@ -313,11 +313,11 @@ v_lshlrev_b64 v[5:6], 0xaf123456, v[2:3] // GFX10: v_lshlrev_b64 v[5:6], 0xaf123456, v[2:3] ; encoding: [0x05,0x00,0xff,0xd6,0xff,0x04,0x02,0x00,0x56,0x34,0x12,0xaf] -// GFX9-ERR: error: invalid literal operand +// GFX9-ERR: error: literal operands are not supported v_lshlrev_b64 v[5:6], v1, 0x3f717273 // GFX10: v_lshlrev_b64 v[5:6], v1, 0x3f717273 ; encoding: [0x05,0x00,0xff,0xd6,0x01,0xff,0x01,0x00,0x73,0x72,0x71,0x3f] -// GFX9-ERR: error: invalid literal operand +// GFX9-ERR: error: literal operands are not supported v_fma_mix_f32 v5, 0x123, v2, v3 // GFX10: v_fma_mix_f32 v5, 0x123, v2, v3 ; encoding: [0x05,0x00,0x20,0xcc,0xff,0x04,0x0e,0x04,0x23,0x01,0x00,0x00] @@ -341,64 +341,64 @@ v_pk_add_f16 v5, 0xaf123456, v2 // GFX10: v_pk_add_f16 v5, 0xaf123456, v2 ; encoding: [0x05,0x00,0x0f,0xcc,0xff,0x04,0x02,0x18,0x56,0x34,0x12,0xaf] -// GFX9-ERR: error: invalid literal operand +// GFX9-ERR: error: literal operands are not supported v_pk_add_f16 v5, v1, 0x3f717273 // GFX10: v_pk_add_f16 v5, v1, 0x3f717273 ; encoding: [0x05,0x00,0x0f,0xcc,0x01,0xff,0x01,0x18,0x73,0x72,0x71,0x3f] -// GFX9-ERR: error: invalid literal operand +// GFX9-ERR: error: literal operands are not supported v_pk_add_f16 v5, 0x3f717273, 0x3f717273 // GFX10: v_pk_add_f16 v5, 0x3f717273, 0x3f717273 ; encoding: [0x05,0x00,0x0f,0xcc,0xff,0xfe,0x01,0x18,0x73,0x72,0x71,0x3f] -// GFX9-ERR: error: invalid literal operand +// GFX9-ERR: error: literal operands are not supported v_pk_add_i16 v5, 0x7b, v2 // GFX10: v_pk_add_i16 v5, 0x7b, v2 ; encoding: [0x05,0x00,0x02,0xcc,0xff,0x04,0x02,0x18,0x7b,0x00,0x00,0x00] -// GFX9-ERR: error: invalid literal operand +// GFX9-ERR: error: literal operands are not supported v_pk_add_i16 v5, v1, 0x7b // GFX10: v_pk_add_i16 v5, v1, 0x7b ; encoding: [0x05,0x00,0x02,0xcc,0x01,0xff,0x01,0x18,0x7b,0x00,0x00,0x00] -// GFX9-ERR: error: invalid literal operand +// GFX9-ERR: error: literal operands are not supported v_pk_add_i16 v5, 0xab7b, 0xab7b // GFX10: v_pk_add_i16 v5, 0xab7b, 0xab7b ; encoding: [0x05,0x00,0x02,0xcc,0xff,0xfe,0x01,0x18,0x7b,0xab,0x00,0x00] -// GFX9-ERR: error: invalid literal operand +// GFX9-ERR: error: literal operands are not supported v_pk_add_i16 v5, 0xab7b, 0xab7a -// GFX10-ERR: error: invalid literal operand -// GFX9-ERR: error: invalid literal operand +// GFX10-ERR: error: only one literal operand is allowed +// GFX9-ERR: error: literal operands are not supported v_div_fmas_f32 v5, v1, 0x123, v3 // GFX10: v_div_fmas_f32 v5, v1, 0x123, v3 ; encoding: [0x05,0x00,0x6f,0xd5,0x01,0xff,0x0d,0x04,0x23,0x01,0x00,0x00] -// GFX9-ERR: error: invalid literal operand +// GFX9-ERR: error: literal operands are not supported 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] -// GFX9-ERR: error: invalid literal operand +// GFX9-ERR: 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] -// GFX9-ERR: error: invalid literal operand +// GFX9-ERR: error: literal operands are not supported v_div_fmas_f64 v[5:6], 0x12345678, v[2:3], v[3:4] // GFX10: v_div_fmas_f64 v[5:6], 0x12345678, v[2:3], v[3:4] ; encoding: [0x05,0x00,0x70,0xd5,0xff,0x04,0x0e,0x04,0x78,0x56,0x34,0x12] -// GFX9-ERR: error: invalid literal operand +// GFX9-ERR: error: literal operands are not supported v_div_fmas_f64 v[5:6], 0x12345678, 0x12345678, 0x12345678 // GFX10: v_div_fmas_f64 v[5:6], 0x12345678, 0x12345678, 0x12345678 ; encoding: [0x05,0x00,0x70,0xd5,0xff,0xfe,0xfd,0x03,0x78,0x56,0x34,0x12] -// GFX9-ERR: error: invalid literal operand +// GFX9-ERR: error: literal operands are not supported v_div_fmas_f64 v[5:6], v[1:2], 0x123457, 0x123456 -// GFX10-ERR: error: invalid literal operand -// GFX9-ERR: error: invalid literal operand +// GFX10-ERR: error: only one literal operand is allowed +// GFX9-ERR: error: literal operands are not supported v_ldexp_f64 v[5:6], 0.12345, v2 // GFX10: v_ldexp_f64 v[5:6], 0x3fbf9a6b, v2 ; encoding: [0x05,0x00,0x68,0xd5,0xff,0x04,0x02,0x00,0x6b,0x9a,0xbf,0x3f] -// GFX9-ERR: error: invalid literal operand +// GFX9-ERR: error: literal operands are not supported v_ldexp_f64 v[5:6], 0.12345, 0x3fbf9a6b // GFX10: v_ldexp_f64 v[5:6], 0x3fbf9a6b, 0x3fbf9a6b ; encoding: [0x05,0x00,0x68,0xd5,0xff,0xfe,0x01,0x00,0x6b,0x9a,0xbf,0x3f] -// GFX9-ERR: error: invalid literal operand +// GFX9-ERR: error: literal operands are not supported v_ldexp_f64 v[5:6], 0.12345, 0x3fbf9a6c -// GFX10-ERR: error: invalid literal operand -// GFX9-ERR: error: invalid literal operand +// GFX10-ERR: error: only one literal operand is allowed +// GFX9-ERR: error: literal operands are not supported diff --git a/llvm/test/MC/AMDGPU/vop3.s b/llvm/test/MC/AMDGPU/vop3.s --- a/llvm/test/MC/AMDGPU/vop3.s +++ b/llvm/test/MC/AMDGPU/vop3.s @@ -599,7 +599,7 @@ // NOSICI: error: instruction not supported on this GPU v_mad_i16 v5, v1, -4.0, v3 -// NOVI: error: invalid literal operand +// NOVI: error: literal operands are not supported // NOSICI: error: instruction not supported on this GPU v_mad_i16 v5, v1, v2, 0 @@ -615,7 +615,7 @@ // NOSICI: error: instruction not supported on this GPU v_mad_u16 v5, v1, v2, -4.0 -// NOVI: error: invalid literal operand +// NOVI: error: literal operands are not supported // NOSICI: error: instruction not supported on this GPU ///===---------------------------------------------------------------------===//