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