[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
llvm-svn: 263008

diff --git a/llvm/include/llvm/IR/IntrinsicsAMDGPU.td b/llvm/include/llvm/IR/IntrinsicsAMDGPU.td
index f32641b..b8479a5 100644 (file)
--- a/llvm/include/llvm/IR/IntrinsicsAMDGPU.td
+++ b/llvm/include/llvm/IR/IntrinsicsAMDGPU.td
@@ -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">,

diff --git a/llvm/lib/Target/AMDGPU/AsmParser/AMDGPUAsmParser.cpp b/llvm/lib/Target/AMDGPU/AsmParser/AMDGPUAsmParser.cpp
index a4d2909..66435a2 100644 (file)
--- a/llvm/lib/Target/AMDGPU/AsmParser/AMDGPUAsmParser.cpp
+++ b/llvm/lib/Target/AMDGPU/AsmParser/AMDGPUAsmParser.cpp
@@ -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(ParseInstructionInfo &Info,
 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 char *Name, OperandVector &Operands,
 
 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(MCInst &Inst, const OperandVector &Operands)
   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.

diff --git a/llvm/lib/Target/AMDGPU/InstPrinter/AMDGPUInstPrinter.cpp b/llvm/lib/Target/AMDGPU/InstPrinter/AMDGPUInstPrinter.cpp
index e78886b..fafa7e5 100644 (file)
--- a/llvm/lib/Target/AMDGPU/InstPrinter/AMDGPUInstPrinter.cpp
+++ b/llvm/lib/Target/AMDGPU/InstPrinter/AMDGPUInstPrinter.cpp
@@ -28,6 +28,11 @@ void AMDGPUInstPrinter::printInst(const MCInst *MI, raw_ostream &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);
@@ -43,6 +48,11 @@ void AMDGPUInstPrinter::printU32ImmOperand(const MCInst *MI, unsigned OpNo,
   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(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 ";
 
@@ -388,6 +400,63 @@ void AMDGPUInstPrinter::printOperandAndMods(const MCInst *MI, unsigned OpNo,
     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();

diff --git a/llvm/lib/Target/AMDGPU/InstPrinter/AMDGPUInstPrinter.h b/llvm/lib/Target/AMDGPU/InstPrinter/AMDGPUInstPrinter.h
index 13335a9..7a1e86b 100644 (file)
--- a/llvm/lib/Target/AMDGPU/InstPrinter/AMDGPUInstPrinter.h
+++ b/llvm/lib/Target/AMDGPU/InstPrinter/AMDGPUInstPrinter.h
@@ -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,

diff --git a/llvm/lib/Target/AMDGPU/SIDefines.h b/llvm/lib/Target/AMDGPU/SIDefines.h
index aa1e352..f3e74ca 100644 (file)
--- a/llvm/lib/Target/AMDGPU/SIDefines.h
+++ b/llvm/lib/Target/AMDGPU/SIDefines.h
@@ -29,16 +29,17 @@ enum {
   VOP2 = 1 << 11,
   VOP3 = 1 << 12,
   VOPC = 1 << 13,
-
-  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
+  DPP = 1 << 14,
+
+  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
 };
 }
 

diff --git a/llvm/lib/Target/AMDGPU/SIInstrFormats.td b/llvm/lib/Target/AMDGPU/SIInstrFormats.td
index a6bed78..ea7b6a1 100644 (file)
--- a/llvm/lib/Target/AMDGPU/SIInstrFormats.td
+++ b/llvm/lib/Target/AMDGPU/SIInstrFormats.td
@@ -31,6 +31,7 @@ class InstSI <dag outs, dag ins, string asm, list<dag> pattern> :
   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 asm, list<dag> pattern> :
   let TSFlags{11} = VOP2;
   let TSFlags{12} = VOP3;
   let TSFlags{13} = VOPC;
-
-  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{14} = DPP;
+
+  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];
 

diff --git a/llvm/lib/Target/AMDGPU/SIInstrInfo.td b/llvm/lib/Target/AMDGPU/SIInstrInfo.td
index 4148a68..6c00d85 100644 (file)
--- a/llvm/lib/Target/AMDGPU/SIInstrInfo.td
+++ b/llvm/lib/Target/AMDGPU/SIInstrInfo.td
@@ -534,6 +534,22 @@ 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";
@@ -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, RegisterClass Src1RC, int NumSrcArgs,
               !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 NumSrcArgs, bit HasModifiers, ValueType DstVT =
                                            " $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, 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 opName, VOPProfile p, list<dag> pattern,
 
 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 opName, VOPProfile p, list <dag> pattern,
 
 }
 
+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 opName, VOPProfile p, list<dag> pat32,
 
   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,

diff --git a/llvm/lib/Target/AMDGPU/VIInstrFormats.td b/llvm/lib/Target/AMDGPU/VIInstrFormats.td
index c211372..8a1cee7 100644 (file)
--- a/llvm/lib/Target/AMDGPU/VIInstrFormats.td
+++ b/llvm/lib/Target/AMDGPU/VIInstrFormats.td
@@ -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;
 

diff --git a/llvm/lib/Target/AMDGPU/VIInstructions.td b/llvm/lib/Target/AMDGPU/VIInstructions.td
index 4b8ce64..ddd164b 100644 (file)
--- a/llvm/lib/Target/AMDGPU/VIInstructions.td
+++ b/llvm/lib/Target/AMDGPU/VIInstructions.td
@@ -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))
 >;
 
 //===----------------------------------------------------------------------===//

diff --git a/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.mov.dpp.ll b/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.mov.dpp.ll
index 72ade21..17b6b7a 100644 (file)
--- a/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.mov.dpp.ll
+++ b/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.mov.dpp.ll
@@ -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 }

diff --git a/llvm/test/MC/AMDGPU/vop_dpp.s b/llvm/test/MC/AMDGPU/vop_dpp.s
new file mode 100644 (file)
index 0000000..2959f3d
--- /dev/null
+++ b/llvm/test/MC/AMDGPU/vop_dpp.s
@@ -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