llvm.org GIT mirror llvm / 6e4c55e
[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 git-svn-id: https://llvm.org/svn/llvm-project/llvm/trunk@263008 91177308-0d34-0410-b5e6-96231b3b80d8 Sam Kolton 4 years ago
11 changed file(s) with 498 addition(s) and 51 deletion(s). Raw diff Collapse all Expand all
274274 // VI Intrinsics
275275 //===----------------------------------------------------------------------===//
276276
277 // llvm.amdgcn.mov.dpp.i32 <bound_ctrl> >
277 // llvm.amdgcn.mov.dpp.i32 <row_mask> >
278278 def int_amdgcn_mov_dpp :
279279 Intrinsic<[llvm_anyint_ty],
280 [LLVMMatchType<0>, llvm_i32_ty, llvm_i1_ty, llvm_i32_ty,
281 llvm_i32_ty], [IntrNoMem, IntrConvergent]>;
280 [LLVMMatchType<0>, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty,
281 llvm_i1_ty], [IntrNoMem, IntrConvergent]>;
282282
283283 def int_amdgcn_s_dcache_wb :
284284 GCCBuiltin<"__builtin_amdgcn_s_dcache_wb">,
6868 ImmTyTFE,
6969 ImmTyClamp,
7070 ImmTyOMod,
71 ImmTyDppCtrl,
72 ImmTyDppRowMask,
73 ImmTyDppBankMask,
74 ImmTyDppBoundCtrl,
7175 ImmTyDMask,
7276 ImmTyUNorm,
7377 ImmTyDA,
143147 bool defaultTokenHasSuffix() const {
144148 StringRef Token(Tok.Data, Tok.Length);
145149
146 return Token.endswith("_e32") || Token.endswith("_e64");
150 return Token.endswith("_e32") || Token.endswith("_e64") ||
151 Token.endswith("_dpp");
147152 }
148153
149154 bool isToken() const override {
233238 bool isSLC() const { return isImmTy(ImmTySLC); }
234239 bool isTFE() const { return isImmTy(ImmTyTFE); }
235240
241 bool isBankMask() const {
242 return isImmTy(ImmTyDppBankMask);
243 }
244
245 bool isRowMask() const {
246 return isImmTy(ImmTyDppRowMask);
247 }
248
249 bool isBoundCtrl() const {
250 return isImmTy(ImmTyDppBoundCtrl);
251 }
252
236253 void setModifiers(unsigned Mods) {
237254 assert(isReg() || (isImm() && Imm.Modifiers == 0));
238255 if (isReg())
390407 bool isMubufOffset() const;
391408 bool isSMRDOffset() const;
392409 bool isSMRDLiteralOffset() const;
410 bool isDPPCtrl() const;
393411 };
394412
395413 class AMDGPUAsmParser : public MCTargetAsmParser {
438456 bool ParseSectionDirectiveHSARodataReadonlyAgent();
439457
440458 public:
441 public:
442459 enum AMDGPUMatchResultTy {
443460 Match_PreferE32 = FIRST_TARGET_MATCH_RESULT_TY
444461 };
537554 void cvtMIMG(MCInst &Inst, const OperandVector &Operands);
538555 void cvtMIMGAtomic(MCInst &Inst, const OperandVector &Operands);
539556 OperandMatchResultTy parseVOP3OptionalOps(OperandVector &Operands);
557
558 OperandMatchResultTy parseDPPCtrlOps(OperandVector &Operands);
559 OperandMatchResultTy parseDPPOptionalOps(OperandVector &Operands);
560 void cvtDPP_mod(MCInst &Inst, const OperandVector &Operands);
561 void cvtDPP_nomod(MCInst &Inst, const OperandVector &Operands);
562 void cvtDPP(MCInst &Inst, const OperandVector &Operands, bool HasMods);
540563 };
541564
542565 struct OptionalOperand {
11461169 AMDGPUAsmParser::OperandMatchResultTy
11471170 AMDGPUAsmParser::parseIntWithPrefix(const char *Prefix, int64_t &Int,
11481171 int64_t Default) {
1149
11501172 // We are at the end of the statement, and this is a default argument, so
11511173 // use a default value.
11521174 if (getLexer().is(AsmToken::EndOfStatement)) {
12261248
12271249 typedef std::map OptionalImmIndexMap;
12281250
1229 void addOptionalImmOperand(MCInst& Inst, const OperandVector& Operands, OptionalImmIndexMap& OptionalIdx, enum AMDGPUOperand::ImmTy ImmT) {
1251 void addOptionalImmOperand(MCInst& Inst, const OperandVector& Operands,
1252 OptionalImmIndexMap& OptionalIdx,
1253 enum AMDGPUOperand::ImmTy ImmT, int64_t Default = 0) {
12301254 auto i = OptionalIdx.find(ImmT);
12311255 if (i != OptionalIdx.end()) {
12321256 unsigned Idx = i->second;
12331257 ((AMDGPUOperand &)*Operands[Idx]).addImmOperands(Inst, 1);
12341258 } else {
1235 Inst.addOperand(MCOperand::createImm(0));
1259 Inst.addOperand(MCOperand::createImm(Default));
12361260 }
12371261 }
12381262
18951919 addOptionalImmOperand(Inst, Operands, OptionalIdx, AMDGPUOperand::ImmTySLC);
18961920 }
18971921
1922 //===----------------------------------------------------------------------===//
1923 // dpp
1924 //===----------------------------------------------------------------------===//
1925
1926 bool AMDGPUOperand::isDPPCtrl() const {
1927 bool result = isImm() && getImmTy() == ImmTyDppCtrl && isUInt<9>(getImm());
1928 if (result) {
1929 int64_t Imm = getImm();
1930 return ((Imm >= 0x000) && (Imm <= 0x0ff)) ||
1931 ((Imm >= 0x101) && (Imm <= 0x10f)) ||
1932 ((Imm >= 0x111) && (Imm <= 0x11f)) ||
1933 ((Imm >= 0x121) && (Imm <= 0x12f)) ||
1934 (Imm == 0x130) ||
1935 (Imm == 0x134) ||
1936 (Imm == 0x138) ||
1937 (Imm == 0x13c) ||
1938 (Imm == 0x140) ||
1939 (Imm == 0x141) ||
1940 (Imm == 0x142) ||
1941 (Imm == 0x143);
1942 }
1943 return false;
1944 }
1945
1946 AMDGPUAsmParser::OperandMatchResultTy
1947 AMDGPUAsmParser::parseDPPCtrlOps(OperandVector &Operands) {
1948 // ToDo: use same syntax as sp3 for dpp_ctrl
1949 SMLoc S = Parser.getTok().getLoc();
1950 StringRef Prefix;
1951 int64_t Int;
1952
1953 switch(getLexer().getKind()) {
1954 default: return MatchOperand_NoMatch;
1955 case AsmToken::Identifier: {
1956 Prefix = Parser.getTok().getString();
1957
1958 Parser.Lex();
1959 if (getLexer().isNot(AsmToken::Colon))
1960 return MatchOperand_ParseFail;
1961
1962 Parser.Lex();
1963 if (getLexer().isNot(AsmToken::Integer))
1964 return MatchOperand_ParseFail;
1965
1966 if (getParser().parseAbsoluteExpression(Int))
1967 return MatchOperand_ParseFail;
1968 break;
1969 }
1970 }
1971
1972 if (Prefix.equals("row_shl")) {
1973 Int |= 0x100;
1974 } else if (Prefix.equals("row_shr")) {
1975 Int |= 0x110;
1976 } else if (Prefix.equals("row_ror")) {
1977 Int |= 0x120;
1978 } else if (Prefix.equals("wave_shl")) {
1979 Int = 0x130;
1980 } else if (Prefix.equals("wave_rol")) {
1981 Int = 0x134;
1982 } else if (Prefix.equals("wave_shr")) {
1983 Int = 0x138;
1984 } else if (Prefix.equals("wave_ror")) {
1985 Int = 0x13C;
1986 } else if (Prefix.equals("row_mirror")) {
1987 Int = 0x140;
1988 } else if (Prefix.equals("row_half_mirror")) {
1989 Int = 0x141;
1990 } else if (Prefix.equals("row_bcast")) {
1991 if (Int == 15) {
1992 Int = 0x142;
1993 } else if (Int == 31) {
1994 Int = 0x143;
1995 }
1996 } else if (!Prefix.equals("quad_perm")) {
1997 return MatchOperand_NoMatch;
1998 }
1999 Operands.push_back(AMDGPUOperand::CreateImm(Int, S,
2000 AMDGPUOperand::ImmTyDppCtrl));
2001 return MatchOperand_Success;
2002 }
2003
2004 static const OptionalOperand DPPOptionalOps [] = {
2005 {"row_mask", AMDGPUOperand::ImmTyDppRowMask, false, 0xf, nullptr},
2006 {"bank_mask", AMDGPUOperand::ImmTyDppBankMask, false, 0xf, nullptr},
2007 {"bound_ctrl", AMDGPUOperand::ImmTyDppBoundCtrl, false, -1, nullptr}
2008 };
2009
2010 AMDGPUAsmParser::OperandMatchResultTy
2011 AMDGPUAsmParser::parseDPPOptionalOps(OperandVector &Operands) {
2012 SMLoc S = Parser.getTok().getLoc();
2013 OperandMatchResultTy Res = parseOptionalOps(DPPOptionalOps, Operands);
2014 // XXX - sp3 use syntax "bound_ctrl:0" to indicate that bound_ctrl bit was set
2015 if (Res == MatchOperand_Success) {
2016 AMDGPUOperand &Op = ((AMDGPUOperand &)*Operands.back());
2017 // If last operand was parsed as bound_ctrl we should replace it with correct value (1)
2018 if (Op.isImmTy(AMDGPUOperand::ImmTyDppBoundCtrl)) {
2019 Operands.pop_back();
2020 Operands.push_back(
2021 AMDGPUOperand::CreateImm(1, S, AMDGPUOperand::ImmTyDppBoundCtrl));
2022 return MatchOperand_Success;
2023 }
2024 }
2025 return Res;
2026 }
2027
2028 void AMDGPUAsmParser::cvtDPP_mod(MCInst &Inst, const OperandVector &Operands) {
2029 cvtDPP(Inst, Operands, true);
2030 }
2031
2032 void AMDGPUAsmParser::cvtDPP_nomod(MCInst &Inst, const OperandVector &Operands) {
2033 cvtDPP(Inst, Operands, false);
2034 }
2035
2036 void AMDGPUAsmParser::cvtDPP(MCInst &Inst, const OperandVector &Operands,
2037 bool HasMods) {
2038 OptionalImmIndexMap OptionalIdx;
2039
2040 unsigned I = 1;
2041 const MCInstrDesc &Desc = MII.get(Inst.getOpcode());
2042 for (unsigned J = 0; J < Desc.getNumDefs(); ++J) {
2043 ((AMDGPUOperand &)*Operands[I++]).addRegOperands(Inst, 1);
2044 }
2045
2046 for (unsigned E = Operands.size(); I != E; ++I) {
2047 AMDGPUOperand &Op = ((AMDGPUOperand &)*Operands[I]);
2048 // Add the register arguments
2049 if (!HasMods && Op.isReg()) {
2050 Op.addRegOperands(Inst, 1);
2051 } else if (HasMods && Op.isRegOrImmWithInputMods()) {
2052 Op.addRegOrImmWithInputModsOperands(Inst, 2);
2053 } else if (Op.isDPPCtrl()) {
2054 Op.addImmOperands(Inst, 1);
2055 } else if (Op.isImm()) {
2056 // Handle optional arguments
2057 OptionalIdx[Op.getImmTy()] = I;
2058 } else {
2059 llvm_unreachable("Invalid operand type");
2060 }
2061 }
2062
2063 // ToDo: fix default values for row_mask and bank_mask
2064 addOptionalImmOperand(Inst, Operands, OptionalIdx, AMDGPUOperand::ImmTyDppRowMask, 0xf);
2065 addOptionalImmOperand(Inst, Operands, OptionalIdx, AMDGPUOperand::ImmTyDppBankMask, 0xf);
2066 addOptionalImmOperand(Inst, Operands, OptionalIdx, AMDGPUOperand::ImmTyDppBoundCtrl);
2067 }
18982068
18992069
19002070 /// Force static initialization.
2727 printAnnotation(OS, Annot);
2828 }
2929
30 void AMDGPUInstPrinter::printU4ImmOperand(const MCInst *MI, unsigned OpNo,
31 raw_ostream &O) {
32 O << formatHex(MI->getOperand(OpNo).getImm() & 0xf);
33 }
34
3035 void AMDGPUInstPrinter::printU8ImmOperand(const MCInst *MI, unsigned OpNo,
3136 raw_ostream &O) {
3237 O << formatHex(MI->getOperand(OpNo).getImm() & 0xff);
4045 void AMDGPUInstPrinter::printU32ImmOperand(const MCInst *MI, unsigned OpNo,
4146 raw_ostream &O) {
4247 O << formatHex(MI->getOperand(OpNo).getImm() & 0xffffffff);
48 }
49
50 void AMDGPUInstPrinter::printU4ImmDecOperand(const MCInst *MI, unsigned OpNo,
51 raw_ostream &O) {
52 O << formatDec(MI->getOperand(OpNo).getImm() & 0xf);
4353 }
4454
4555 void AMDGPUInstPrinter::printU8ImmDecOperand(const MCInst *MI, unsigned OpNo,
250260 raw_ostream &O) {
251261 if (MII.get(MI->getOpcode()).TSFlags & SIInstrFlags::VOP3)
252262 O << "_e64 ";
263 else if (MII.get(MI->getOpcode()).TSFlags & SIInstrFlags::DPP)
264 O << "_dpp ";
253265 else
254266 O << "_e32 ";
255267
387399 O << '|';
388400 }
389401
402
403 void AMDGPUInstPrinter::printDPPCtrlOperand(const MCInst *MI, unsigned OpNo,
404 raw_ostream &O) {
405 unsigned Imm = MI->getOperand(OpNo).getImm();
406 if ((Imm >= 0x000) && (Imm <= 0x0ff)) {
407 O << " quad_perm:";
408 printU8ImmDecOperand(MI, OpNo, O);
409 } else if ((Imm >= 0x101) && (Imm <= 0x10f)) {
410 O << " row_shl:";
411 printU4ImmDecOperand(MI, OpNo, O);
412 } else if ((Imm >= 0x111) && (Imm <= 0x11f)) {
413 O << " row_shr:";
414 printU4ImmDecOperand(MI, OpNo, O);
415 } else if ((Imm >= 0x121) && (Imm <= 0x12f)) {
416 O << " row_ror:";
417 printU4ImmDecOperand(MI, OpNo, O);
418 } else if (Imm == 0x130) {
419 O << " wave_shl:1";
420 } else if (Imm == 0x134) {
421 O << " wave_rol:1";
422 } else if (Imm == 0x138) {
423 O << " wave_shr:1";
424 } else if (Imm == 0x13c) {
425 O << " wave_ror:1";
426 } else if (Imm == 0x140) {
427 O << " row_mirror:1";
428 } else if (Imm == 0x141) {
429 O << " row_half_mirror:1";
430 } else if (Imm == 0x142) {
431 O << " row_bcast:15";
432 } else if (Imm == 0x143) {
433 O << " row_bcast:31";
434 } else {
435 llvm_unreachable("Invalid dpp_ctrl value");
436 }
437 }
438
439 void AMDGPUInstPrinter::printRowMaskOperand(const MCInst *MI, unsigned OpNo,
440 raw_ostream &O) {
441 O << " row_mask:";
442 printU4ImmOperand(MI, OpNo, O);
443 }
444
445 void AMDGPUInstPrinter::printBankMaskOperand(const MCInst *MI, unsigned OpNo,
446 raw_ostream &O) {
447 O << " bank_mask:";
448 printU4ImmOperand(MI, OpNo, O);
449 }
450
451 void AMDGPUInstPrinter::printBoundCtrlOperand(const MCInst *MI, unsigned OpNo,
452 raw_ostream &O) {
453 unsigned Imm = MI->getOperand(OpNo).getImm();
454 if (Imm) {
455 O << " bound_ctrl:0"; // XXX - this syntax is used in sp3
456 }
457 }
458
390459 void AMDGPUInstPrinter::printInterpSlot(const MCInst *MI, unsigned OpNum,
391460 raw_ostream &O) {
392461 unsigned Imm = MI->getOperand(OpNum).getImm();
3232 const MCRegisterInfo &MRI);
3333
3434 private:
35 void printU4ImmOperand(const MCInst *MI, unsigned OpNo, raw_ostream &O);
3536 void printU8ImmOperand(const MCInst *MI, unsigned OpNo, raw_ostream &O);
3637 void printU16ImmOperand(const MCInst *MI, unsigned OpNo, raw_ostream &O);
38 void printU4ImmDecOperand(const MCInst *MI, unsigned OpNo, raw_ostream &O);
3739 void printU8ImmDecOperand(const MCInst *MI, unsigned OpNo, raw_ostream &O);
3840 void printU16ImmDecOperand(const MCInst *MI, unsigned OpNo, raw_ostream &O);
3941 void printU32ImmOperand(const MCInst *MI, unsigned OpNo, raw_ostream &O);
6062 void printImmediate64(uint64_t I, raw_ostream &O);
6163 void printOperand(const MCInst *MI, unsigned OpNo, raw_ostream &O);
6264 void printOperandAndMods(const MCInst *MI, unsigned OpNo, raw_ostream &O);
65 void printDPPCtrlOperand(const MCInst *MI, unsigned OpNo, raw_ostream &O);
66 void printRowMaskOperand(const MCInst *MI, unsigned OpNo, raw_ostream &O);
67 void printBankMaskOperand(const MCInst *MI, unsigned OpNo, raw_ostream &O);
68 void printBoundCtrlOperand(const MCInst *MI, unsigned OpNo, raw_ostream &O);
6369 static void printInterpSlot(const MCInst *MI, unsigned OpNum, raw_ostream &O);
6470 void printMemOperand(const MCInst *MI, unsigned OpNo, raw_ostream &O);
6571 static void printIfSet(const MCInst *MI, unsigned OpNo, raw_ostream &O,
2828 VOP2 = 1 << 11,
2929 VOP3 = 1 << 12,
3030 VOPC = 1 << 13,
31 DPP = 1 << 14,
3132
32 MUBUF = 1 << 14,
33 MTBUF = 1 << 15,
34 SMRD = 1 << 16,
35 DS = 1 << 17,
36 MIMG = 1 << 18,
37 FLAT = 1 << 19,
38 WQM = 1 << 20,
39 VGPRSpill = 1 << 21,
40 VOPAsmPrefer32Bit = 1 << 22
33 MUBUF = 1 << 15,
34 MTBUF = 1 << 16,
35 SMRD = 1 << 17,
36 DS = 1 << 18,
37 MIMG = 1 << 19,
38 FLAT = 1 << 20,
39 WQM = 1 << 21,
40 VGPRSpill = 1 << 22,
41 VOPAsmPrefer32Bit = 1 << 23
4142 };
4243 }
4344
3030 field bits<1> VOP2 = 0;
3131 field bits<1> VOP3 = 0;
3232 field bits<1> VOPC = 0;
33 field bits<1> DPP = 0;
3334
3435 field bits<1> MUBUF = 0;
3536 field bits<1> MTBUF = 0;
6263 let TSFlags{11} = VOP2;
6364 let TSFlags{12} = VOP3;
6465 let TSFlags{13} = VOPC;
65
66 let TSFlags{14} = MUBUF;
67 let TSFlags{15} = MTBUF;
68 let TSFlags{16} = SMRD;
69 let TSFlags{17} = DS;
70 let TSFlags{18} = MIMG;
71 let TSFlags{19} = FLAT;
72 let TSFlags{20} = WQM;
73 let TSFlags{21} = VGPRSpill;
74 let TSFlags{22} = VOPAsmPrefer32Bit;
66 let TSFlags{14} = DPP;
67
68 let TSFlags{15} = MUBUF;
69 let TSFlags{16} = MTBUF;
70 let TSFlags{17} = SMRD;
71 let TSFlags{18} = DS;
72 let TSFlags{19} = MIMG;
73 let TSFlags{20} = FLAT;
74 let TSFlags{21} = WQM;
75 let TSFlags{22} = VGPRSpill;
76 let TSFlags{23} = VOPAsmPrefer32Bit;
7577
7678 let SchedRW = [Write32Bit];
7779
533533 "isSMRDLiteralOffset"
534534 >;
535535
536 def DPPCtrlMatchClass : AsmOperandClass {
537 let Name = "DPPCtrl";
538 let PredicateMethod = "isDPPCtrl";
539 let ParserMethod = "parseDPPCtrlOps";
540 let RenderMethod = "addImmOperands";
541 let IsOptional = 0;
542 }
543
544 class DPPOptionalMatchClass : AsmOperandClass {
545 let Name = "DPPOptional"#OpName;
546 let PredicateMethod = "is"#OpName;
547 let ParserMethod = "parseDPPOptionalOps";
548 let RenderMethod = "addImmOperands";
549 let IsOptional = 1;
550 }
551
536552 class OptionalImmAsmOperand : AsmOperandClass {
537553 let Name = "Imm"#OpName;
538554 let PredicateMethod = "isImm";
665681
666682 def lwe : NamedBitOperand<"LWE"> {
667683 let ParserMatchClass = NamedBitMatchClass<"LWE">;
684 }
685
686 def dpp_ctrl : Operand {
687 let PrintMethod = "printDPPCtrlOperand";
688 let ParserMatchClass = DPPCtrlMatchClass;
689 }
690
691 def row_mask : Operand {
692 let PrintMethod = "printRowMaskOperand";
693 let ParserMatchClass = DPPOptionalMatchClass<"RowMask">;
694 }
695
696 def bank_mask : Operand {
697 let PrintMethod = "printBankMaskOperand";
698 let ParserMatchClass = DPPOptionalMatchClass<"BankMask">;
699 }
700
701 def bound_ctrl : Operand {
702 let PrintMethod = "printBoundCtrlOperand";
703 let ParserMatchClass = DPPOptionalMatchClass<"BoundCtrl">;
668704 }
669705
670706 } // End OperandType = "OPERAND_IMMEDIATE"
12791315 !if (!eq(HasModifiers, 1),
12801316 // VOP1_DPP with modifiers
12811317 (ins InputModsNoDefault:$src0_modifiers, Src0RC:$src0,
1282 i32imm:$dpp_ctrl, i1imm:$bound_ctrl,
1283 i32imm:$bank_mask, i32imm:$row_mask)
1318 dpp_ctrl:$dpp_ctrl, row_mask:$row_mask,
1319 bank_mask:$bank_mask, bound_ctrl:$bound_ctrl)
12841320 /* else */,
12851321 // VOP1_DPP without modifiers
1286 (ins Src0RC:$src0, i32imm:$dpp_ctrl, i1imm:$bound_ctrl,
1287 i32imm:$bank_mask, i32imm:$row_mask)
1322 (ins Src0RC:$src0, dpp_ctrl:$dpp_ctrl, row_mask:$row_mask,
1323 bank_mask:$bank_mask, bound_ctrl:$bound_ctrl)
12881324 /* endif */)
1289 /* NumSrcArgs == 2 */,
1325 /* NumSrcArgs == 2 */,
12901326 !if (!eq(HasModifiers, 1),
12911327 // VOP2_DPP with modifiers
12921328 (ins InputModsNoDefault:$src0_modifiers, Src0RC:$src0,
1293 InputModsNoDefault:$src1_modifiers, Src1RC:$src1,
1294 i32imm:$dpp_ctrl, i1imm:$bound_ctrl,
1295 i32imm:$bank_mask, i32imm:$row_mask)
1329 InputModsNoDefault:$src1_modifiers, Src1RC:$src1,
1330 dpp_ctrl:$dpp_ctrl, row_mask:$row_mask,
1331 bank_mask:$bank_mask, bound_ctrl:$bound_ctrl)
12961332 /* else */,
12971333 // VOP2_DPP without modifiers
1298 (ins Src0RC:$src0, Src1RC:$src1, i32imm:$dpp_ctrl, i1imm:$bound_ctrl,
1299 i32imm:$bank_mask, i32imm:$row_mask)
1334 (ins Src0RC:$src0, Src1RC:$src1, dpp_ctrl:$dpp_ctrl,
1335 row_mask:$row_mask, bank_mask:$bank_mask,
1336 bound_ctrl:$bound_ctrl)
13001337 /* endif */));
13011338 }
13021339
13371374 " $src1_modifiers,"));
13381375 string args = !if(!eq(HasModifiers, 0),
13391376 getAsm32<0, NumSrcArgs, DstVT>.ret,
1340 src0#src1);
1341 string ret = " "#dst#args#", $dpp_ctrl, "#"$bound_ctrl, "#"$bank_mask, "#"$row_mask";
1377 ", "#src0#src1);
1378 string ret = dst#args#" $dpp_ctrl $row_mask $bank_mask $bound_ctrl";
13421379 }
13431380
13441381 class VOPProfile _ArgVT> {
13501387 field ValueType Src1VT = ArgVT[2];
13511388 field ValueType Src2VT = ArgVT[3];
13521389 field RegisterOperand DstRC = getVALUDstForVT.ret;
1353 field RegisterClass DstRCDPP = !if(!eq(DstVT.Size, 64), VReg_64, VGPR_32);
1390 field RegisterOperand DstRCDPP = getVALUDstForVT.ret;
13541391 field RegisterOperand Src0RC32 = getVOPSrc0ForVT.ret;
13551392 field RegisterClass Src1RC32 = getVOPSrc1ForVT.ret;
13561393 field RegisterOperand Src0RC64 = getVOP3SrcForVT.ret;
14961533 let Ins32 = (ins Src0RC32:$src0, Src1RC32:$src1, VGPR_32:$src2);
14971534 let Ins64 = getIns64, 3,
14981535 HasModifiers>.ret;
1536 let InsDPP = (ins InputModsNoDefault:$src0_modifiers, Src0RC32:$src0,
1537 InputModsNoDefault:$src1_modifiers, Src1RC32:$src1,
1538 VGPR_32:$src2, // stub argument
1539 dpp_ctrl:$dpp_ctrl, row_mask:$row_mask,
1540 bank_mask:$bank_mask, bound_ctrl:$bound_ctrl);
14991541 let Asm32 = getAsm32<1, 2, f32>.ret;
15001542 let Asm64 = getAsm64<1, 2, HasModifiers, f32>.ret;
1543 let AsmDPP = getAsmDPP<1, 2, HasModifiers, f32>.ret;
15011544 }
15021545 def VOP_F64_F64_F64_F64 : VOPProfile <[f64, f64, f64, f64]>;
15031546 def VOP_I32_I32_I32_I32 : VOPProfile <[i32, i32, i32, i32]>;
16061649
16071650 class VOP1_DPP :
16081651 VOP1_DPPe ,
1609 VOP_DPP > {
1652 VOP_DPP , p.HasModifiers> {
16101653 let AssemblerPredicates = [isVI];
1611 let src0_modifiers = !if(p.HasModifiers, ?, 0);
1654 let src0_modifiers = !if(p.HasModifiers, ?, 0);
16121655 let src1_modifiers = 0;
16131656 }
16141657
16641707
16651708 def _vi : VOP2_Real_vi ;
16661709
1710 }
1711
1712 class VOP2_DPP :
1713 VOP2_DPPe ,
1714 VOP_DPP {
1715 let AssemblerPredicates = [isVI];
1716 let src0_modifiers = !if(p.HasModifiers, ?, 0);
1717 let src1_modifiers = !if(p.HasModifiers, ?, 0);
16671718 }
16681719
16691720 class VOP3DisableFields {
19281979
19291980 defm _e64 : VOP3_2_m
19301981 revOp, p.HasModifiers>;
1982
1983 def _dpp : VOP2_DPP ;
19311984 }
19321985
19331986 multiclass VOP2Inst
168168 let Inst{63} = src2_modifiers{0};
169169 }
170170
171 class VOP_DPP pattern> :
171 class VOP_DPP pattern, bit HasMods = 0> :
172172 VOPAnyCommon {
173 let DPP = 1;
173174 let Size = 8;
175
176 let AsmMatchConverter = !if(!eq(HasMods,1), "cvtDPP_mod", "cvtDPP_nomod");
174177 }
175178
176179 class VOP_DPPe : Enc64 {
202205 let Inst{31-25} = 0x3f; //encoding
203206 }
204207
205 class VOP2_DPPe op> : Enc32 {
208 class VOP2_DPPe op> : VOP_DPPe {
206209 bits<8> vdst;
207210 bits<8> src1;
208211
120120 //===----------------------------------------------------------------------===//
121121
122122 def : Pat <
123 (int_amdgcn_mov_dpp i32:$src, imm:$dpp_ctrl, imm:$bound_ctrl,
124 imm:$bank_mask, imm:$row_mask),
125 (V_MOV_B32_dpp $src, (as_i32imm $dpp_ctrl), (as_i1imm $bound_ctrl),
126 (as_i32imm $bank_mask), (as_i32imm $row_mask))
123 (int_amdgcn_mov_dpp i32:$src, imm:$dpp_ctrl, imm:$row_mask, imm:$bank_mask,
124 imm:$bound_ctrl),
125 (V_MOV_B32_dpp $src, (as_i32imm $dpp_ctrl), (as_i32imm $row_mask),
126 (as_i32imm $bank_mask), (as_i1imm $bound_ctrl))
127127 >;
128128
129129 //===----------------------------------------------------------------------===//
0 ; RUN: llc -march=amdgcn -mcpu=tonga -verify-machineinstrs -show-mc-encoding < %s | FileCheck -check-prefix=VI %s
11
22 ; VI-LABEL: {{^}}dpp_test:
3 ; VI: v_mov_b32 v0, v0, 1, -1, 1, 1 ; encoding: [0xfa,0x02,0x00,0x7e,0x00,0x01,0x08,0x11]
3 ; 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]
44 define void @dpp_test(i32 addrspace(1)* %out, i32 %in) {
5 %tmp0 = call i32 @llvm.amdgcn.mov.dpp.i32(i32 %in, i32 1, i1 1, i32 1, i32 1) #0
5 %tmp0 = call i32 @llvm.amdgcn.mov.dpp.i32(i32 %in, i32 1, i32 1, i32 1, i1 1) #0
66 store i32 %tmp0, i32 addrspace(1)* %out
77 ret void
88 }
99
10 declare i32 @llvm.amdgcn.mov.dpp.i32(i32, i32, i1, i32, i32) #0
10 declare i32 @llvm.amdgcn.mov.dpp.i32(i32, i32, i32, i32, i1) #0
1111
1212 attributes #0 = { nounwind readnone convergent }
0 // RUN: llvm-mc -arch=amdgcn -mcpu=tonga -show-encoding %s | FileCheck %s --check-prefix=GCN --check-prefix=CIVI --check-prefix=VI
1 // RUN: not llvm-mc -arch=amdgcn -show-encoding %s 2>&1 | FileCheck %s --check-prefix=NOSI --check-prefix=NOSICI
2 // RUN: not llvm-mc -arch=amdgcn -mcpu=SI -show-encoding %s 2>&1 | FileCheck %s --check-prefix=NOSI --check-prefix=NOSICI
3 // RUN: not llvm-mc -arch=amdgcn -mcpu=bonaire -show-encoding %s 2>&1 | FileCheck %s --check-prefix=NOSICI
4
5 //===----------------------------------------------------------------------===//
6 // Check dpp_ctrl values
7 //===----------------------------------------------------------------------===//
8
9 // NOSICI: error:
10 // 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]
11 v_mov_b32 v0, v0 quad_perm:37
12
13 // NOSICI: error:
14 // 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]
15 v_mov_b32 v0, v0 row_shl:1
16
17 // NOSICI: error:
18 // 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]
19 v_mov_b32 v0, v0 row_shr:0xf
20
21 // NOSICI: error:
22 // 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]
23 v_mov_b32 v0, v0 row_ror:0xc
24
25 // NOSICI: error:
26 // 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]
27 v_mov_b32 v0, v0 wave_shl:1
28
29 // NOSICI: error:
30 // 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]
31 v_mov_b32 v0, v0 wave_rol:1
32
33 // NOSICI: error:
34 // 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]
35 v_mov_b32 v0, v0 wave_shr:1
36
37 // NOSICI: error:
38 // 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]
39 v_mov_b32 v0, v0 wave_ror:1
40
41 // NOSICI: error:
42 // 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]
43 v_mov_b32 v0, v0 row_mirror:1
44
45 // NOSICI: error:
46 // 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]
47 v_mov_b32 v0, v0 row_half_mirror:1
48
49 // NOSICI: error:
50 // 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]
51 v_mov_b32 v0, v0 row_bcast:15
52
53 // NOSICI: error:
54 // 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]
55 v_mov_b32 v0, v0 row_bcast:31
56
57 //===----------------------------------------------------------------------===//
58 // Check optional fields
59 //===----------------------------------------------------------------------===//
60
61 // NOSICI: error:
62 // 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]
63 v_mov_b32 v0, v0 quad_perm:1 row_mask:0xa bank_mask:0x1 bound_ctrl:0
64
65 // NOSICI: error:
66 // 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]
67 v_mov_b32 v0, v0 quad_perm:1 row_mask:0xa
68
69 // NOSICI: error:
70 // 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]
71 v_mov_b32 v0, v0 quad_perm:1 bank_mask:0x1
72
73 // NOSICI: error:
74 // 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]
75 v_mov_b32 v0, v0 quad_perm:1 bound_ctrl:0
76
77 // NOSICI: error:
78 // 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]
79 v_mov_b32 v0, v0 quad_perm:1 row_mask:0xa bank_mask:0x1
80
81 // NOSICI: error:
82 // 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]
83 v_mov_b32 v0, v0 quad_perm:1 row_mask:0xa bound_ctrl:0
84
85 // NOSICI: error:
86 // 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]
87 v_mov_b32 v0, v0 quad_perm:1 bank_mask:0x1 bound_ctrl:0
88
89 //===----------------------------------------------------------------------===//
90 // Check VOP1 opcodes
91 //===----------------------------------------------------------------------===//
92 // ToDo: v_nop
93
94 // NOSICI: error:
95 // 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]
96 v_cvt_u32_f32 v0, v0 row_shl:1 row_mask:0xa bank_mask:0x1 bound_ctrl:0
97
98 // NOSICI: error:
99 // 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]
100 v_fract_f32 v0, v0 row_shl:1 row_mask:0xa bank_mask:0x1 bound_ctrl:0
101
102 // NOSICI: error:
103 // 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]
104 v_sin_f32 v0, v0 row_shl:1 row_mask:0xa bank_mask:0x1 bound_ctrl:0
105
106 //===----------------------------------------------------------------------===//
107 // Check VOP2 opcodes
108 //===----------------------------------------------------------------------===//
109 // ToDo: VOP2bInst instructions: v_add_u32, v_sub_u32 ... (vcc and ApplyMnemonic in AsmMatcherEmitter.cpp)
110 // ToDo: v_mac_f32 (VOP_MAC)
111
112 // NOSICI: error:
113 // 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]
114 v_add_f32 v0, v0, v0 row_shl:1 row_mask:0xa bank_mask:0x1 bound_ctrl:0
115
116 // NOSICI: error:
117 // 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]
118 v_min_f32 v0, v0, v0 row_shl:1 row_mask:0xa bank_mask:0x1 bound_ctrl:0
119
120 // NOSICI: error:
121 // 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]
122 v_and_b32 v0, v0, v0 row_shl:1 row_mask:0xa bank_mask:0x1 bound_ctrl:0
123
124 //===----------------------------------------------------------------------===//
125 // Check modifiers
126 //===----------------------------------------------------------------------===//
127
128 // NOSICI: error:
129 // 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]
130 v_add_f32 v0, -v0, v0 row_shl:1 row_mask:0xa bank_mask:0x1 bound_ctrl:0
131
132 // NOSICI: error:
133 // 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]
134 v_add_f32 v0, v0, |v0| row_shl:1 row_mask:0xa bank_mask:0x1 bound_ctrl:0
135
136 // NOSICI: error:
137 // 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]
138 v_add_f32 v0, -v0, |v0| row_shl:1 row_mask:0xa bank_mask:0x1 bound_ctrl:0
139
140 // NOSICI: error:
141 // 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]
142 v_add_f32 v0, |v0|, -v0 row_shl:1 row_mask:0xa bank_mask:0x1 bound_ctrl:0