[AMDGPU] Assembler: Support DPP instructions.

[SamWot]

Summary

Added support for DPP instruction in assembler for VOP1 and VOP2 formats. Added tests for assembler.

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

[nhaustov]

lib/Target/AMDGPU/AsmParser/AMDGPUAsmParser.cpp

242

Use IsImmTy(...) here and below.

lib/Target/AMDGPU/VIInstrFormats.td

174

Missing space

[SamWot]

Fixed IsImmTy and space

[SamWot]

Submitted r263008

include/llvm/IR/IntrinsicsAMDGPU.td

 def int_amdgcn_buffer_wbinvl1_vol :
   GCCBuiltin<"__builtin_amdgcn_buffer_wbinvl1_vol">,
   Intrinsic<[], [], []>;
 
 //===----------------------------------------------------------------------===//
 // 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">,
   Intrinsic<[], [], []>;
 
 def int_amdgcn_s_dcache_wb_vol :
   GCCBuiltin<"__builtin_amdgcn_s_dcache_wb_vol">,
   Intrinsic<[], [], []>;

lib/Target/AMDGPU/AsmParser/AMDGPUAsmParser.cpp

enum ImmTy {
     ImmTyDSOffset1,
     ImmTyGDS,
     ImmTyOffset,
     ImmTyGLC,
     ImmTySLC,
     ImmTyTFE,
     ImmTyClamp,
     ImmTyOMod,
+    ImmTyDppCtrl,
+    ImmTyDppRowMask,
+    ImmTyDppBankMask,
+    ImmTyDppBoundCtrl,
     ImmTyDMask,
     ImmTyUNorm,
     ImmTyDA,
     ImmTyR128,
     ImmTyLWE,
   };
 
   struct TokOp {
else {
       assert(isExpr());
       Inst.addOperand(MCOperand::createExpr(Expr));
     }
   }
 
   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 {
     return Kind == Token;
   }
 
   bool isImm() const override {
     return Kind == Immediate;
bool isMod() const {
     return isClamp() || isOMod();
   }
 
   bool isGDS() const { return isImmTy(ImmTyGDS); }
   bool isGLC() const { return isImmTy(ImmTyGLC); }
   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())
       Reg.Modifiers = Mods;
     else
       Imm.Modifiers = Mods;
   }
 
public:
   }
 
   bool isDSOffset() const;
   bool isDSOffset01() const;
   bool isSWaitCnt() const;
   bool isMubufOffset() const;
   bool isSMRDOffset() const;
   bool isSMRDLiteralOffset() const;
