[llvm] r263008 - [AMDGPU] Assembler: Support DPP instructions.
Teresa Johnson via llvm-commits
llvm-commits at lists.llvm.org
Wed Mar 9 07:04:02 PST 2016
I'm getting the following error building this commit with clang:
/usr/local/google/home/tejohnson/llvm/llvm_15/lib/Target/AMDGPU/InstPrinter/AMDGPUInstPrinter.cpp:407:12:
error: comparison of unsigned expression >= 0 is always true
[-Werror,-Wtautological-compare]
if ((Imm >= 0x000) && (Imm <= 0x0ff)) {
~~~ ^ ~~~~~
(went ahead and committed a fix in r263014)
Thanks,
Teresa
On Wed, Mar 9, 2016 at 4:29 AM, Sam Kolton via llvm-commits <
llvm-commits at lists.llvm.org> wrote:
> Author: skolton
> Date: Wed Mar 9 06:29:31 2016
> New Revision: 263008
>
> URL: http://llvm.org/viewvc/llvm-project?rev=263008&view=rev
> Log:
> [AMDGPU] Assembler: Support DPP instructions.
>
> Supprot DPP syntax as used in SP3 (except several operands syntax).
> Added dpp-specific operands in td-files.
> Added DPP flag to TSFlags to determine if instruction is dpp in
> InstPrinter.
> Support for VOP2 DPP instructions in td-files.
> Some tests for DPP instructions.
>
> ToDo:
> - VOP2bInst:
> - vcc is considered as operand
> - AsmMatcher doesn't apply mnemonic aliases when parsing operands
> - v_mac_f32
> - v_nop
> - disable instructions with 64-bit operands
> - change dpp_ctrl assembler representation to conform sp3
>
> Review: http://reviews.llvm.org/D17804
>
> Added:
> llvm/trunk/test/MC/AMDGPU/vop_dpp.s
> Modified:
> llvm/trunk/include/llvm/IR/IntrinsicsAMDGPU.td
> llvm/trunk/lib/Target/AMDGPU/AsmParser/AMDGPUAsmParser.cpp
> llvm/trunk/lib/Target/AMDGPU/InstPrinter/AMDGPUInstPrinter.cpp
> llvm/trunk/lib/Target/AMDGPU/InstPrinter/AMDGPUInstPrinter.h
> llvm/trunk/lib/Target/AMDGPU/SIDefines.h
> llvm/trunk/lib/Target/AMDGPU/SIInstrFormats.td
> llvm/trunk/lib/Target/AMDGPU/SIInstrInfo.td
> llvm/trunk/lib/Target/AMDGPU/VIInstrFormats.td
> llvm/trunk/lib/Target/AMDGPU/VIInstructions.td
> llvm/trunk/test/CodeGen/AMDGPU/llvm.amdgcn.mov.dpp.ll
>
> Modified: llvm/trunk/include/llvm/IR/IntrinsicsAMDGPU.td
> URL:
> http://llvm.org/viewvc/llvm-project/llvm/trunk/include/llvm/IR/IntrinsicsAMDGPU.td?rev=263008&r1=263007&r2=263008&view=diff
>
> ==============================================================================
> --- llvm/trunk/include/llvm/IR/IntrinsicsAMDGPU.td (original)
> +++ llvm/trunk/include/llvm/IR/IntrinsicsAMDGPU.td Wed Mar 9 06:29:31 2016
> @@ -275,11 +275,11 @@ def int_amdgcn_buffer_wbinvl1_vol :
> // VI Intrinsics
>
> //===----------------------------------------------------------------------===//
>
> -// llvm.amdgcn.mov.dpp.i32 <src> <dpp_ctrl> <bound_ctrl> <bank_mask>
> <row_mask>
> +// llvm.amdgcn.mov.dpp.i32 <src> <dpp_ctrl> <row_mask> <bank_mask>
> <bound_ctrl>
> def int_amdgcn_mov_dpp :
> Intrinsic<[llvm_anyint_ty],
> - [LLVMMatchType<0>, llvm_i32_ty, llvm_i1_ty, llvm_i32_ty,
> - llvm_i32_ty], [IntrNoMem, IntrConvergent]>;
> + [LLVMMatchType<0>, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty,
> + llvm_i1_ty], [IntrNoMem, IntrConvergent]>;
>
> def int_amdgcn_s_dcache_wb :
> GCCBuiltin<"__builtin_amdgcn_s_dcache_wb">,
>
> Modified: llvm/trunk/lib/Target/AMDGPU/AsmParser/AMDGPUAsmParser.cpp
> URL:
> http://llvm.org/viewvc/llvm-project/llvm/trunk/lib/Target/AMDGPU/AsmParser/AMDGPUAsmParser.cpp?rev=263008&r1=263007&r2=263008&view=diff
>
> ==============================================================================
> --- llvm/trunk/lib/Target/AMDGPU/AsmParser/AMDGPUAsmParser.cpp (original)
> +++ llvm/trunk/lib/Target/AMDGPU/AsmParser/AMDGPUAsmParser.cpp Wed Mar 9
> 06:29:31 2016
> @@ -69,6 +69,10 @@ public:
> ImmTyTFE,
> ImmTyClamp,
> ImmTyOMod,
> + ImmTyDppCtrl,
> + ImmTyDppRowMask,
> + ImmTyDppBankMask,
> + ImmTyDppBoundCtrl,
> ImmTyDMask,
> ImmTyUNorm,
> ImmTyDA,
> @@ -144,7 +148,8 @@ public:
> bool defaultTokenHasSuffix() const {
> StringRef Token(Tok.Data, Tok.Length);
>
> - return Token.endswith("_e32") || Token.endswith("_e64");
> + return Token.endswith("_e32") || Token.endswith("_e64") ||
> + Token.endswith("_dpp");
> }
>
> bool isToken() const override {
> @@ -234,6 +239,18 @@ public:
> bool isSLC() const { return isImmTy(ImmTySLC); }
> bool isTFE() const { return isImmTy(ImmTyTFE); }
>
> + bool isBankMask() const {
> + return isImmTy(ImmTyDppBankMask);
> + }
> +
> + bool isRowMask() const {
> + return isImmTy(ImmTyDppRowMask);
> + }
> +
> + bool isBoundCtrl() const {
> + return isImmTy(ImmTyDppBoundCtrl);
> + }
> +
> void setModifiers(unsigned Mods) {
> assert(isReg() || (isImm() && Imm.Modifiers == 0));
> if (isReg())
> @@ -391,6 +408,7 @@ public:
> bool isMubufOffset() const;
> bool isSMRDOffset() const;
> bool isSMRDLiteralOffset() const;
> + bool isDPPCtrl() const;
> };
>
> class AMDGPUAsmParser : public MCTargetAsmParser {
> @@ -439,7 +457,6 @@ private:
> bool ParseSectionDirectiveHSARodataReadonlyAgent();
>
> public:
> -public:
> enum AMDGPUMatchResultTy {
> Match_PreferE32 = FIRST_TARGET_MATCH_RESULT_TY
> };
> @@ -538,6 +555,12 @@ public:
> void cvtMIMG(MCInst &Inst, const OperandVector &Operands);
> void cvtMIMGAtomic(MCInst &Inst, const OperandVector &Operands);
> OperandMatchResultTy parseVOP3OptionalOps(OperandVector &Operands);
> +
> + OperandMatchResultTy parseDPPCtrlOps(OperandVector &Operands);
> + OperandMatchResultTy parseDPPOptionalOps(OperandVector &Operands);
> + void cvtDPP_mod(MCInst &Inst, const OperandVector &Operands);
> + void cvtDPP_nomod(MCInst &Inst, const OperandVector &Operands);
> + void cvtDPP(MCInst &Inst, const OperandVector &Operands, bool HasMods);
> };
>
> struct OptionalOperand {
> @@ -1147,7 +1170,6 @@ bool AMDGPUAsmParser::ParseInstruction(P
> AMDGPUAsmParser::OperandMatchResultTy
> AMDGPUAsmParser::parseIntWithPrefix(const char *Prefix, int64_t &Int,
> int64_t Default) {
> -
> // We are at the end of the statement, and this is a default argument,
> so
> // use a default value.
> if (getLexer().is(AsmToken::EndOfStatement)) {
> @@ -1227,13 +1249,15 @@ AMDGPUAsmParser::parseNamedBit(const cha
>
> typedef std::map<enum AMDGPUOperand::ImmTy, unsigned> OptionalImmIndexMap;
>
> -void addOptionalImmOperand(MCInst& Inst, const OperandVector& Operands,
> OptionalImmIndexMap& OptionalIdx, enum AMDGPUOperand::ImmTy ImmT) {
> +void addOptionalImmOperand(MCInst& Inst, const OperandVector& Operands,
> + OptionalImmIndexMap& OptionalIdx,
> + enum AMDGPUOperand::ImmTy ImmT, int64_t
> Default = 0) {
> auto i = OptionalIdx.find(ImmT);
> if (i != OptionalIdx.end()) {
> unsigned Idx = i->second;
> ((AMDGPUOperand &)*Operands[Idx]).addImmOperands(Inst, 1);
> } else {
> - Inst.addOperand(MCOperand::createImm(0));
> + Inst.addOperand(MCOperand::createImm(Default));
> }
> }
>
> @@ -1896,6 +1920,152 @@ void AMDGPUAsmParser::cvtMIMGAtomic(MCIn
> addOptionalImmOperand(Inst, Operands, OptionalIdx,
> AMDGPUOperand::ImmTySLC);
> }
>
>
> +//===----------------------------------------------------------------------===//
> +// dpp
>
> +//===----------------------------------------------------------------------===//
> +
> +bool AMDGPUOperand::isDPPCtrl() const {
> + bool result = isImm() && getImmTy() == ImmTyDppCtrl &&
> isUInt<9>(getImm());
> + if (result) {
> + int64_t Imm = getImm();
> + return ((Imm >= 0x000) && (Imm <= 0x0ff)) ||
> + ((Imm >= 0x101) && (Imm <= 0x10f)) ||
> + ((Imm >= 0x111) && (Imm <= 0x11f)) ||
> + ((Imm >= 0x121) && (Imm <= 0x12f)) ||
> + (Imm == 0x130) ||
> + (Imm == 0x134) ||
> + (Imm == 0x138) ||
> + (Imm == 0x13c) ||
> + (Imm == 0x140) ||
> + (Imm == 0x141) ||
> + (Imm == 0x142) ||
> + (Imm == 0x143);
> + }
> + return false;
> +}
> +
> +AMDGPUAsmParser::OperandMatchResultTy
> +AMDGPUAsmParser::parseDPPCtrlOps(OperandVector &Operands) {
> + // ToDo: use same syntax as sp3 for dpp_ctrl
> + SMLoc S = Parser.getTok().getLoc();
> + StringRef Prefix;
> + int64_t Int;
> +
> + switch(getLexer().getKind()) {
> + default: return MatchOperand_NoMatch;
> + case AsmToken::Identifier: {
> + Prefix = Parser.getTok().getString();
> +
> + Parser.Lex();
> + if (getLexer().isNot(AsmToken::Colon))
> + return MatchOperand_ParseFail;
> +
> + Parser.Lex();
> + if (getLexer().isNot(AsmToken::Integer))
> + return MatchOperand_ParseFail;
> +
> + if (getParser().parseAbsoluteExpression(Int))
> + return MatchOperand_ParseFail;
> + break;
> + }
> + }
> +
> + if (Prefix.equals("row_shl")) {
> + Int |= 0x100;
> + } else if (Prefix.equals("row_shr")) {
> + Int |= 0x110;
> + } else if (Prefix.equals("row_ror")) {
> + Int |= 0x120;
> + } else if (Prefix.equals("wave_shl")) {
> + Int = 0x130;
> + } else if (Prefix.equals("wave_rol")) {
> + Int = 0x134;
> + } else if (Prefix.equals("wave_shr")) {
> + Int = 0x138;
> + } else if (Prefix.equals("wave_ror")) {
> + Int = 0x13C;
> + } else if (Prefix.equals("row_mirror")) {
> + Int = 0x140;
> + } else if (Prefix.equals("row_half_mirror")) {
> + Int = 0x141;
> + } else if (Prefix.equals("row_bcast")) {
> + if (Int == 15) {
> + Int = 0x142;
> + } else if (Int == 31) {
> + Int = 0x143;
> + }
> + } else if (!Prefix.equals("quad_perm")) {
> + return MatchOperand_NoMatch;
> + }
> + Operands.push_back(AMDGPUOperand::CreateImm(Int, S,
> +
> AMDGPUOperand::ImmTyDppCtrl));
> + return MatchOperand_Success;
> +}
> +
> +static const OptionalOperand DPPOptionalOps [] = {
> + {"row_mask", AMDGPUOperand::ImmTyDppRowMask, false, 0xf, nullptr},
> + {"bank_mask", AMDGPUOperand::ImmTyDppBankMask, false, 0xf, nullptr},
> + {"bound_ctrl", AMDGPUOperand::ImmTyDppBoundCtrl, false, -1, nullptr}
> +};
> +
> +AMDGPUAsmParser::OperandMatchResultTy
> +AMDGPUAsmParser::parseDPPOptionalOps(OperandVector &Operands) {
> + SMLoc S = Parser.getTok().getLoc();
> + OperandMatchResultTy Res = parseOptionalOps(DPPOptionalOps, Operands);
> + // XXX - sp3 use syntax "bound_ctrl:0" to indicate that bound_ctrl bit
> was set
> + if (Res == MatchOperand_Success) {
> + AMDGPUOperand &Op = ((AMDGPUOperand &)*Operands.back());
> + // If last operand was parsed as bound_ctrl we should replace it with
> correct value (1)
> + if (Op.isImmTy(AMDGPUOperand::ImmTyDppBoundCtrl)) {
> + Operands.pop_back();
> + Operands.push_back(
> + AMDGPUOperand::CreateImm(1, S, AMDGPUOperand::ImmTyDppBoundCtrl));
> + return MatchOperand_Success;
> + }
> + }
> + return Res;
> +}
> +
> +void AMDGPUAsmParser::cvtDPP_mod(MCInst &Inst, const OperandVector
> &Operands) {
> + cvtDPP(Inst, Operands, true);
> +}
> +
> +void AMDGPUAsmParser::cvtDPP_nomod(MCInst &Inst, const OperandVector
> &Operands) {
> + cvtDPP(Inst, Operands, false);
> +}
> +
> +void AMDGPUAsmParser::cvtDPP(MCInst &Inst, const OperandVector &Operands,
> + bool HasMods) {
> + OptionalImmIndexMap OptionalIdx;
> +
> + unsigned I = 1;
> + const MCInstrDesc &Desc = MII.get(Inst.getOpcode());
> + for (unsigned J = 0; J < Desc.getNumDefs(); ++J) {
> + ((AMDGPUOperand &)*Operands[I++]).addRegOperands(Inst, 1);
> + }
> +
> + for (unsigned E = Operands.size(); I != E; ++I) {
> + AMDGPUOperand &Op = ((AMDGPUOperand &)*Operands[I]);
> + // Add the register arguments
> + if (!HasMods && Op.isReg()) {
> + Op.addRegOperands(Inst, 1);
> + } else if (HasMods && Op.isRegOrImmWithInputMods()) {
> + Op.addRegOrImmWithInputModsOperands(Inst, 2);
> + } else if (Op.isDPPCtrl()) {
> + Op.addImmOperands(Inst, 1);
> + } else if (Op.isImm()) {
> + // Handle optional arguments
> + OptionalIdx[Op.getImmTy()] = I;
> + } else {
> + llvm_unreachable("Invalid operand type");
> + }
> + }
> +
> + // ToDo: fix default values for row_mask and bank_mask
> + addOptionalImmOperand(Inst, Operands, OptionalIdx,
> AMDGPUOperand::ImmTyDppRowMask, 0xf);
> + addOptionalImmOperand(Inst, Operands, OptionalIdx,
> AMDGPUOperand::ImmTyDppBankMask, 0xf);
> + addOptionalImmOperand(Inst, Operands, OptionalIdx,
> AMDGPUOperand::ImmTyDppBoundCtrl);
> +}
>
>
> /// Force static initialization.
>
> Modified: llvm/trunk/lib/Target/AMDGPU/InstPrinter/AMDGPUInstPrinter.cpp
> URL:
> http://llvm.org/viewvc/llvm-project/llvm/trunk/lib/Target/AMDGPU/InstPrinter/AMDGPUInstPrinter.cpp?rev=263008&r1=263007&r2=263008&view=diff
>
> ==============================================================================
> --- llvm/trunk/lib/Target/AMDGPU/InstPrinter/AMDGPUInstPrinter.cpp
> (original)
> +++ llvm/trunk/lib/Target/AMDGPU/InstPrinter/AMDGPUInstPrinter.cpp Wed
> Mar 9 06:29:31 2016
> @@ -28,6 +28,11 @@ void AMDGPUInstPrinter::printInst(const
> printAnnotation(OS, Annot);
> }
>
> +void AMDGPUInstPrinter::printU4ImmOperand(const MCInst *MI, unsigned OpNo,
> + raw_ostream &O) {
> + O << formatHex(MI->getOperand(OpNo).getImm() & 0xf);
> +}
> +
> void AMDGPUInstPrinter::printU8ImmOperand(const MCInst *MI, unsigned OpNo,
> raw_ostream &O) {
> O << formatHex(MI->getOperand(OpNo).getImm() & 0xff);
> @@ -43,6 +48,11 @@ void AMDGPUInstPrinter::printU32ImmOpera
> O << formatHex(MI->getOperand(OpNo).getImm() & 0xffffffff);
> }
>
> +void AMDGPUInstPrinter::printU4ImmDecOperand(const MCInst *MI, unsigned
> OpNo,
> + raw_ostream &O) {
> + O << formatDec(MI->getOperand(OpNo).getImm() & 0xf);
> +}
> +
> void AMDGPUInstPrinter::printU8ImmDecOperand(const MCInst *MI, unsigned
> OpNo,
> raw_ostream &O) {
> O << formatDec(MI->getOperand(OpNo).getImm() & 0xff);
> @@ -251,6 +261,8 @@ void AMDGPUInstPrinter::printVOPDst(cons
> raw_ostream &O) {
> if (MII.get(MI->getOpcode()).TSFlags & SIInstrFlags::VOP3)
> O << "_e64 ";
> + else if (MII.get(MI->getOpcode()).TSFlags & SIInstrFlags::DPP)
> + O << "_dpp ";
> else
> O << "_e32 ";
>
> @@ -388,6 +400,63 @@ void AMDGPUInstPrinter::printOperandAndM
> O << '|';
> }
>
> +
> +void AMDGPUInstPrinter::printDPPCtrlOperand(const MCInst *MI, unsigned
> OpNo,
> + raw_ostream &O) {
> + unsigned Imm = MI->getOperand(OpNo).getImm();
> + if ((Imm >= 0x000) && (Imm <= 0x0ff)) {
> + O << " quad_perm:";
> + printU8ImmDecOperand(MI, OpNo, O);
> + } else if ((Imm >= 0x101) && (Imm <= 0x10f)) {
> + O << " row_shl:";
> + printU4ImmDecOperand(MI, OpNo, O);
> + } else if ((Imm >= 0x111) && (Imm <= 0x11f)) {
> + O << " row_shr:";
> + printU4ImmDecOperand(MI, OpNo, O);
> + } else if ((Imm >= 0x121) && (Imm <= 0x12f)) {
> + O << " row_ror:";
> + printU4ImmDecOperand(MI, OpNo, O);
> + } else if (Imm == 0x130) {
> + O << " wave_shl:1";
> + } else if (Imm == 0x134) {
> + O << " wave_rol:1";
> + } else if (Imm == 0x138) {
> + O << " wave_shr:1";
> + } else if (Imm == 0x13c) {
> + O << " wave_ror:1";
> + } else if (Imm == 0x140) {
> + O << " row_mirror:1";
> + } else if (Imm == 0x141) {
> + O << " row_half_mirror:1";
> + } else if (Imm == 0x142) {
> + O << " row_bcast:15";
> + } else if (Imm == 0x143) {
> + O << " row_bcast:31";
> + } else {
> + llvm_unreachable("Invalid dpp_ctrl value");
> + }
> +}
> +
> +void AMDGPUInstPrinter::printRowMaskOperand(const MCInst *MI, unsigned
> OpNo,
> + raw_ostream &O) {
> + O << " row_mask:";
> + printU4ImmOperand(MI, OpNo, O);
> +}
> +
> +void AMDGPUInstPrinter::printBankMaskOperand(const MCInst *MI, unsigned
> OpNo,
> + raw_ostream &O) {
> + O << " bank_mask:";
> + printU4ImmOperand(MI, OpNo, O);
> +}
> +
> +void AMDGPUInstPrinter::printBoundCtrlOperand(const MCInst *MI, unsigned
> OpNo,
> + raw_ostream &O) {
> + unsigned Imm = MI->getOperand(OpNo).getImm();
> + if (Imm) {
> + O << " bound_ctrl:0"; // XXX - this syntax is used in sp3
> + }
> +}
> +
> void AMDGPUInstPrinter::printInterpSlot(const MCInst *MI, unsigned OpNum,
> raw_ostream &O) {
> unsigned Imm = MI->getOperand(OpNum).getImm();
>
> Modified: llvm/trunk/lib/Target/AMDGPU/InstPrinter/AMDGPUInstPrinter.h
> URL:
> http://llvm.org/viewvc/llvm-project/llvm/trunk/lib/Target/AMDGPU/InstPrinter/AMDGPUInstPrinter.h?rev=263008&r1=263007&r2=263008&view=diff
>
> ==============================================================================
> --- llvm/trunk/lib/Target/AMDGPU/InstPrinter/AMDGPUInstPrinter.h (original)
> +++ llvm/trunk/lib/Target/AMDGPU/InstPrinter/AMDGPUInstPrinter.h Wed Mar
> 9 06:29:31 2016
> @@ -33,8 +33,10 @@ public:
> const MCRegisterInfo &MRI);
>
> private:
> + void printU4ImmOperand(const MCInst *MI, unsigned OpNo, raw_ostream &O);
> void printU8ImmOperand(const MCInst *MI, unsigned OpNo, raw_ostream &O);
> void printU16ImmOperand(const MCInst *MI, unsigned OpNo, raw_ostream
> &O);
> + void printU4ImmDecOperand(const MCInst *MI, unsigned OpNo, raw_ostream
> &O);
> void printU8ImmDecOperand(const MCInst *MI, unsigned OpNo, raw_ostream
> &O);
> void printU16ImmDecOperand(const MCInst *MI, unsigned OpNo, raw_ostream
> &O);
> void printU32ImmOperand(const MCInst *MI, unsigned OpNo, raw_ostream
> &O);
> @@ -61,6 +63,10 @@ private:
> void printImmediate64(uint64_t I, raw_ostream &O);
> void printOperand(const MCInst *MI, unsigned OpNo, raw_ostream &O);
> void printOperandAndMods(const MCInst *MI, unsigned OpNo, raw_ostream
> &O);
> + void printDPPCtrlOperand(const MCInst *MI, unsigned OpNo, raw_ostream
> &O);
> + void printRowMaskOperand(const MCInst *MI, unsigned OpNo, raw_ostream
> &O);
> + void printBankMaskOperand(const MCInst *MI, unsigned OpNo, raw_ostream
> &O);
> + void printBoundCtrlOperand(const MCInst *MI, unsigned OpNo, raw_ostream
> &O);
> static void printInterpSlot(const MCInst *MI, unsigned OpNum,
> raw_ostream &O);
> void printMemOperand(const MCInst *MI, unsigned OpNo, raw_ostream &O);
> static void printIfSet(const MCInst *MI, unsigned OpNo, raw_ostream &O,
>
> Modified: llvm/trunk/lib/Target/AMDGPU/SIDefines.h
> URL:
> http://llvm.org/viewvc/llvm-project/llvm/trunk/lib/Target/AMDGPU/SIDefines.h?rev=263008&r1=263007&r2=263008&view=diff
>
> ==============================================================================
> --- llvm/trunk/lib/Target/AMDGPU/SIDefines.h (original)
> +++ llvm/trunk/lib/Target/AMDGPU/SIDefines.h Wed Mar 9 06:29:31 2016
> @@ -29,16 +29,17 @@ enum {
> VOP2 = 1 << 11,
> VOP3 = 1 << 12,
> VOPC = 1 << 13,
> + DPP = 1 << 14,
>
> - MUBUF = 1 << 14,
> - MTBUF = 1 << 15,
> - SMRD = 1 << 16,
> - DS = 1 << 17,
> - MIMG = 1 << 18,
> - FLAT = 1 << 19,
> - WQM = 1 << 20,
> - VGPRSpill = 1 << 21,
> - VOPAsmPrefer32Bit = 1 << 22
> + MUBUF = 1 << 15,
> + MTBUF = 1 << 16,
> + SMRD = 1 << 17,
> + DS = 1 << 18,
> + MIMG = 1 << 19,
> + FLAT = 1 << 20,
> + WQM = 1 << 21,
> + VGPRSpill = 1 << 22,
> + VOPAsmPrefer32Bit = 1 << 23
> };
> }
>
>
> Modified: llvm/trunk/lib/Target/AMDGPU/SIInstrFormats.td
> URL:
> http://llvm.org/viewvc/llvm-project/llvm/trunk/lib/Target/AMDGPU/SIInstrFormats.td?rev=263008&r1=263007&r2=263008&view=diff
>
> ==============================================================================
> --- llvm/trunk/lib/Target/AMDGPU/SIInstrFormats.td (original)
> +++ llvm/trunk/lib/Target/AMDGPU/SIInstrFormats.td Wed Mar 9 06:29:31 2016
> @@ -31,6 +31,7 @@ class InstSI <dag outs, dag ins, string
> field bits<1> VOP2 = 0;
> field bits<1> VOP3 = 0;
> field bits<1> VOPC = 0;
> + field bits<1> DPP = 0;
>
> field bits<1> MUBUF = 0;
> field bits<1> MTBUF = 0;
> @@ -63,16 +64,17 @@ class InstSI <dag outs, dag ins, string
> let TSFlags{11} = VOP2;
> let TSFlags{12} = VOP3;
> let TSFlags{13} = VOPC;
> + let TSFlags{14} = DPP;
>
> - let TSFlags{14} = MUBUF;
> - let TSFlags{15} = MTBUF;
> - let TSFlags{16} = SMRD;
> - let TSFlags{17} = DS;
> - let TSFlags{18} = MIMG;
> - let TSFlags{19} = FLAT;
> - let TSFlags{20} = WQM;
> - let TSFlags{21} = VGPRSpill;
> - let TSFlags{22} = VOPAsmPrefer32Bit;
> + let TSFlags{15} = MUBUF;
> + let TSFlags{16} = MTBUF;
> + let TSFlags{17} = SMRD;
> + let TSFlags{18} = DS;
> + let TSFlags{19} = MIMG;
> + let TSFlags{20} = FLAT;
> + let TSFlags{21} = WQM;
> + let TSFlags{22} = VGPRSpill;
> + let TSFlags{23} = VOPAsmPrefer32Bit;
>
> let SchedRW = [Write32Bit];
>
>
> Modified: llvm/trunk/lib/Target/AMDGPU/SIInstrInfo.td
> URL:
> http://llvm.org/viewvc/llvm-project/llvm/trunk/lib/Target/AMDGPU/SIInstrInfo.td?rev=263008&r1=263007&r2=263008&view=diff
>
> ==============================================================================
> --- llvm/trunk/lib/Target/AMDGPU/SIInstrInfo.td (original)
> +++ llvm/trunk/lib/Target/AMDGPU/SIInstrInfo.td Wed Mar 9 06:29:31 2016
> @@ -534,6 +534,22 @@ def SMRDLiteralOffsetMatchClass : SMRDOf
> "isSMRDLiteralOffset"
> >;
>
> +def DPPCtrlMatchClass : AsmOperandClass {
> + let Name = "DPPCtrl";
> + let PredicateMethod = "isDPPCtrl";
> + let ParserMethod = "parseDPPCtrlOps";
> + let RenderMethod = "addImmOperands";
> + let IsOptional = 0;
> +}
> +
> +class DPPOptionalMatchClass <string OpName>: AsmOperandClass {
> + let Name = "DPPOptional"#OpName;
> + let PredicateMethod = "is"#OpName;
> + let ParserMethod = "parseDPPOptionalOps";
> + let RenderMethod = "addImmOperands";
> + let IsOptional = 1;
> +}
> +
> class OptionalImmAsmOperand <string OpName> : AsmOperandClass {
> let Name = "Imm"#OpName;
> let PredicateMethod = "isImm";
> @@ -668,6 +684,26 @@ def lwe : NamedBitOperand<"LWE"> {
> let ParserMatchClass = NamedBitMatchClass<"LWE">;
> }
>
> +def dpp_ctrl : Operand <i32> {
> + let PrintMethod = "printDPPCtrlOperand";
> + let ParserMatchClass = DPPCtrlMatchClass;
> +}
> +
> +def row_mask : Operand <i32> {
> + let PrintMethod = "printRowMaskOperand";
> + let ParserMatchClass = DPPOptionalMatchClass<"RowMask">;
> +}
> +
> +def bank_mask : Operand <i32> {
> + let PrintMethod = "printBankMaskOperand";
> + let ParserMatchClass = DPPOptionalMatchClass<"BankMask">;
> +}
> +
> +def bound_ctrl : Operand <i1> {
> + let PrintMethod = "printBoundCtrlOperand";
> + let ParserMatchClass = DPPOptionalMatchClass<"BoundCtrl">;
> +}
> +
> } // End OperandType = "OPERAND_IMMEDIATE"
>
>
> @@ -1280,24 +1316,25 @@ class getInsDPP <RegisterClass Src0RC, R
> !if (!eq(HasModifiers, 1),
> // VOP1_DPP with modifiers
> (ins InputModsNoDefault:$src0_modifiers, Src0RC:$src0,
> - i32imm:$dpp_ctrl, i1imm:$bound_ctrl,
> - i32imm:$bank_mask, i32imm:$row_mask)
> + dpp_ctrl:$dpp_ctrl, row_mask:$row_mask,
> + bank_mask:$bank_mask, bound_ctrl:$bound_ctrl)
> /* else */,
> // VOP1_DPP without modifiers
> - (ins Src0RC:$src0, i32imm:$dpp_ctrl, i1imm:$bound_ctrl,
> - i32imm:$bank_mask, i32imm:$row_mask)
> + (ins Src0RC:$src0, dpp_ctrl:$dpp_ctrl, row_mask:$row_mask,
> + bank_mask:$bank_mask, bound_ctrl:$bound_ctrl)
> /* endif */)
> - /* NumSrcArgs == 2 */,
> + /* NumSrcArgs == 2 */,
> !if (!eq(HasModifiers, 1),
> // VOP2_DPP with modifiers
> (ins InputModsNoDefault:$src0_modifiers, Src0RC:$src0,
> - InputModsNoDefault:$src1_modifiers, Src1RC:$src1,
> - i32imm:$dpp_ctrl, i1imm:$bound_ctrl,
> - i32imm:$bank_mask, i32imm:$row_mask)
> + InputModsNoDefault:$src1_modifiers, Src1RC:$src1,
> + dpp_ctrl:$dpp_ctrl, row_mask:$row_mask,
> + bank_mask:$bank_mask, bound_ctrl:$bound_ctrl)
> /* else */,
> // VOP2_DPP without modifiers
> - (ins Src0RC:$src0, Src1RC:$src1, i32imm:$dpp_ctrl,
> i1imm:$bound_ctrl,
> - i32imm:$bank_mask, i32imm:$row_mask)
> + (ins Src0RC:$src0, Src1RC:$src1, dpp_ctrl:$dpp_ctrl,
> + row_mask:$row_mask, bank_mask:$bank_mask,
> + bound_ctrl:$bound_ctrl)
> /* endif */));
> }
>
> @@ -1338,8 +1375,8 @@ class getAsmDPP <bit HasDst, int NumSrcA
> " $src1_modifiers,"));
> string args = !if(!eq(HasModifiers, 0),
> getAsm32<0, NumSrcArgs, DstVT>.ret,
> - src0#src1);
> - string ret = " "#dst#args#", $dpp_ctrl, "#"$bound_ctrl, "#"$bank_mask,
> "#"$row_mask";
> + ", "#src0#src1);
> + string ret = dst#args#" $dpp_ctrl $row_mask $bank_mask $bound_ctrl";
> }
>
> class VOPProfile <list<ValueType> _ArgVT> {
> @@ -1351,7 +1388,7 @@ class VOPProfile <list<ValueType> _ArgVT
> field ValueType Src1VT = ArgVT[2];
> field ValueType Src2VT = ArgVT[3];
> field RegisterOperand DstRC = getVALUDstForVT<DstVT>.ret;
> - field RegisterClass DstRCDPP = !if(!eq(DstVT.Size, 64), VReg_64,
> VGPR_32);
> + field RegisterOperand DstRCDPP = getVALUDstForVT<DstVT>.ret;
> field RegisterOperand Src0RC32 = getVOPSrc0ForVT<Src0VT>.ret;
> field RegisterClass Src1RC32 = getVOPSrc1ForVT<Src1VT>.ret;
> field RegisterOperand Src0RC64 = getVOP3SrcForVT<Src0VT>.ret;
> @@ -1497,8 +1534,14 @@ def VOP_MAC : VOPProfile <[f32, f32, f32
> let Ins32 = (ins Src0RC32:$src0, Src1RC32:$src1, VGPR_32:$src2);
> let Ins64 = getIns64<Src0RC64, Src1RC64, RegisterOperand<VGPR_32>, 3,
> HasModifiers>.ret;
> + let InsDPP = (ins InputModsNoDefault:$src0_modifiers, Src0RC32:$src0,
> + InputModsNoDefault:$src1_modifiers, Src1RC32:$src1,
> + VGPR_32:$src2, // stub argument
> + dpp_ctrl:$dpp_ctrl, row_mask:$row_mask,
> + bank_mask:$bank_mask, bound_ctrl:$bound_ctrl);
> let Asm32 = getAsm32<1, 2, f32>.ret;
> let Asm64 = getAsm64<1, 2, HasModifiers, f32>.ret;
> + let AsmDPP = getAsmDPP<1, 2, HasModifiers, f32>.ret;
> }
> def VOP_F64_F64_F64_F64 : VOPProfile <[f64, f64, f64, f64]>;
> def VOP_I32_I32_I32_I32 : VOPProfile <[i32, i32, i32, i32]>;
> @@ -1607,9 +1650,9 @@ multiclass VOP1_m <vop1 op, string opNam
>
> class VOP1_DPP <vop1 op, string opName, VOPProfile p> :
> VOP1_DPPe <op.VI>,
> - VOP_DPP <p.OutsDPP, p.InsDPP, opName#p.AsmDPP, []> {
> + VOP_DPP <p.OutsDPP, p.InsDPP, opName#p.AsmDPP, [], p.HasModifiers> {
> let AssemblerPredicates = [isVI];
> - let src0_modifiers = !if(p.HasModifiers, ?, 0);
> + let src0_modifiers = !if(p.HasModifiers, ?, 0);
> let src1_modifiers = 0;
> }
>
> @@ -1667,6 +1710,14 @@ multiclass VOP2_m <vop2 op, string opNam
>
> }
>
> +class VOP2_DPP <vop2 op, string opName, VOPProfile p> :
> + VOP2_DPPe <op.VI>,
> + VOP_DPP <p.OutsDPP, p.InsDPP, opName#p.AsmDPP, [], p.HasModifiers> {
> + let AssemblerPredicates = [isVI];
> + let src0_modifiers = !if(p.HasModifiers, ?, 0);
> + let src1_modifiers = !if(p.HasModifiers, ?, 0);
> +}
> +
> class VOP3DisableFields <bit HasSrc1, bit HasSrc2, bit HasModifiers> {
>
> bits<2> src0_modifiers = !if(HasModifiers, ?, 0);
> @@ -1929,6 +1980,8 @@ multiclass VOP2_Helper <vop2 op, string
>
> defm _e64 : VOP3_2_m <op, p.Outs, p.Ins64, opName#p.Asm64, pat64,
> opName,
> revOp, p.HasModifiers>;
> +
> + def _dpp : VOP2_DPP <op, opName, p>;
> }
>
> multiclass VOP2Inst <vop2 op, string opName, VOPProfile P,
>
> Modified: llvm/trunk/lib/Target/AMDGPU/VIInstrFormats.td
> URL:
> http://llvm.org/viewvc/llvm-project/llvm/trunk/lib/Target/AMDGPU/VIInstrFormats.td?rev=263008&r1=263007&r2=263008&view=diff
>
> ==============================================================================
> --- llvm/trunk/lib/Target/AMDGPU/VIInstrFormats.td (original)
> +++ llvm/trunk/lib/Target/AMDGPU/VIInstrFormats.td Wed Mar 9 06:29:31 2016
> @@ -169,9 +169,12 @@ class VOP3be_vi <bits<10> op> : Enc64 {
> let Inst{63} = src2_modifiers{0};
> }
>
> -class VOP_DPP <dag outs, dag ins, string asm, list<dag> pattern> :
> +class VOP_DPP <dag outs, dag ins, string asm, list<dag> pattern, bit
> HasMods = 0> :
> VOPAnyCommon <outs, ins, asm, pattern> {
> + let DPP = 1;
> let Size = 8;
> +
> + let AsmMatchConverter = !if(!eq(HasMods,1), "cvtDPP_mod",
> "cvtDPP_nomod");
> }
>
> class VOP_DPPe : Enc64 {
> @@ -203,7 +206,7 @@ class VOP1_DPPe <bits<8> op> : VOP_DPPe
> let Inst{31-25} = 0x3f; //encoding
> }
>
> -class VOP2_DPPe <bits<6> op> : Enc32 {
> +class VOP2_DPPe <bits<6> op> : VOP_DPPe {
> bits<8> vdst;
> bits<8> src1;
>
>
> Modified: llvm/trunk/lib/Target/AMDGPU/VIInstructions.td
> URL:
> http://llvm.org/viewvc/llvm-project/llvm/trunk/lib/Target/AMDGPU/VIInstructions.td?rev=263008&r1=263007&r2=263008&view=diff
>
> ==============================================================================
> --- llvm/trunk/lib/Target/AMDGPU/VIInstructions.td (original)
> +++ llvm/trunk/lib/Target/AMDGPU/VIInstructions.td Wed Mar 9 06:29:31 2016
> @@ -121,10 +121,10 @@ def : Pat <
>
> //===----------------------------------------------------------------------===//
>
> def : Pat <
> - (int_amdgcn_mov_dpp i32:$src, imm:$dpp_ctrl, imm:$bound_ctrl,
> - imm:$bank_mask, imm:$row_mask),
> - (V_MOV_B32_dpp $src, (as_i32imm $dpp_ctrl), (as_i1imm $bound_ctrl),
> - (as_i32imm $bank_mask), (as_i32imm $row_mask))
> + (int_amdgcn_mov_dpp i32:$src, imm:$dpp_ctrl, imm:$row_mask,
> imm:$bank_mask,
> + imm:$bound_ctrl),
> + (V_MOV_B32_dpp $src, (as_i32imm $dpp_ctrl), (as_i32imm $row_mask),
> + (as_i32imm $bank_mask), (as_i1imm $bound_ctrl))
> >;
>
>
> //===----------------------------------------------------------------------===//
>
> Modified: llvm/trunk/test/CodeGen/AMDGPU/llvm.amdgcn.mov.dpp.ll
> URL:
> http://llvm.org/viewvc/llvm-project/llvm/trunk/test/CodeGen/AMDGPU/llvm.amdgcn.mov.dpp.ll?rev=263008&r1=263007&r2=263008&view=diff
>
> ==============================================================================
> --- llvm/trunk/test/CodeGen/AMDGPU/llvm.amdgcn.mov.dpp.ll (original)
> +++ llvm/trunk/test/CodeGen/AMDGPU/llvm.amdgcn.mov.dpp.ll Wed Mar 9
> 06:29:31 2016
> @@ -1,13 +1,13 @@
> ; RUN: llc -march=amdgcn -mcpu=tonga -verify-machineinstrs
> -show-mc-encoding < %s | FileCheck -check-prefix=VI %s
>
> ; VI-LABEL: {{^}}dpp_test:
> -; VI: v_mov_b32 v0, v0, 1, -1, 1, 1 ; encoding:
> [0xfa,0x02,0x00,0x7e,0x00,0x01,0x08,0x11]
> +; VI: v_mov_b32_dpp v0, v0 quad_perm:1 row_mask:0x1 bank_mask:0x1
> bound_ctrl:0 ; encoding: [0xfa,0x02,0x00,0x7e,0x00,0x01,0x08,0x11]
> define void @dpp_test(i32 addrspace(1)* %out, i32 %in) {
> - %tmp0 = call i32 @llvm.amdgcn.mov.dpp.i32(i32 %in, i32 1, i1 1, i32 1,
> i32 1) #0
> + %tmp0 = call i32 @llvm.amdgcn.mov.dpp.i32(i32 %in, i32 1, i32 1, i32 1,
> i1 1) #0
> store i32 %tmp0, i32 addrspace(1)* %out
> ret void
> }
>
> -declare i32 @llvm.amdgcn.mov.dpp.i32(i32, i32, i1, i32, i32) #0
> +declare i32 @llvm.amdgcn.mov.dpp.i32(i32, i32, i32, i32, i1) #0
>
> attributes #0 = { nounwind readnone convergent }
>
> Added: llvm/trunk/test/MC/AMDGPU/vop_dpp.s
> URL:
> http://llvm.org/viewvc/llvm-project/llvm/trunk/test/MC/AMDGPU/vop_dpp.s?rev=263008&view=auto
>
> ==============================================================================
> --- llvm/trunk/test/MC/AMDGPU/vop_dpp.s (added)
> +++ llvm/trunk/test/MC/AMDGPU/vop_dpp.s Wed Mar 9 06:29:31 2016
> @@ -0,0 +1,143 @@
> +// RUN: llvm-mc -arch=amdgcn -mcpu=tonga -show-encoding %s | FileCheck %s
> --check-prefix=GCN --check-prefix=CIVI --check-prefix=VI
> +// RUN: not llvm-mc -arch=amdgcn -show-encoding %s 2>&1 | FileCheck %s
> --check-prefix=NOSI --check-prefix=NOSICI
> +// RUN: not llvm-mc -arch=amdgcn -mcpu=SI -show-encoding %s 2>&1 |
> FileCheck %s --check-prefix=NOSI --check-prefix=NOSICI
> +// RUN: not llvm-mc -arch=amdgcn -mcpu=bonaire -show-encoding %s 2>&1 |
> FileCheck %s --check-prefix=NOSICI
> +
>
> +//===----------------------------------------------------------------------===//
> +// Check dpp_ctrl values
>
> +//===----------------------------------------------------------------------===//
> +
> +// NOSICI: error:
> +// VI: v_mov_b32_dpp v0, v0 quad_perm:37 row_mask:0xf bank_mask:0xf ;
> encoding: [0xfa,0x02,0x00,0x7e,0x00,0x25,0x00,0xff]
> +v_mov_b32 v0, v0 quad_perm:37
> +
> +// NOSICI: error:
> +// VI: v_mov_b32_dpp v0, v0 row_shl:1 row_mask:0xf bank_mask:0xf ;
> encoding: [0xfa,0x02,0x00,0x7e,0x00,0x01,0x01,0xff]
> +v_mov_b32 v0, v0 row_shl:1
> +
> +// NOSICI: error:
> +// VI: v_mov_b32_dpp v0, v0 row_shr:15 row_mask:0xf bank_mask:0xf ;
> encoding: [0xfa,0x02,0x00,0x7e,0x00,0x1f,0x01,0xff]
> +v_mov_b32 v0, v0 row_shr:0xf
> +
> +// NOSICI: error:
> +// VI: v_mov_b32_dpp v0, v0 row_ror:12 row_mask:0xf bank_mask:0xf ;
> encoding: [0xfa,0x02,0x00,0x7e,0x00,0x2c,0x01,0xff]
> +v_mov_b32 v0, v0 row_ror:0xc
> +
> +// NOSICI: error:
> +// VI: v_mov_b32_dpp v0, v0 wave_shl:1 row_mask:0xf bank_mask:0xf ;
> encoding: [0xfa,0x02,0x00,0x7e,0x00,0x30,0x01,0xff]
> +v_mov_b32 v0, v0 wave_shl:1
> +
> +// NOSICI: error:
> +// VI: v_mov_b32_dpp v0, v0 wave_rol:1 row_mask:0xf bank_mask:0xf ;
> encoding: [0xfa,0x02,0x00,0x7e,0x00,0x34,0x01,0xff]
> +v_mov_b32 v0, v0 wave_rol:1
> +
> +// NOSICI: error:
> +// VI: v_mov_b32_dpp v0, v0 wave_shr:1 row_mask:0xf bank_mask:0xf ;
> encoding: [0xfa,0x02,0x00,0x7e,0x00,0x38,0x01,0xff]
> +v_mov_b32 v0, v0 wave_shr:1
> +
> +// NOSICI: error:
> +// VI: v_mov_b32_dpp v0, v0 wave_ror:1 row_mask:0xf bank_mask:0xf ;
> encoding: [0xfa,0x02,0x00,0x7e,0x00,0x3c,0x01,0xff]
> +v_mov_b32 v0, v0 wave_ror:1
> +
> +// NOSICI: error:
> +// VI: v_mov_b32_dpp v0, v0 row_mirror:1 row_mask:0xf bank_mask:0xf ;
> encoding: [0xfa,0x02,0x00,0x7e,0x00,0x40,0x01,0xff]
> +v_mov_b32 v0, v0 row_mirror:1
> +
> +// NOSICI: error:
> +// VI: v_mov_b32_dpp v0, v0 row_half_mirror:1 row_mask:0xf bank_mask:0xf
> ; encoding: [0xfa,0x02,0x00,0x7e,0x00,0x41,0x01,0xff]
> +v_mov_b32 v0, v0 row_half_mirror:1
> +
> +// NOSICI: error:
> +// VI: v_mov_b32_dpp v0, v0 row_bcast:15 row_mask:0xf bank_mask:0xf ;
> encoding: [0xfa,0x02,0x00,0x7e,0x00,0x42,0x01,0xff]
> +v_mov_b32 v0, v0 row_bcast:15
> +
> +// NOSICI: error:
> +// VI: v_mov_b32_dpp v0, v0 row_bcast:31 row_mask:0xf bank_mask:0xf ;
> encoding: [0xfa,0x02,0x00,0x7e,0x00,0x43,0x01,0xff]
> +v_mov_b32 v0, v0 row_bcast:31
> +
>
> +//===----------------------------------------------------------------------===//
> +// Check optional fields
>
> +//===----------------------------------------------------------------------===//
> +
> +// NOSICI: error:
> +// VI: v_mov_b32_dpp v0, v0 quad_perm:1 row_mask:0xa bank_mask:0x1
> bound_ctrl:0 ; encoding: [0xfa,0x02,0x00,0x7e,0x00,0x01,0x08,0xa1]
> +v_mov_b32 v0, v0 quad_perm:1 row_mask:0xa bank_mask:0x1 bound_ctrl:0
> +
> +// NOSICI: error:
> +// VI: v_mov_b32_dpp v0, v0 quad_perm:1 row_mask:0xa bank_mask:0xf ;
> encoding: [0xfa,0x02,0x00,0x7e,0x00,0x01,0x00,0xaf]
> +v_mov_b32 v0, v0 quad_perm:1 row_mask:0xa
> +
> +// NOSICI: error:
> +// VI: v_mov_b32_dpp v0, v0 quad_perm:1 row_mask:0xf bank_mask:0x1 ;
> encoding: [0xfa,0x02,0x00,0x7e,0x00,0x01,0x00,0xf1]
> +v_mov_b32 v0, v0 quad_perm:1 bank_mask:0x1
> +
> +// NOSICI: error:
> +// VI: v_mov_b32_dpp v0, v0 quad_perm:1 row_mask:0xf bank_mask:0xf
> bound_ctrl:0 ; encoding: [0xfa,0x02,0x00,0x7e,0x00,0x01,0x08,0xff]
> +v_mov_b32 v0, v0 quad_perm:1 bound_ctrl:0
> +
> +// NOSICI: error:
> +// VI: v_mov_b32_dpp v0, v0 quad_perm:1 row_mask:0xa bank_mask:0x1 ;
> encoding: [0xfa,0x02,0x00,0x7e,0x00,0x01,0x00,0xa1]
> +v_mov_b32 v0, v0 quad_perm:1 row_mask:0xa bank_mask:0x1
> +
> +// NOSICI: error:
> +// VI: v_mov_b32_dpp v0, v0 quad_perm:1 row_mask:0xa bank_mask:0xf
> bound_ctrl:0 ; encoding: [0xfa,0x02,0x00,0x7e,0x00,0x01,0x08,0xaf]
> +v_mov_b32 v0, v0 quad_perm:1 row_mask:0xa bound_ctrl:0
> +
> +// NOSICI: error:
> +// VI: v_mov_b32_dpp v0, v0 quad_perm:1 row_mask:0xf bank_mask:0x1
> bound_ctrl:0 ; encoding: [0xfa,0x02,0x00,0x7e,0x00,0x01,0x08,0xf1]
> +v_mov_b32 v0, v0 quad_perm:1 bank_mask:0x1 bound_ctrl:0
> +
>
> +//===----------------------------------------------------------------------===//
> +// Check VOP1 opcodes
>
> +//===----------------------------------------------------------------------===//
> +// ToDo: v_nop
> +
> +// NOSICI: error:
> +// VI: v_cvt_u32_f32_dpp v0, v0 row_shl:1 row_mask:0xa bank_mask:0x1
> bound_ctrl:0 ; encoding: [0xfa,0x0e,0x00,0x7e,0x00,0x01,0x09,0xa1]
> +v_cvt_u32_f32 v0, v0 row_shl:1 row_mask:0xa bank_mask:0x1 bound_ctrl:0
> +
> +// NOSICI: error:
> +// VI: v_fract_f32_dpp v0, v0 row_shl:1 row_mask:0xa bank_mask:0x1
> bound_ctrl:0 ; encoding: [0xfa,0x36,0x00,0x7e,0x00,0x01,0x09,0xa1]
> +v_fract_f32 v0, v0 row_shl:1 row_mask:0xa bank_mask:0x1 bound_ctrl:0
> +
> +// NOSICI: error:
> +// VI: v_sin_f32_dpp v0, v0 row_shl:1 row_mask:0xa bank_mask:0x1
> bound_ctrl:0 ; encoding: [0xfa,0x52,0x00,0x7e,0x00,0x01,0x09,0xa1]
> +v_sin_f32 v0, v0 row_shl:1 row_mask:0xa bank_mask:0x1 bound_ctrl:0
> +
>
> +//===----------------------------------------------------------------------===//
> +// Check VOP2 opcodes
>
> +//===----------------------------------------------------------------------===//
> +// ToDo: VOP2bInst instructions: v_add_u32, v_sub_u32 ... (vcc and
> ApplyMnemonic in AsmMatcherEmitter.cpp)
> +// ToDo: v_mac_f32 (VOP_MAC)
> +
> +// NOSICI: error:
> +// VI: v_add_f32_dpp v0, v0, v0 row_shl:1 row_mask:0xa bank_mask:0x1
> bound_ctrl:0 ; encoding: [0xfa,0x00,0x00,0x02,0x00,0x01,0x09,0xa1]
> +v_add_f32 v0, v0, v0 row_shl:1 row_mask:0xa bank_mask:0x1 bound_ctrl:0
> +
> +// NOSICI: error:
> +// VI: v_min_f32_dpp v0, v0, v0 row_shl:1 row_mask:0xa bank_mask:0x1
> bound_ctrl:0 ; encoding: [0xfa,0x00,0x00,0x14,0x00,0x01,0x09,0xa1]
> +v_min_f32 v0, v0, v0 row_shl:1 row_mask:0xa bank_mask:0x1 bound_ctrl:0
> +
> +// NOSICI: error:
> +// VI: v_and_b32_dpp v0, v0, v0 row_shl:1 row_mask:0xa bank_mask:0x1
> bound_ctrl:0 ; encoding: [0xfa,0x00,0x00,0x26,0x00,0x01,0x09,0xa1]
> +v_and_b32 v0, v0, v0 row_shl:1 row_mask:0xa bank_mask:0x1 bound_ctrl:0
> +
>
> +//===----------------------------------------------------------------------===//
> +// Check modifiers
>
> +//===----------------------------------------------------------------------===//
> +
> +// NOSICI: error:
> +// VI: v_add_f32_dpp v0, -v0, v0 row_shl:1 row_mask:0xa bank_mask:0x1
> bound_ctrl:0 ; encoding: [0xfa,0x00,0x00,0x02,0x00,0x01,0x19,0xa1]
> +v_add_f32 v0, -v0, v0 row_shl:1 row_mask:0xa bank_mask:0x1 bound_ctrl:0
> +
> +// NOSICI: error:
> +// VI: v_add_f32_dpp v0, v0, |v0| row_shl:1 row_mask:0xa bank_mask:0x1
> bound_ctrl:0 ; encoding: [0xfa,0x00,0x00,0x02,0x00,0x01,0x89,0xa1]
> +v_add_f32 v0, v0, |v0| row_shl:1 row_mask:0xa bank_mask:0x1 bound_ctrl:0
> +
> +// NOSICI: error:
> +// VI: v_add_f32_dpp v0, -v0, |v0| row_shl:1 row_mask:0xa bank_mask:0x1
> bound_ctrl:0 ; encoding: [0xfa,0x00,0x00,0x02,0x00,0x01,0x99,0xa1]
> +v_add_f32 v0, -v0, |v0| row_shl:1 row_mask:0xa bank_mask:0x1 bound_ctrl:0
> +
> +// NOSICI: error:
> +// VI: v_add_f32_dpp v0, |v0|, -v0 row_shl:1 row_mask:0xa bank_mask:0x1
> bound_ctrl:0 ; encoding: [0xfa,0x00,0x00,0x02,0x00,0x01,0x69,0xa1]
> +v_add_f32 v0, |v0|, -v0 row_shl:1 row_mask:0xa bank_mask:0x1 bound_ctrl:0
>
>
> _______________________________________________
> llvm-commits mailing list
> llvm-commits at lists.llvm.org
> http://lists.llvm.org/cgi-bin/mailman/listinfo/llvm-commits
>
--
Teresa Johnson | Software Engineer | tejohnson at google.com | 408-460-2413
-------------- next part --------------
An HTML attachment was scrubbed...
URL: <http://lists.llvm.org/pipermail/llvm-commits/attachments/20160309/28ade5c5/attachment.html>
More information about the llvm-commits
mailing list