maxjhandsome
/
llvm-project
forked from OSchip/llvm-project
1
0
Fork 0
Code Issues Pull Requests Packages Releases Wiki Activity

[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
This commit is contained in:
Sam Kolton 2016-03-09 12:29:31 +00:00
parent 10d6f9ac04
commit dfa29f7c5b
11 changed files with 497 additions and 50 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

6
llvm/include/llvm/IR/IntrinsicsAMDGPU.td
View File

@ -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">,

180
llvm/lib/Target/AMDGPU/AsmParser/AMDGPUAsmParser.cpp
View File

@ -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 {
@ -438,7 +456,6 @@ private:
bool ParseSectionDirectiveHSADataGlobalProgram();
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.

69
llvm/lib/Target/AMDGPU/InstPrinter/AMDGPUInstPrinter.cpp
View File

@ -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();

6
llvm/lib/Target/AMDGPU/InstPrinter/AMDGPUInstPrinter.h
View File

@ -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,

19
llvm/lib/Target/AMDGPU/SIDefines.h
View File

@ -29,16 +29,17 @@ enum {
VOP2 = 1 << 11,
VOP3 = 1 << 12,
VOPC = 1 << 13,
DPP = 1 << 14,
MUBUF = 1 << 14,
MTBUF = 1 << 15,
SMRD = 1 << 16,
DS = 1 << 17,
MIMG = 1 << 18,
FLAT = 1 << 19,
WQM = 1 << 20,
VGPRSpill = 1 << 21,
VOPAsmPrefer32Bit = 1 << 22
MUBUF = 1 << 15,
MTBUF = 1 << 16,
SMRD = 1 << 17,
DS = 1 << 18,
MIMG = 1 << 19,
FLAT = 1 << 20,
WQM = 1 << 21,
VGPRSpill = 1 << 22,
VOPAsmPrefer32Bit = 1 << 23
};
}

20
llvm/lib/Target/AMDGPU/SIInstrFormats.td
View File

@ -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} = 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];

83
llvm/lib/Target/AMDGPU/SIInstrInfo.td
View File

@ -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,

7
llvm/lib/Target/AMDGPU/VIInstrFormats.td
View File

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

8
llvm/lib/Target/AMDGPU/VIInstructions.td
View File

@ -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))
>;
//===----------------------------------------------------------------------===//

6
llvm/test/CodeGen/AMDGPU/llvm.amdgcn.mov.dpp.ll
View File

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

143
llvm/test/MC/AMDGPU/vop_dpp.s Normal file
View File

@ -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