+  bool isDPPCtrl() const;
 };
 
 class AMDGPUAsmParser : public MCTargetAsmParser {
   const MCInstrInfo &MII;
   MCAsmParser &Parser;
 
   unsigned ForcedEncodingSize;
 
private:
   bool ParseDirectiveAMDGPUHsaKernel();
   bool ParseDirectiveAMDGPUHsaModuleGlobal();
   bool ParseDirectiveAMDGPUHsaProgramGlobal();
   bool ParseSectionDirectiveHSADataGlobalAgent();
   bool ParseSectionDirectiveHSADataGlobalProgram();
   bool ParseSectionDirectiveHSARodataReadonlyAgent();
 
 public:
-public:
   enum AMDGPUMatchResultTy {
     Match_PreferE32 = FIRST_TARGET_MATCH_RESULT_TY
   };
 
   AMDGPUAsmParser(const MCSubtargetInfo &STI, MCAsmParser &_Parser,
                const MCInstrInfo &MII,
                const MCTargetOptions &Options)
       : MCTargetAsmParser(Options, STI), MII(MII), Parser(_Parser),
public:
   void cvtId(MCInst &Inst, const OperandVector &Operands);
   void cvtVOP3_2_mod(MCInst &Inst, const OperandVector &Operands);
   void cvtVOP3_2_nomod(MCInst &Inst, const OperandVector &Operands);
   void cvtVOP3_only(MCInst &Inst, const OperandVector &Operands);
   void cvtVOP3(MCInst &Inst, const OperandVector &Operands);
 
   void cvtMIMG(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 {
   const char *Name;
   AMDGPUOperand::ImmTy Type;
   bool IsBit;
   int64_t Default;
   bool (*ConvertResult)(int64_t&);

 
 //===----------------------------------------------------------------------===//
 // Utility functions
 //===----------------------------------------------------------------------===//
 
 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)) {
     Int = Default;
     return MatchOperand_Success;
   }
 
   switch(getLexer().getKind()) {
AMDGPUAsmParser::parseNamedBit(const char *Name, OperandVector &Operands,
   }
 
   Operands.push_back(AMDGPUOperand::CreateImm(Bit, S, ImmTy));
   return MatchOperand_Success;
 }
 
 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));
   }
 }
 
 static bool operandsHasOptionalOp(const OperandVector &Operands,
                                   const OptionalOperand &OOp) {
   for (unsigned i = 0; i < Operands.size(); i++) {
     const AMDGPUOperand &ParsedOp = ((const AMDGPUOperand &)*Operands[i]);
     if ((ParsedOp.isImm() && ParsedOp.getImmTy() == OOp.Type) ||
void AMDGPUAsmParser::cvtMIMG(MCInst &Inst, const OperandVector &Operands) {
   addOptionalImmOperand(Inst, Operands, OptionalIdx, AMDGPUOperand::ImmTyGLC);
   addOptionalImmOperand(Inst, Operands, OptionalIdx, AMDGPUOperand::ImmTyDA);
   addOptionalImmOperand(Inst, Operands, OptionalIdx, AMDGPUOperand::ImmTyR128);
   addOptionalImmOperand(Inst, Operands, OptionalIdx, AMDGPUOperand::ImmTyTFE);
   addOptionalImmOperand(Inst, Operands, OptionalIdx, AMDGPUOperand::ImmTyLWE);
   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.
 extern "C" void LLVMInitializeAMDGPUAsmParser() {
   RegisterMCAsmParser<AMDGPUAsmParser> A(TheAMDGPUTarget);
   RegisterMCAsmParser<AMDGPUAsmParser> B(TheGCNTarget);
 }
 
 #define GET_REGISTER_MATCHER
 #define GET_MATCHER_IMPLEMENTATION
 #include "AMDGPUGenAsmMatcher.inc"

lib/Target/AMDGPU/InstPrinter/AMDGPUInstPrinter.h

public:
   static const char *getRegisterName(unsigned RegNo);
 
   void printInst(const MCInst *MI, raw_ostream &O, StringRef Annot,
                  const MCSubtargetInfo &STI) override;
   static void printRegOperand(unsigned RegNo, raw_ostream &O,
                               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);
   void printNamedBit(const MCInst* MI, unsigned OpNo, raw_ostream& O, const char* BitName);
   void printOffen(const MCInst *MI, unsigned OpNo, raw_ostream &O);
   void printIdxen(const MCInst *MI, unsigned OpNo, raw_ostream &O);
   void printAddr64(const MCInst *MI, unsigned OpNo, raw_ostream &O);
   void printMBUFOffset(const MCInst *MI, unsigned OpNo, raw_ostream &O);
private:
   void printR128(const MCInst *MI, unsigned OpNo, raw_ostream &O);
   void printLWE(const MCInst *MI, unsigned OpNo, raw_ostream &O);
   void printRegOperand(unsigned RegNo, raw_ostream &O);
   void printVOPDst(const MCInst *MI, unsigned OpNo, raw_ostream &O);
   void printImmediate32(uint32_t I, raw_ostream &O);
   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,
                          StringRef Asm, StringRef Default = "");
   static void printAbs(const MCInst *MI, unsigned OpNo, raw_ostream &O);
   static void printClamp(const MCInst *MI, unsigned OpNo, raw_ostream &O);
   static void printClampSI(const MCInst *MI, unsigned OpNo, raw_ostream &O);
   static void printOModSI(const MCInst *MI, unsigned OpNo, raw_ostream &O);

lib/Target/AMDGPU/InstPrinter/AMDGPUInstPrinter.cpp

 void AMDGPUInstPrinter::printInst(const MCInst *MI, raw_ostream &OS,
                                   StringRef Annot, const MCSubtargetInfo &STI) {
   OS.flush();
   printInstruction(MI, OS);
 
   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);
 }
 
 void AMDGPUInstPrinter::printU16ImmOperand(const MCInst *MI, unsigned OpNo,
                                            raw_ostream &O) {
   O << formatHex(MI->getOperand(OpNo).getImm() & 0xffff);
 }
 
 void AMDGPUInstPrinter::printU32ImmOperand(const MCInst *MI, unsigned OpNo,
                                            raw_ostream &O) {
   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);
 }
 
 void AMDGPUInstPrinter::printU16ImmDecOperand(const MCInst *MI, unsigned OpNo,
                                               raw_ostream &O) {
   O << formatDec(MI->getOperand(OpNo).getImm() & 0xffff);
void AMDGPUInstPrinter::printRegOperand(unsigned reg, raw_ostream &O,
 
   O << Type << '[' << RegIdx << ':' << (RegIdx + NumRegs - 1) << ']';
 }
 
 void AMDGPUInstPrinter::printVOPDst(const MCInst *MI, unsigned OpNo,
                                     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 ";
 
   printOperand(MI, OpNo, O);
 }
 
 void AMDGPUInstPrinter::printImmediate32(uint32_t Imm, raw_ostream &O) {
   int32_t SImm = static_cast<int32_t>(Imm);
if (InputModifiers & SISrcMods::NEG)
     O << '-';
   if (InputModifiers & SISrcMods::ABS)
     O << '|';
   printOperand(MI, OpNo + 1, O);
   if (InputModifiers & SISrcMods::ABS)
     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();
 
   if (Imm == 2) {
     O << "P0";
   } else if (Imm == 1) {
     O << "P20";

lib/Target/AMDGPU/SIDefines.h

enum {
   SOPC = 1 << 7,
   SOPK = 1 << 8,
   SOPP = 1 << 9,
 
   VOP1 = 1 << 10,
   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
 };
 }
 
 namespace llvm {
 namespace AMDGPU {
   enum OperandType {
     /// Operand with register or 32-bit immediate
     OPERAND_REG_IMM32 = llvm::MCOI::OPERAND_FIRST_TARGET,

lib/Target/AMDGPU/SIInstrFormats.td

class InstSI <dag outs, dag ins, string asm, list<dag> pattern> :
   field bits<1> SOPC = 0;
   field bits<1> SOPK = 0;
   field bits<1> SOPP = 0;
 
   field bits<1> VOP1 = 0;
   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;
   field bits<1> SMRD = 0;
   field bits<1> DS = 0;
   field bits<1> MIMG = 0;
   field bits<1> FLAT = 0;
   field bits<1> WQM = 0;
class InstSI <dag outs, dag ins, string asm, list<dag> pattern> :
   let TSFlags{7} = SOPC;
   let TSFlags{8} = SOPK;
   let TSFlags{9} = SOPP;
 
   let TSFlags{10} = VOP1;
   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];
 
   field bits<1> DisableSIDecoder = 0;
   field bits<1> DisableVIDecoder = 0;
   field bits<1> DisableDecoder = 0;
 
   let isAsmParserOnly = !if(!eq(DisableDecoder{0}, {0}), 0, 1);

lib/Target/AMDGPU/SIInstrInfo.td

class SMRDOffsetBaseMatchClass <string predicate> : AsmOperandClass {
   let RenderMethod = "addImmOperands";
 }
 
 def SMRDOffsetMatchClass : SMRDOffsetBaseMatchClass <"isSMRDOffset">;
 def SMRDLiteralOffsetMatchClass : SMRDOffsetBaseMatchClass <
   "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";
   let RenderMethod = "addImmOperands";
   let IsOptional = 1;
 }
 
 def DMaskMatchClass : AsmOperandClass {

 def r128 : NamedBitOperand<"R128"> {
   let ParserMatchClass = NamedBitMatchClass<"R128">;
 }
 
 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"
 
 
 def VOPDstS64 : VOPDstOperand <SReg_64>;
 
 //===----------------------------------------------------------------------===//
 // Complex patterns
 //===----------------------------------------------------------------------===//

 
 class getInsDPP <RegisterClass Src0RC, RegisterClass Src1RC, int NumSrcArgs,
                                                              bit HasModifiers> {
 
   dag ret = !if (!eq(NumSrcArgs, 1),
               !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 */,
               !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)
+                      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 */));
 }
 
 // Returns the assembly string for the inputs and outputs of a VOP[12C]
 // instruction.  This does not add the _e32 suffix, so it can be reused
 // by getAsm64.
 class getAsm32 <bit HasDst, int NumSrcArgs, ValueType DstVT = i32> {
   string dst = !if(!eq(DstVT.Size, 1), "$sdst", "$vdst"); // use $sdst for VOPC

 class getAsmDPP <bit HasDst, int NumSrcArgs, bit HasModifiers, ValueType DstVT = i32> {
   string dst = !if(!eq(DstVT.Size, 1), "$sdst", "$vdst"); // use $sdst for VOPC
   string src0 = !if(!eq(NumSrcArgs, 1), "$src0_modifiers", "$src0_modifiers,");
   string src1 = !if(!eq(NumSrcArgs, 1), "",
                    !if(!eq(NumSrcArgs, 2), " $src1_modifiers",
                                            " $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> {
 
   field list<ValueType> ArgVT = _ArgVT;
 
   field ValueType DstVT = ArgVT[0];
   field ValueType Src0VT = ArgVT[1];
   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;
   field RegisterOperand Src1RC64 = getVOP3SrcForVT<Src1VT>.ret;
   field RegisterOperand Src2RC64 = getVOP3SrcForVT<Src2VT>.ret;
   field RegisterClass Src0DPP = getDPPSrcForVT<Src0VT>.ret;
   field RegisterClass Src1DPP = getDPPSrcForVT<Src1VT>.ret;
 

 def VOP_MADK : VOPProfile <[f32, f32, f32, f32]> {
   field dag Ins = (ins VCSrc_32:$src0, VGPR_32:$vsrc1, u32imm:$src2);
   field string Asm = "$vdst, $src0, $vsrc1, $src2";
 }
 def VOP_MAC : VOPProfile <[f32, 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]>;
 def VOP_I64_I32_I32_I64 : VOPProfile <[i64, i32, i32, i64]>;
 
 // This class is used only with VOPC instructions. Use $sdst for out operand
 class SIInstAlias <string asm, Instruction inst, VOPProfile p> :
     InstAlias <asm, (inst)>, PredicateControl {
multiclass VOP1_m <vop1 op, string opName, VOPProfile p, list<dag> pattern,
   def _si : VOP1_Real_si <opName, op, p.Outs, p.Ins32, asm>;
 
   def _vi : VOP1_Real_vi <opName, op, p.Outs, p.Ins32, asm>;
 
 }
 
 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;
 }
 
 multiclass VOP1SI_m <vop1 op, string opName, VOPProfile p, list<dag> pattern,
                      string asm = opName#p.Asm32> {
 
   def "" : VOP1_Pseudo <p.Outs, p.Ins32, pattern, opName>;
 
def "" : VOP2_Pseudo <p.Outs32, p.Ins32, pattern, opName>,
            VOP2_REV<revOp#"_e32", !eq(revOp, opName)>;
 
   def _si : VOP2_Real_si <opName, op, p.Outs32, p.Ins32, p.Asm32>;
 
   def _vi : VOP2_Real_vi <opName, op, p.Outs32, p.Ins32, p.Asm32>;
 
 }
 
+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);
   bits<2> src1_modifiers = !if(HasModifiers, !if(HasSrc1, ?, 0), 0);
   bits<2> src2_modifiers = !if(HasModifiers, !if(HasSrc2, ?, 0), 0);
   bits<2> omod = !if(HasModifiers, ?, 0);
   bits<1> clamp = !if(HasModifiers, ?, 0);
   bits<9> src1 = !if(HasSrc1, ?, 0);

 
 multiclass VOP2_Helper <vop2 op, string opName, VOPProfile p, list<dag> pat32,
                         list<dag> pat64, string revOp> {
 
   defm _e32 : VOP2_m <op, opName, p, pat32, revOp>;
 
   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,
                      SDPatternOperator node = null_frag,
                      string revOp = opName> : VOP2_Helper <
   op, opName, P, [],
   !if(P.HasModifiers,
       [(set P.DstVT:$vdst,

lib/Target/AMDGPU/VIInstrFormats.td

class VOP3be_vi <bits<10> op> : Enc64 {
   let Inst{49-41} = src1;
   let Inst{58-50} = src2;
   let Inst{60-59} = omod;
   let Inst{61} = src0_modifiers{0};
   let Inst{62} = src1_modifiers{0};
   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 {
   bits<2> src0_modifiers;
   bits<8> src0;
   bits<2> src1_modifiers;
   bits<9> dpp_ctrl;
   bits<1> bound_ctrl;
class VOP1_DPPe <bits<8> op> : VOP_DPPe {
   bits<8> vdst;
 
   let Inst{8-0} = 0xfa; // dpp
   let Inst{16-9} = op;
   let Inst{24-17} = vdst;
   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;
 
   let Inst{8-0} = 0xfa; //dpp
   let Inst{16-9} = src1;
   let Inst{24-17} = vdst;
   let Inst{30-25} = op;
   let Inst{31} = 0x0; //encoding

lib/Target/AMDGPU/VIInstructions.td

def : Pat <
   (S_BUFFER_LOAD_DWORD_IMM $sbase, (as_i32imm $offset))
 >;
 
 //===----------------------------------------------------------------------===//
 // DPP Patterns
 //===----------------------------------------------------------------------===//
 
 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))
 >;
 
 //===----------------------------------------------------------------------===//
 // Misc Patterns
 //===----------------------------------------------------------------------===//
 
 def : Pat <
   (i64 (readcyclecounter)),

test/CodeGen/AMDGPU/llvm.amdgcn.mov.dpp.ll

 ; 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 }

test/MC/AMDGPU/vop_dpp.s

+// 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