diff --git a/amd/comgr/src/comgr-hotswap-b0a0.cpp b/amd/comgr/src/comgr-hotswap-b0a0.cpp index 9a60572cb3989..7cb4f2f21be96 100755 --- a/amd/comgr/src/comgr-hotswap-b0a0.cpp +++ b/amd/comgr/src/comgr-hotswap-b0a0.cpp @@ -4083,6 +4083,24 @@ struct FiniteControlFlowAudit { bool HasUnboundedIndirectEntries = false; }; +// Some B0-only vector encodings are intentionally absent from the A0 MC +// decoder used by hotswap. The legacy VOP3 encoding has the exact six-bit +// major 0x34 (Inst[31:26]): it cannot transfer control or write the scalar +// MODE register. An undecoded instance therefore remains opaque to dataflow, +// but it is not an object-wide indirect-entry source. Keep this whitelist on +// the exact encoding class; every other undecoded encoding retains the +// fail-closed behavior. +static bool +isProvablyNonControlFlowUndecodedVectorInst(const InternalDecodedInst &DI, + ArrayRef Text) { + if (DI.DecodeSucceeded || DI.Offset > Text.size() || + MinInstSize > Text.size() - DI.Offset) + return false; + uint32_t Word = + support::endian::read32le(Text.data() + static_cast(DI.Offset)); + return (Word & 0xfc000000u) == (0x34u << 26); +} + static FiniteControlFlowAudit auditFiniteIndirectControlFlow( ArrayRef Decoded, const LLVMState &LS, uint64_t TextAddr, uint64_t TextSize, @@ -4091,7 +4109,8 @@ static FiniteControlFlowAudit auditFiniteIndirectControlFlow( const ControlFlowScanIndex &Index, ArrayRef FiniteSetPcTransfers, ArrayRef BoundedReturns, - ArrayRef SymbolLessRegions) { + ArrayRef SymbolLessRegions, + ArrayRef Text) { FiniteControlFlowAudit Audit{BitVector(FiniteSetPcTransfers.size()), true}; auto markUnboundedIndirectEntry = [&]() { Audit.Closed = false; @@ -4301,7 +4320,9 @@ static FiniteControlFlowAudit auditFiniteIndirectControlFlow( markUnboundedIndirectEntry(); } for (int I = Reachable.find_first(); I >= 0; I = Reachable.find_next(I)) - if (!Decoded[static_cast(I)].DecodeSucceeded) + if (!Decoded[static_cast(I)].DecodeSucceeded && + !isProvablyNonControlFlowUndecodedVectorInst( + Decoded[static_cast(I)], Text)) markUnboundedIndirectEntry(); return Audit; } @@ -5937,7 +5958,7 @@ std::optional collectDirectBranchTargets( FiniteControlFlowAudit Audit = auditFiniteIndirectControlFlow( Decoded, LS, TextAddr, TextSize, FunctionRanges, DeclaredEntries, ExternalEntries, *Index, EnabledSetPcTransfers, AllBoundedReturns, - SymbolLessRegions); + SymbolLessRegions, Text); if (Audit.InvalidSetPcCandidates.any()) { for (size_t I = 0; I != EnabledSetPcTransfers.size(); ++I) { if (!Audit.InvalidSetPcCandidates.test(I)) @@ -5965,7 +5986,7 @@ std::optional collectDirectBranchTargets( Audit = auditFiniteIndirectControlFlow( Decoded, LS, TextAddr, TextSize, FunctionRanges, DeclaredEntries, ExternalEntries, *Index, EnabledSetPcTransfers, AllBoundedReturns, - SymbolLessRegions); + SymbolLessRegions, Text); } if (!Audit.Closed && !EnabledSetPcTransfers.empty()) { for (const FiniteSetPcTransfer &Enabled : EnabledSetPcTransfers) diff --git a/amd/comgr/src/comgr-hotswap-internal.h b/amd/comgr/src/comgr-hotswap-internal.h index c6e695b4d5b78..095a82c0589a9 100644 --- a/amd/comgr/src/comgr-hotswap-internal.h +++ b/amd/comgr/src/comgr-hotswap-internal.h @@ -1467,6 +1467,41 @@ struct PatchContext { llvm::StringMap KernelVgprGranuleCache; }; +/// One node in the all-path proof that an incoming physical VGPR value is +/// killed before it can be observed. Opaque nodes and unsafe exits observe +/// every still-live value conservatively. A safe terminal (s_endpgm or the +/// patched site on a later loop iteration) observes none. +struct ForwardVgprProofNode { + llvm::BitVector Uses; + llvm::BitVector FullDefs; + llvm::SmallVector Successors; + bool Opaque = false; + bool HasUnsafeExit = false; + bool SafeTerminal = false; + + explicit ForwardVgprProofNode(unsigned MaxVgprs = 0) + : Uses(MaxVgprs), FullDefs(MaxVgprs) {} +}; + +/// Return physical VGPR values whose incoming contents are killed on every +/// path before a use, opaque instruction, unsafe exit, or non-killing cycle. +/// Malformed graph inputs fail closed with std::nullopt. +std::optional +computeForwardDeadVgprs(llvm::ArrayRef Nodes, + size_t EntryNode, unsigned MaxVgprs); + +/// True when [Base, Base + Width) is non-empty, within MaxVgprs, and does not +/// cross one of gfx1250's 256-register physical VGPR banks. +bool physicalVgprRangeFitsOneBank(unsigned Base, unsigned Width, + unsigned MaxVgprs); + +/// Return true when \p Reg or one of its aliases belongs to a physical vector +/// register file. Physical-VGPR proofs use this after encoded-range recovery +/// fails: a true result must invalidate the proof rather than silently treating +/// the operand as scalar. +bool isVectorRegisterOrAlias(llvm::MCRegister Reg, + const llvm::MCRegisterInfo &MRI); + enum class VgprMsbOperand : unsigned { Src0 = 0, Src1 = 2, @@ -1490,6 +1525,13 @@ void ensureVgprMsbModes(PatchContext &Ctx); [[nodiscard]] std::optional getLocallyEstablishedVgprMsbMode(PatchContext &Ctx, size_t Idx); +/// Apply one instruction's persistent VGPR-MSB transfer to an exact packed +/// mode. Returns VgprMsbUnknown when the incoming state or the instruction's +/// MODE effect is ambiguous; an exact setter can recover an exact mode. +int16_t transferExactVgprMsbMode(int16_t Incoming, + const InternalDecodedInst &DI, + const LLVMState &LS); + unsigned getVgprMsbBank(unsigned Mode, VgprMsbOperand Operand); void setVgprMsbBank(unsigned &Mode, VgprMsbOperand Operand, unsigned Bank); diff --git a/amd/comgr/src/comgr-hotswap-patch-wmma-scale16.cpp b/amd/comgr/src/comgr-hotswap-patch-wmma-scale16.cpp index 0b676fc5bee7c..b4cb2235a0fb1 100644 --- a/amd/comgr/src/comgr-hotswap-patch-wmma-scale16.cpp +++ b/amd/comgr/src/comgr-hotswap-patch-wmma-scale16.cpp @@ -46,8 +46,11 @@ /// patch failed so the rewrite returns an error instead of a miscompile. A loud /// failure beats silent wrong results. /// -/// The 32x16x128_f4 (M=32) variant also needs an M-split; it is not lowered -/// exactly yet and fails closed. +/// The 32x16x128_f4 (M=32) variant is split into two M=16 halves, and each +/// resulting half is K-split as above, for four exact block-32 WMMAs total. +/// Scratch reuse inside a fully allocated kernel is allowed only when exact +/// all-path physical-register liveness proves each of its four scratch values +/// dead. /// //===----------------------------------------------------------------------===// @@ -57,8 +60,10 @@ #include "llvm/ADT/StringExtras.h" #include "llvm/ADT/StringRef.h" #include "llvm/ADT/Twine.h" +#include "llvm/Support/FormatVariadic.h" #include "llvm/Support/raw_ostream.h" +#include #include using namespace llvm; @@ -75,6 +80,12 @@ static constexpr unsigned VOP3PXSize = 16; static constexpr unsigned VgprEncBase = 256; static constexpr unsigned VgprBankSize = 256; +bool physicalVgprRangeFitsOneBank(unsigned Base, unsigned Width, + unsigned MaxVgprs) { + return Width != 0 && Base < MaxVgprs && Width <= MaxVgprs - Base && + Base / VgprBankSize == (Base + Width - 1) / VgprBankSize; +} + static std::string vgprName(unsigned N) { return ("v" + Twine(N)).str(); } static std::string encodedVgprName(unsigned Physical) { @@ -182,7 +193,7 @@ using VgprBankRequirement = std::pair; static void emitModeForOperands(raw_string_ostream &OS, unsigned &CurrentMode, - std::initializer_list Requirements) { + ArrayRef Requirements) { unsigned NewMode = CurrentMode; for (const VgprBankRequirement &Requirement : Requirements) setVgprMsbBank(NewMode, Requirement.first, Requirement.second); @@ -338,6 +349,52 @@ static void emitVgprCopy(raw_string_ostream &OS, unsigned DstBase, } } +static void emitScalePairSplitInPlace(raw_string_ostream &OS, unsigned Lo, + unsigned Hi, unsigned Tmp, + unsigned &CurrentMode) { + unsigned PairBank = Lo / VgprBankSize; + unsigned TmpBank = Tmp / VgprBankSize; + emitVgprCopy(OS, Tmp, Lo, /*W=*/1, TmpBank, CurrentMode); + emitModeForOperands(OS, CurrentMode, + {{VgprMsbOperand::Dst, PairBank}, + {VgprMsbOperand::Src0, TmpBank}, + {VgprMsbOperand::Src1, PairBank}}); + OS << "v_perm_b32 " << encodedVgprName(Lo) << ", " << encodedVgprName(Tmp) + << ", " << encodedVgprName(Hi) << ", 0x06040200\n"; + OS << "v_perm_b32 " << encodedVgprName(Hi) << ", " << encodedVgprName(Tmp) + << ", " << encodedVgprName(Hi) << ", 0x07050301\n"; +} + +static void emitScalePairRestoreInPlace(raw_string_ostream &OS, unsigned Lo, + unsigned Hi, unsigned Tmp, + unsigned &CurrentMode) { + unsigned PairBank = Lo / VgprBankSize; + unsigned TmpBank = Tmp / VgprBankSize; + emitVgprCopy(OS, Tmp, Lo, /*W=*/1, TmpBank, CurrentMode); + emitModeForOperands(OS, CurrentMode, + {{VgprMsbOperand::Dst, PairBank}, + {VgprMsbOperand::Src0, TmpBank}, + {VgprMsbOperand::Src1, PairBank}}); + OS << "v_perm_b32 " << encodedVgprName(Lo) << ", " << encodedVgprName(Tmp) + << ", " << encodedVgprName(Hi) << ", 0x05010400\n"; + OS << "v_perm_b32 " << encodedVgprName(Hi) << ", " << encodedVgprName(Tmp) + << ", " << encodedVgprName(Hi) << ", 0x07030602\n"; +} + +static void emitMaskVgprsInPlace(raw_string_ostream &OS, bool KeepLow, + unsigned Base, unsigned W, unsigned SubW, + unsigned &CurrentMode) { + for (unsigned I = 0; I != W; ++I) { + bool IsLow = ((I / SubW) % 2) == 0; + if (IsLow == KeepLow) + continue; + unsigned Physical = Base + I; + emitModeForOperands(OS, CurrentMode, + {{VgprMsbOperand::Dst, Physical / VgprBankSize}}); + OS << "v_mov_b32 " << encodedVgprName(Physical) << ", 0\n"; + } +} + // Parse a matrix VGPR range from the printer's canonical form. struct VgprRange { unsigned Base; @@ -376,6 +433,757 @@ matrixOperandRange(PatchContext &Ctx, const InternalDecodedInst &DI, return VgprRange{Lo, Hi - Lo + 1}; } +struct Scale32PrintedAsm { + std::string Operands[6]; // dst, matrix A/B, src2, scale A/B + std::string ModifierSuffix; +}; + +// Parse the six positional operands of the M=32 scaled WMMA while preserving +// the printer's exact inline-immediate spelling and modifier ordering. +static std::optional +parseScale32PrintedAsm(PatchContext &Ctx, const InternalDecodedInst &DI) { + SmallString<256> Buf; + raw_svector_ostream OS(Buf); + Ctx.LS.MCIP->printInst(&DI.Inst, /*Address=*/0, /*Annot=*/"", *Ctx.LS.STI, + OS); + StringRef S = StringRef(Buf).trim(); + size_t MnemEnd = S.find_first_of(" \t"); + if (MnemEnd == StringRef::npos) + return std::nullopt; + + Scale32PrintedAsm Result; + StringRef Rest = S.substr(MnemEnd).ltrim(); + for (unsigned I = 0; I != 5; ++I) { + size_t Comma = Rest.find(','); + if (Comma == StringRef::npos) + return std::nullopt; + Result.Operands[I] = Rest.substr(0, Comma).trim().str(); + Rest = Rest.substr(Comma + 1).ltrim(); + } + size_t ModBegin = Rest.find_first_of(" \t"); + if (ModBegin == StringRef::npos) { + Result.Operands[5] = Rest.str(); + } else { + Result.Operands[5] = Rest.substr(0, ModBegin).str(); + Result.ModifierSuffix = Rest.substr(ModBegin).str(); + } + return Result; +} + +static SmallVector tokenizeScaleModifiers(StringRef Suffix) { + SmallVector Result; + StringRef Rest = Suffix.ltrim(); + while (!Rest.empty()) { + size_t Space = Rest.find_first_of(" \t"); + if (Space == StringRef::npos) { + Result.push_back(Rest); + break; + } + Result.push_back(Rest.substr(0, Space)); + Rest = Rest.substr(Space + 1).ltrim(); + } + return Result; +} + +static bool parsePackedScaleModifier(StringRef Token, StringRef Name, + std::array &Bits) { + if (!Token.starts_with(Name) || !Token.ends_with("]")) + return false; + Token = Token.drop_front(Name.size()); + if (!Token.starts_with(":[")) + return false; + SmallVector Parts; + Token.drop_front(2).drop_back(1).split(Parts, ","); + if (Parts.size() != 3) + return false; + for (unsigned I = 0; I != 3; ++I) { + Bits[I] = Parts[I].trim(); + if (Bits[I] != "0" && Bits[I] != "1") + return false; + } + return true; +} + +static bool isKnownScale32Modifier(StringRef Token) { + if (Token == "matrix_a_reuse" || Token == "matrix_b_reuse" || + Token == "matrix_a_scale:MATRIX_SCALE_ROW1" || + Token == "matrix_b_scale:MATRIX_SCALE_ROW1" || + Token == "matrix_a_scale_fmt:MATRIX_SCALE_FMT_E8" || + Token == "matrix_a_scale_fmt:MATRIX_SCALE_FMT_E5M3" || + Token == "matrix_a_scale_fmt:MATRIX_SCALE_FMT_E4M3" || + Token == "matrix_b_scale_fmt:MATRIX_SCALE_FMT_E8" || + Token == "matrix_b_scale_fmt:MATRIX_SCALE_FMT_E5M3" || + Token == "matrix_b_scale_fmt:MATRIX_SCALE_FMT_E4M3") + return true; + std::array Bits; + return parsePackedScaleModifier(Token, "neg_lo", Bits) || + parsePackedScaleModifier(Token, "neg_hi", Bits); +} + +// The split changes both matrix register layout and the K-pass accumulator. +// Reuse promises therefore no longer describe the generated sequence and are +// removed. On each high-K pass src2 is the low-pass result, not the original +// C, so clear the src2 neg/abs bit instead of applying C's modifier twice. +static std::optional +transformScale32ModifierSuffix(StringRef Suffix, bool HighKPass) { + std::string Result; + for (StringRef Token : tokenizeScaleModifiers(Suffix)) { + if (!isKnownScale32Modifier(Token)) { + log() << "hotswap: error: wmma_scale16: unsupported M=32 modifier token " + "\"" + << Token << "\"\n"; + return std::nullopt; + } + if (Token == "matrix_a_reuse" || Token == "matrix_b_reuse") + continue; + + std::array Bits; + if (HighKPass && (parsePackedScaleModifier(Token, "neg_lo", Bits) || + parsePackedScaleModifier(Token, "neg_hi", Bits))) { + if (Bits[0] == "0" && Bits[1] == "0") + continue; + StringRef Name = Token.take_front(Token.find(':')); + Result += (" " + Name + ":[" + Bits[0] + "," + Bits[1] + ",0]").str(); + continue; + } + Result += ' '; + Result += Token.str(); + } + return Result; +} + +static std::string encodedVgprRange(unsigned PhysicalBase, unsigned Width) { + assert(Width > 0); + unsigned EncodedBase = PhysicalBase % VgprBankSize; + return formatv("v[{0}:{1}]", EncodedBase, EncodedBase + Width - 1).str(); +} + +static void emitScale32Half(raw_string_ostream &OS, unsigned DstBase, + unsigned MatrixABase, unsigned MatrixBBase, + StringRef Src2, unsigned ScaleAReg, + unsigned ScaleBReg, StringRef ModifierSuffix) { + OS << "v_wmma_scale_f32_16x16x128_f8f6f4 " << encodedVgprRange(DstBase, 8) + << ", " << encodedVgprRange(MatrixABase, 8) << ", " + << encodedVgprRange(MatrixBBase, 8) << ", " << Src2 << ", " + << encodedVgprName(ScaleAReg) << ", " << encodedVgprName(ScaleBReg) + << " matrix_a_fmt:MATRIX_FMT_FP4 matrix_b_fmt:MATRIX_FMT_FP4" + << ModifierSuffix << "\n"; +} + +// Prefer one exact-liveness-proven dead block inside the kernel's declared +// allocation. Fully allocated production kernels often have no above-KD +// headroom even though a particular WMMA site has a large dead interval. +// Search high-to-low within each physical bank so every generated range keeps +// one VGPR-MSB setting. Fall back to the allocator's normal above-KD growth. +static std::optional +allocContiguousDeadOrAboveInBank(VgprAllocator &Alloc, unsigned Count, + unsigned Align, unsigned BankSize, + bool AllowDeadReuse) { + if (Count == 0 || Count > BankSize || Align == 0) + return std::nullopt; + + if (AllowDeadReuse) { + unsigned BankCount = (Alloc.KdAllocatedVgprs + BankSize - 1) / BankSize; + for (unsigned ReverseBank = BankCount; ReverseBank != 0; --ReverseBank) { + unsigned Bank = ReverseBank - 1; + unsigned BankBegin = Bank * BankSize; + unsigned BankEnd = std::min(Alloc.KdAllocatedVgprs, BankBegin + BankSize); + if (BankEnd < BankBegin + Count) + continue; + + unsigned Base = BankEnd - Count; + Base -= Base % Align; + while (Base >= BankBegin) { + bool AllDead = true; + for (unsigned V = Base; V != Base + Count; ++V) { + if (V >= static_cast(Alloc.LiveAtPoint.size()) || + Alloc.LiveAtPoint.test(V)) { + AllDead = false; + break; + } + } + if (AllDead) { + Alloc.LiveAtPoint.set(Base, Base + Count); + return Base; + } + if (Base < BankBegin + Align) + break; + Base -= Align; + } + } + } + return Alloc.allocContiguousAboveKdInBank(Count, Align, BankSize); +} + +static bool rangesOverlap(unsigned ABase, unsigned AWidth, unsigned BBase, + unsigned BWidth) { + return ABase < BBase + BWidth && BBase < ABase + AWidth; +} + +static bool sameRange(unsigned ABase, unsigned AWidth, unsigned BBase, + unsigned BWidth) { + return ABase == BBase && AWidth == BWidth; +} + +static void reserveVgprRange(VgprAllocator &Alloc, unsigned Base, + unsigned Width) { + assert(Width > 0 && Base <= Alloc.LiveAtPoint.size() && + Width <= Alloc.LiveAtPoint.size() - Base); + Alloc.LiveAtPoint.set(Base, Base + Width); +} + +struct EncodedVgprRange { + unsigned Base = 0; + unsigned Width = 0; + bool FullDwords = false; +}; + +static std::optional parseScalarVgprName(StringRef Name, + bool &IsPartial) { + IsPartial = false; + if (!Name.consume_front("VGPR")) + return std::nullopt; + size_t Digits = Name.find_first_not_of("0123456789"); + StringRef Number = Digits == StringRef::npos ? Name : Name.take_front(Digits); + unsigned Index = 0; + if (Number.empty() || Number.getAsInteger(10, Index)) + return std::nullopt; + if (Digits == StringRef::npos) + return Index; + StringRef Suffix = Name.drop_front(Digits); + if (Suffix == "_LO16" || Suffix == "_HI16") { + IsPartial = true; + return Index; + } + return std::nullopt; +} + +// Convert an explicit MC VGPR or VGPR tuple to its encoded v0..v255 interval. +// True16 operands are identified but never accepted as a full-value kill. +static std::optional +getEncodedVgprRange(MCRegister Reg, const MCRegisterInfo &MRI) { + if (!Reg) + return std::nullopt; + + bool IsPartial = false; + if (std::optional Scalar = + parseScalarVgprName(MRI.getName(Reg), IsPartial)) + return EncodedVgprRange{*Scalar, 1, !IsPartial}; + + SmallVector Scalars; + for (MCPhysReg Sub : MRI.subregs(Reg)) { + bool SubPartial = false; + std::optional Index = + parseScalarVgprName(MRI.getName(Sub), SubPartial); + if (Index && !SubPartial) + Scalars.push_back(*Index); + } + if (Scalars.empty()) + return std::nullopt; + llvm::sort(Scalars); + Scalars.erase(std::unique(Scalars.begin(), Scalars.end()), Scalars.end()); + for (unsigned I = 1; I != Scalars.size(); ++I) + if (Scalars[I] != Scalars.front() + I) + return std::nullopt; + return EncodedVgprRange{Scalars.front(), + static_cast(Scalars.size()), true}; +} + +bool isVectorRegisterOrAlias(MCRegister Reg, const MCRegisterInfo &MRI) { + if (!Reg) + return false; + for (MCRegAliasIterator Alias(Reg, &MRI, /*IncludeSelf=*/true); + Alias.isValid(); ++Alias) { + StringRef Name = MRI.getName(*Alias); + if (Name.contains("VGPR") || Name.contains("AGPR")) + return true; + } + return false; +} + +static bool setPhysicalVgprRange(BitVector &Out, const EncodedVgprRange &Range, + unsigned Bank, unsigned MaxVgprs) { + if (Range.Width == 0 || Range.Base >= VgprBankSize || + Range.Width > VgprBankSize - Range.Base) + return false; + unsigned Base = Range.Base + Bank * VgprBankSize; + if (Base >= MaxVgprs || Range.Width > MaxVgprs - Base) + return false; + Out.set(Base, Base + Range.Width); + return true; +} + +struct PhysicalVgprAccess { + BitVector Uses; + BitVector FullDefs; + bool Valid = true; + + explicit PhysicalVgprAccess(unsigned MaxVgprs) + : Uses(MaxVgprs), FullDefs(MaxVgprs) {} +}; + +// The M=32 Scale16 MC layout is mirrored and runtime-validated by the +// lowering below. Its two scale operands reuse the matrix source banks: +// matrix-A/scale-A use src0, matrix-B/scale-B use src1, and the accumulator +// uses src2. Other instruction layouts remain conservative-all-banks unless +// their role is structurally unambiguous. +static std::optional +getExactSourceRole(const InternalDecodedInst &DI, unsigned OperandIndex, + unsigned NumDefs) { + if (DI.Mnemonic == "v_wmma_scale16_f32_32x16x128_f4") { + switch (OperandIndex) { + case 1: + case 5: + return VgprMsbOperand::Src0; + case 2: + case 6: + return VgprMsbOperand::Src1; + case 4: + return VgprMsbOperand::Src2; + default: + return std::nullopt; + } + } + if (StringRef(DI.Mnemonic).starts_with("ds_") && OperandIndex == NumDefs) + return VgprMsbOperand::Src0; + return std::nullopt; +} + +// Resolve every explicit access through the exact persistent VGPR-MSB mode. +// Sources with a validated architectural role use that role's bank. Every +// other source maps through the union of src0/src1/src2 banks, which can only +// add uses, never hide one. Explicit full-width definitions use the +// architectural dst bank. Tied definitions are reads of their incoming +// destination. Implicit/partial VGPR operands cannot prove a kill and +// conservatively block the encoded range in every bank. +static PhysicalVgprAccess getPhysicalVgprAccess(const InternalDecodedInst &DI, + const LLVMState &LS, + unsigned Mode, + unsigned MaxVgprs) { + PhysicalVgprAccess Result(MaxVgprs); + const MCInstrDesc &Desc = LS.MCII->get(DI.Inst.getOpcode()); + const MCRegisterInfo &MRI = *LS.MRI; + unsigned DstBank = getVgprMsbBank(Mode, VgprMsbOperand::Dst); + SmallVector SrcBanks = { + getVgprMsbBank(Mode, VgprMsbOperand::Src0), + getVgprMsbBank(Mode, VgprMsbOperand::Src1), + getVgprMsbBank(Mode, VgprMsbOperand::Src2)}; + llvm::sort(SrcBanks); + SrcBanks.erase(std::unique(SrcBanks.begin(), SrcBanks.end()), SrcBanks.end()); + + auto AddEveryBankUse = [&](const EncodedVgprRange &Range) { + for (unsigned Bank = 0; Bank * VgprBankSize < MaxVgprs; ++Bank) + if (!setPhysicalVgprRange(Result.Uses, Range, Bank, MaxVgprs)) + Result.Valid = false; + }; + + unsigned NumDefs = Desc.getNumDefs(); + for (unsigned I = 0, E = DI.Inst.getNumOperands(); I != E; ++I) { + const MCOperand &Op = DI.Inst.getOperand(I); + if (!Op.isReg() || !Op.getReg()) + continue; + std::optional Range = + getEncodedVgprRange(MCRegister(Op.getReg()), MRI); + if (!Range) { + if (isVectorRegisterOrAlias(MCRegister(Op.getReg()), MRI)) + Result.Valid = false; + continue; + } + + bool IsDef = I < NumDefs; + if (!IsDef) { + if (std::optional Role = + getExactSourceRole(DI, I, NumDefs)) { + unsigned Bank = getVgprMsbBank(Mode, *Role); + if (!setPhysicalVgprRange(Result.Uses, *Range, Bank, MaxVgprs)) + Result.Valid = false; + } else { + for (unsigned Bank : SrcBanks) + if (!setPhysicalVgprRange(Result.Uses, *Range, Bank, MaxVgprs)) + Result.Valid = false; + } + continue; + } + + int TiedTo = Desc.getOperandConstraint(I, MCOI::TIED_TO); + if (TiedTo >= 0) { + if (!setPhysicalVgprRange(Result.Uses, *Range, DstBank, MaxVgprs)) + Result.Valid = false; + } + if (!Range->FullDwords) { + AddEveryBankUse(*Range); + continue; + } + if (!setPhysicalVgprRange(Result.FullDefs, *Range, DstBank, MaxVgprs)) + Result.Valid = false; + } + + for (MCPhysReg Implicit : Desc.implicit_uses()) { + std::optional Range = + getEncodedVgprRange(MCRegister(Implicit), MRI); + if (Range) + AddEveryBankUse(*Range); + else if (isVectorRegisterOrAlias(MCRegister(Implicit), MRI)) + Result.Valid = false; + } + for (MCPhysReg Implicit : Desc.implicit_defs()) { + std::optional Range = + getEncodedVgprRange(MCRegister(Implicit), MRI); + if (Range) + AddEveryBankUse(*Range); + else if (isVectorRegisterOrAlias(MCRegister(Implicit), MRI)) + Result.Valid = false; + } + return Result; +} + +static bool hasDynamicVgprAddressing(ArrayRef Decoded, + size_t Begin, size_t End) { + for (size_t I = Begin; I != End; ++I) { + StringRef Mnemonic = Decoded[I].Mnemonic; + if (Mnemonic.contains("movrel") || Mnemonic.contains("gpr_idx") || + Mnemonic.starts_with("s_setreg")) + return true; + } + return false; +} + +std::optional +computeForwardDeadVgprs(ArrayRef Nodes, size_t EntryNode, + unsigned MaxVgprs) { + if (Nodes.empty() || EntryNode >= Nodes.size() || MaxVgprs == 0) + return std::nullopt; + for (const ForwardVgprProofNode &Node : Nodes) { + if (Node.Uses.size() != MaxVgprs || Node.FullDefs.size() != MaxVgprs) + return std::nullopt; + for (size_t Successor : Node.Successors) + if (Successor >= Nodes.size()) + return std::nullopt; + } + + std::vector AliveAt(Nodes.size(), BitVector(MaxVgprs)); + AliveAt[EntryNode].set(); + BitVector Unsafe(MaxVgprs); + SmallVector Worklist; + Worklist.push_back(EntryNode); + + while (!Worklist.empty()) { + size_t Index = Worklist.pop_back_val(); + BitVector Alive = AliveAt[Index]; + if (Alive.none()) + continue; + + const ForwardVgprProofNode &Node = Nodes[Index]; + if (Node.Opaque) { + Unsafe |= Alive; + continue; + } + + BitVector UsedAlive = Alive; + UsedAlive &= Node.Uses; + Unsafe |= UsedAlive; + Alive.reset(Node.Uses); + Alive.reset(Node.FullDefs); + + if (Node.HasUnsafeExit) + Unsafe |= Alive; + if (Node.Successors.empty()) { + if (!Node.SafeTerminal) + Unsafe |= Alive; + continue; + } + + for (size_t Successor : Node.Successors) { + BitVector NewBits = Alive; + NewBits.reset(AliveAt[Successor]); + if (NewBits.none()) + continue; + AliveAt[Successor] |= Alive; + Worklist.push_back(Successor); + } + } + + // A value that can circulate around a cycle without a use or full kill is + // not accepted as scratch. Detect such cycles in the per-value subgraph: + // Kahn removal leaves exactly the nodes belonging to, or fed only by, a + // surviving cycle. This is intentionally conservative for non-terminating + // paths and makes loop handling independent of worklist visitation order. + SmallVector InDegree(Nodes.size()); + SmallVector Queue; + BitVector Included(Nodes.size()); + for (unsigned V = 0; V != MaxVgprs; ++V) { + if (Unsafe.test(V)) + continue; + Included.reset(); + unsigned IncludedCount = 0; + for (size_t I = 0; I != Nodes.size(); ++I) { + const ForwardVgprProofNode &Node = Nodes[I]; + if (AliveAt[I].test(V) && !Node.Opaque && !Node.Uses.test(V) && + !Node.FullDefs.test(V)) { + Included.set(I); + ++IncludedCount; + } + } + if (IncludedCount == 0) + continue; + + llvm::fill(InDegree, 0); + for (int I = Included.find_first(); I >= 0; I = Included.find_next(I)) + for (size_t Successor : Nodes[static_cast(I)].Successors) + if (Included.test(Successor)) + ++InDegree[Successor]; + Queue.clear(); + for (int I = Included.find_first(); I >= 0; I = Included.find_next(I)) + if (InDegree[static_cast(I)] == 0) + Queue.push_back(static_cast(I)); + + unsigned Removed = 0; + while (!Queue.empty()) { + size_t I = Queue.pop_back_val(); + ++Removed; + for (size_t Successor : Nodes[I].Successors) + if (Included.test(Successor) && --InDegree[Successor] == 0) + Queue.push_back(Successor); + } + if (Removed != IncludedCount) + Unsafe.set(V); + } + + BitVector Safe(MaxVgprs); + Safe.set(); + Safe.reset(Unsafe); + return Safe; +} + +struct ScaleForwardGraph { + std::vector Nodes; + std::vector GlobalIndices; + std::vector ModeBefore; + BitVector UndecodedVop3Heads; + size_t EntryNode = 0; +}; + +// The B0 f4gemm object contains a legacy VOP3 opcode that the A0 gfx1250 MC +// decoder intentionally does not recognize. The failed decode advances only +// one dword, so its second dword appears as a separate unknown instruction. +// +// Recognize only the exact observed encoding: +// * legacy VOP3 major 0x34 and opcode 0x31; +// * vdst encoding zero, with only the observed bit-14 modifier variation; +// * scalar source encodings, in either observed second-dword spelling. +// +// We do not assign the unknown opcode semantics. Instead, conservatively model +// encoded v0 as a use in every physical bank (so it cannot be scratch), model +// no kill, and skip the split continuation dword. This is enough to traverse +// the instruction without relying on whether opcode 0x31 has a vector or +// scalar destination on B0. +static bool recognizeUndecodedB0Vop3ScalarSources( + PatchContext &Ctx, size_t Global, unsigned MaxVgprs, + BitVector &ConservativeUses) { + if (Global + 1 >= Ctx.Decoded.size()) + return false; + const InternalDecodedInst &Head = Ctx.Decoded[Global]; + const InternalDecodedInst &Tail = Ctx.Decoded[Global + 1]; + if (Head.DecodeSucceeded || Head.Size != MinInstSize || + Tail.Offset != Head.Offset + MinInstSize || + Head.Offset > Ctx.TextSize || 2 * MinInstSize > Ctx.TextSize - Head.Offset) + return false; + + const uint8_t *Raw = Ctx.Text + Head.Offset; + uint32_t Word0 = support::endian::read32le(Raw); + uint32_t Word1 = support::endian::read32le(Raw + MinInstSize); + constexpr uint32_t Bit14 = 1u << 14; + if ((Word0 & ~Bit14) != 0xd0310000u || + (Word1 != 0 && Word1 != 0x00100000u)) + return false; + + // A direct/declared entry into the decoder-created continuation slot would + // make treating the pair as one instruction unsound. + if (Ctx.DirectControlFlow.Targets.contains(Tail.Offset) || + llvm::is_contained(Ctx.DeclaredEntries, Tail.Offset)) + return false; + + for (unsigned Physical = 0; Physical < MaxVgprs; + Physical += VgprBankSize) + ConservativeUses.set(Physical); + return true; +} + +static std::optional +buildScaleForwardGraph(PatchContext &Ctx, size_t SiteIdx, unsigned EntryMode, + unsigned MaxVgprs) { + if (!Ctx.LS.MIA || !Ctx.LS.MCII || !Ctx.LS.MRI || + SiteIdx >= Ctx.Decoded.size()) + return std::nullopt; + + std::optional Owner = + Ctx.Elf.findFunctionTextRangeAtOffset(Ctx.Decoded[SiteIdx].Offset); + if (!Owner || SiteIdx + 1 >= Ctx.Decoded.size()) + return std::nullopt; + + size_t BeginIndex = SiteIdx; + while (BeginIndex > 0 && Ctx.Decoded[BeginIndex - 1].Offset >= Owner->Begin) + --BeginIndex; + size_t EndIndex = SiteIdx + 1; + while (EndIndex < Ctx.Decoded.size() && + Ctx.Decoded[EndIndex].Offset < Owner->End) + ++EndIndex; + if (BeginIndex == EndIndex || SiteIdx + 1 >= EndIndex) + return std::nullopt; + + ScaleForwardGraph Graph; + size_t Count = EndIndex - BeginIndex; + Graph.Nodes.reserve(Count); + Graph.GlobalIndices.reserve(Count); + for (size_t I = BeginIndex; I != EndIndex; ++I) { + Graph.Nodes.emplace_back(MaxVgprs); + Graph.GlobalIndices.push_back(I); + } + Graph.UndecodedVop3Heads.resize(Count); + Graph.EntryNode = SiteIdx + 1 - BeginIndex; + + DenseMap IndexAtOffset; + for (size_t Local = 0; Local != Count; ++Local) + IndexAtOffset[Ctx.Decoded[Graph.GlobalIndices[Local]].Offset] = Local; + + auto AddFallthrough = [&](size_t Local) { + if (Local + 1 < Count) + Graph.Nodes[Local].Successors.push_back(Local + 1); + else + Graph.Nodes[Local].HasUnsafeExit = true; + }; + + for (size_t Local = 0; Local != Count; ++Local) { + size_t Global = Graph.GlobalIndices[Local]; + const InternalDecodedInst &DI = Ctx.Decoded[Global]; + ForwardVgprProofNode &Node = Graph.Nodes[Local]; + + // A later loop iteration reaches the replacement itself. Scratch excludes + // every original operand, and the replacement defines its scratch before + // reading it, so no incoming scratch value can be observed there. + if (Global == SiteIdx) { + Node.SafeTerminal = true; + continue; + } + if (!DI.DecodeSucceeded && + recognizeUndecodedB0Vop3ScalarSources(Ctx, Global, MaxVgprs, + Node.Uses)) { + if (Local + 2 >= Count) + Node.HasUnsafeExit = true; + else + Node.Successors.push_back(Local + 2); + Graph.UndecodedVop3Heads.set(Local); + continue; + } + if (!DI.DecodeSucceeded || + hasDynamicVgprAddressing(Ctx.Decoded, Global, Global + 1)) { + Node.Opaque = true; + continue; + } + if (DI.Inst.getOpcode() == Ctx.LS.SEndPgmOpcode || + DI.Inst.getOpcode() == Ctx.LS.SEndPgmSavedOpcode) { + Node.SafeTerminal = true; + continue; + } + if (Ctx.LS.MIA->isCall(DI.Inst) || Ctx.LS.MIA->isIndirectBranch(DI.Inst) || + Ctx.LS.MIA->isReturn(DI.Inst)) { + Node.Opaque = true; + continue; + } + if (Ctx.LS.MIA->isBranch(DI.Inst)) { + uint64_t Target = 0; + if (!Ctx.LS.MIA->evaluateBranch(DI.Inst, DI.Offset, DI.Size, Target)) { + Node.Opaque = true; + continue; + } + DenseMap::const_iterator TargetIt = + IndexAtOffset.find(Target); + if (TargetIt == IndexAtOffset.end()) + Node.HasUnsafeExit = true; + else + Node.Successors.push_back(TargetIt->second); + if (Ctx.LS.MIA->isConditionalBranch(DI.Inst)) + AddFallthrough(Local); + else if (!Ctx.LS.MIA->isUnconditionalBranch(DI.Inst)) + Node.Opaque = true; + continue; + } + if (Ctx.LS.MIA->mayAffectControlFlow(DI.Inst, *Ctx.LS.MRI)) { + Node.Opaque = true; + continue; + } + AddFallthrough(Local); + } + + Graph.ModeBefore.assign(Count, VgprMsbUnreachable); + Graph.ModeBefore[Graph.EntryNode] = static_cast(EntryMode & 0xff); + SmallVector Worklist; + Worklist.push_back(Graph.EntryNode); + for (size_t Next = 0; Next != Worklist.size(); ++Next) { + size_t Local = Worklist[Next]; + const ForwardVgprProofNode &Node = Graph.Nodes[Local]; + if (Node.Opaque || Node.SafeTerminal) + continue; + int16_t Out = + Graph.UndecodedVop3Heads.test(Local) + ? Graph.ModeBefore[Local] + : transferExactVgprMsbMode( + Graph.ModeBefore[Local], + Ctx.Decoded[Graph.GlobalIndices[Local]], Ctx.LS); + for (size_t Successor : Node.Successors) { + int16_t Old = Graph.ModeBefore[Successor]; + int16_t Merged = Old == VgprMsbUnreachable ? Out + : Old == Out ? Old + : VgprMsbUnknown; + if (Merged != Old) { + Graph.ModeBefore[Successor] = Merged; + Worklist.push_back(Successor); + } + } + } + return Graph; +} + +// Return physical VGPR values whose incoming contents cannot be observed on +// any continuation path after SiteIdx. This deliberately does not consume the +// generic LivenessInfo: its weak in-tree implementation is encoded-v0..v255 +// conservative liveness, not a proof over gfx1250's four physical banks. +static std::optional +computeForwardDeadPhysicalVgprs(PatchContext &Ctx, size_t SiteIdx, + unsigned EntryMode, unsigned MaxVgprs) { + if (Ctx.DirectControlFlow.HasUnresolvedTargets || + Ctx.DirectControlFlow.HasUnboundedIndirectEntries || !Ctx.LS.MIA || + !Ctx.LS.MCII || !Ctx.LS.MRI || SiteIdx >= Ctx.Decoded.size()) + return std::nullopt; + + std::optional Graph = + buildScaleForwardGraph(Ctx, SiteIdx, EntryMode, MaxVgprs); + if (!Graph) + return std::nullopt; + + for (size_t Local = 0; Local != Graph->Nodes.size(); ++Local) { + ForwardVgprProofNode &Node = Graph->Nodes[Local]; + if (Node.Opaque || Node.SafeTerminal) + continue; + if (Graph->UndecodedVop3Heads.test(Local)) + continue; + + int16_t Mode = Graph->ModeBefore[Local]; + if (Mode == VgprMsbUnreachable) + continue; + unsigned AccessMode = Mode >= 0 ? static_cast(Mode) : 0; + PhysicalVgprAccess Access = getPhysicalVgprAccess( + Ctx.Decoded[Graph->GlobalIndices[Local]], Ctx.LS, AccessMode, MaxVgprs); + if (!Access.Valid) + return std::nullopt; + if (Mode < 0 && (Access.Uses.any() || Access.FullDefs.any())) + return std::nullopt; + Node.Uses = std::move(Access.Uses); + Node.FullDefs = std::move(Access.FullDefs); + } + return computeForwardDeadVgprs(Graph->Nodes, Graph->EntryNode, MaxVgprs); +} + // Matrix-A K-subblock masking scheme, chosen by the matrix-A data format. // The K-split must isolate each 16-K subblock, and how a subblock maps to // lanes/VGPRs is format-dependent: @@ -674,7 +1482,8 @@ static uint32_t patchWmmaScale16_16x16(PatchContext &Ctx, size_t Idx) { ScratchPatchInfo Info; Info.Offset = DI.Offset; - Info.ScratchRegs = Alloc.LiveAtPoint; + Info.ScratchRegs.resize(Ctx.Config.MaxVgprs); + Info.ScratchRegs.set(*ScratchBase, *ScratchBase + ScratchCount); Ctx.OutScratchPatches.push_back(std::move(Info)); log() << "hotswap: wmma_scale16: exact K-split at offset 0x" @@ -687,6 +1496,334 @@ static uint32_t patchWmmaScale16_16x16(PatchContext &Ctx, size_t Idx) { return 1; } +// --------------------------------------------------------------------------- +// v_wmma_scale16_f32_32x16x128_f4 -> exact M+K split +// --------------------------------------------------------------------------- + +static uint32_t patchWmmaScale16_32x16(PatchContext &Ctx, size_t Idx) { + const InternalDecodedInst &DI = Ctx.Decoded[Idx]; + + if (DI.Size != VOP3PXSize) + return failClosed(Ctx, DI, "unexpected instruction size " + Twine(DI.Size)); + for (const Trampoline &T : Ctx.OutTrampolines) + if (T.OriginalOffset == DI.Offset) + return 0; + + const uint8_t *Raw = Ctx.Text + DI.Offset; + std::optional ScaleABase = + decodeVgprEncoding(extractScaleSrc0(Raw)); + std::optional ScaleBBase = + decodeVgprEncoding(extractScaleSrc1(Raw)); + if (!ScaleABase || !ScaleBBase) + return failClosed(Ctx, DI, "non-VGPR block-16 scale operand"); + + std::optional ActiveMode = getActiveVgprMsbMode(Ctx, Idx); + if (!ActiveMode) + ActiveMode = getLocallyEstablishedVgprMsbMode(Ctx, Idx); + if (!ActiveMode) { + std::string Detail = "cannot determine active VGPR-MSB mode"; + if (Ctx.DirectControlFlow.HasUnresolvedTargets) + Detail += " (unresolved control-flow target)"; + if (Ctx.DirectControlFlow.HasUnboundedIndirectEntries) + Detail += " (unbounded indirect entry)"; + return failClosed(Ctx, DI, Detail); + } + + std::optional Printed = parseScale32PrintedAsm(Ctx, DI); + if (!Printed) + return failClosed(Ctx, DI, "could not parse canonical instruction"); + std::optional LowSuffix = + transformScale32ModifierSuffix(Printed->ModifierSuffix, + /*HighKPass=*/false); + std::optional HighSuffix = + transformScale32ModifierSuffix(Printed->ModifierSuffix, + /*HighKPass=*/true); + if (!LowSuffix || !HighSuffix) + return failClosed(Ctx, DI, "unsupported modifier combination"); + + std::optional DRange = + matrixOperandRange(Ctx, DI, /*OperandIndex=*/0); + std::optional ARange = + matrixOperandRange(Ctx, DI, /*OperandIndex=*/1); + std::optional BRange = + matrixOperandRange(Ctx, DI, /*OperandIndex=*/2); + if (!DRange || !ARange || !BRange || DRange->Width != 16 || + ARange->Width != 16 || BRange->Width != 8) + return failClosed(Ctx, DI, "unexpected M=32 matrix operand widths/layout"); + + // The M=32 profile mirrors the common VOP3P layout in its first five MC + // operands: vdst, src0, src1, src2 modifiers, src2. Validate that mirror at + // runtime so a TableGen layout change fails closed. + if (DI.Inst.getNumOperands() < 5) + return failClosed(Ctx, DI, "truncated M=32 MC operand layout"); + const MCOperand &Src2Op = DI.Inst.getOperand(4); + bool Src2IsImm = Src2Op.isImm(); + std::optional CRange; + if (Src2Op.isReg()) + CRange = matrixOperandRange(Ctx, DI, /*OperandIndex=*/3); + else if (!Src2IsImm) + return failClosed(Ctx, DI, "unsupported non-VGPR/non-immediate src2"); + if (CRange && CRange->Width != 16) + return failClosed(Ctx, DI, "src2 and destination widths differ"); + if (!Src2IsImm && !CRange) + return failClosed(Ctx, DI, "could not determine src2 VGPR range"); + + unsigned Src0Bank = getVgprMsbBank(*ActiveMode, VgprMsbOperand::Src0); + unsigned Src1Bank = getVgprMsbBank(*ActiveMode, VgprMsbOperand::Src1); + unsigned Src2Bank = getVgprMsbBank(*ActiveMode, VgprMsbOperand::Src2); + unsigned DstBank = getVgprMsbBank(*ActiveMode, VgprMsbOperand::Dst); + + unsigned DBase = DRange->Base + DstBank * VgprBankSize; + unsigned ABase = ARange->Base + Src0Bank * VgprBankSize; + unsigned BBase = BRange->Base + Src1Bank * VgprBankSize; + unsigned CBase = CRange ? CRange->Base + Src2Bank * VgprBankSize : 0; + unsigned ScaleALo = *ScaleABase + Src0Bank * VgprBankSize; + unsigned ScaleAHi = ScaleALo + 1; + unsigned ScaleBLo = *ScaleBBase + Src1Bank * VgprBankSize; + unsigned ScaleBHi = ScaleBLo + 1; + + if (!physicalVgprRangeFitsOneBank(DBase, 16, Ctx.Config.MaxVgprs) || + !physicalVgprRangeFitsOneBank(ABase, 16, Ctx.Config.MaxVgprs) || + !physicalVgprRangeFitsOneBank(BBase, 8, Ctx.Config.MaxVgprs) || + (CRange && + !physicalVgprRangeFitsOneBank(CBase, 16, Ctx.Config.MaxVgprs)) || + !physicalVgprRangeFitsOneBank(ScaleALo, 2, Ctx.Config.MaxVgprs) || + !physicalVgprRangeFitsOneBank(ScaleBLo, 2, Ctx.Config.MaxVgprs)) + return failClosed(Ctx, DI, + "M=32 operand exceeds or crosses a physical VGPR bank"); + + // The split reads A in stages after its first D-half write. The fused source + // instruction reads all of A before writing D, so any D/A overlap would + // otherwise let the replacement destroy a later A read. + if (rangesOverlap(DBase, 16, ABase, 16)) + return failClosed(Ctx, DI, + "destination overlaps matrix A across staged reads"); + if (rangesOverlap(DBase, 16, BBase, 8)) + return failClosed(Ctx, DI, + "destination overlaps matrix B across staged reads"); + if (rangesOverlap(ABase, 16, BBase, 8)) + return failClosed(Ctx, DI, + "matrix A overlaps matrix B during in-place masking"); + + // Exact D==C is the ordinary in-place accumulator form. A disjoint C is + // likewise safe. Reject partial/cross-half overlap: writing DLo could + // otherwise destroy CHi before the second low-K pass consumes it. + if (CRange && rangesOverlap(DBase, 16, CBase, 16) && + !sameRange(DBase, 16, CBase, 16)) + return failClosed(Ctx, DI, + "partial destination/src2 overlap across staged reads"); + + auto ScaleOverlaps = [&](unsigned Base, unsigned Width) { + return rangesOverlap(ScaleALo, 2, Base, Width) || + rangesOverlap(ScaleBLo, 2, Base, Width); + }; + if (ScaleOverlaps(DBase, 16) || ScaleOverlaps(ABase, 16) || + ScaleOverlaps(BBase, 8) || (CRange && ScaleOverlaps(CBase, 16)) || + rangesOverlap(ScaleALo, 2, ScaleBLo, 2)) + return failClosed( + Ctx, DI, + "scale pair overlaps a staged matrix operand or the other scale"); + if (CRange && rangesOverlap(ABase, 16, CBase, 16)) + return failClosed(Ctx, DI, + "matrix A overlaps src2 during in-place masking"); + + std::string KernelName = + Ctx.Elf.findKernelAtAddress(DI.Offset + Ctx.Elf.textAddr()); + std::optional KdVgprs = Ctx.Elf.getKernelVgprCount( + KernelName, getKernelVgprGranuleSize(Ctx, KernelName)); + unsigned KdCount = KdVgprs.value_or(Ctx.Config.MaxVgprs); + VgprAllocator Alloc(Ctx.Liveness.liveBefore(Idx), KdCount, + Ctx.Config.MaxVgprs); + + // Replace generic encoded-register liveness with a physical-bank all-path + // proof when available. An unset bit is the only state the in-KD allocator + // accepts; failure leaves the allocator conservative-all-live and therefore + // permits only ordinary above-KD growth. + std::optional ForwardDead = computeForwardDeadPhysicalVgprs( + Ctx, Idx, *ActiveMode, Ctx.Config.MaxVgprs); + if (ForwardDead) { + for (int V = ForwardDead->find_first(); V >= 0; + V = ForwardDead->find_next(V)) + Alloc.LiveAtPoint.reset(static_cast(V)); + unsigned BestBase = 0; + unsigned BestWidth = 0; + for (unsigned BankBase = 0; BankBase < Ctx.Config.MaxVgprs; + BankBase += VgprBankSize) { + unsigned BankEnd = std::min(Ctx.Config.MaxVgprs, BankBase + VgprBankSize); + for (unsigned V = BankBase; V != BankEnd;) { + if (!ForwardDead->test(V)) { + ++V; + continue; + } + unsigned Begin = V; + while (V != BankEnd && ForwardDead->test(V)) + ++V; + if (V - Begin > BestWidth) { + BestBase = Begin; + BestWidth = V - Begin; + } + } + } + log() << "hotswap: wmma_scale16: physical forward-dead proof at offset 0x" + << utohexstr(DI.Offset) << " found " << ForwardDead->count() + << " VGPRs; longest single-bank run "; + if (BestWidth) + log() << "v" << BestBase << ":" << (BestBase + BestWidth - 1) << " (" + << BestWidth << ")\n"; + else + log() << "\n"; + } else { + log() << "hotswap: wmma_scale16: physical forward-dead proof unavailable " + "at offset 0x" + << utohexstr(DI.Offset) << "\n"; + } + + // Liveness describes values entering the original instruction. Its + // destination can therefore appear dead even though every replacement + // writes it, and tied/overlapping inputs need the same protection. Reserve + // every physical VGPR range decoded from the original instruction before + // considering an in-KD scratch block. + reserveVgprRange(Alloc, DBase, 16); + reserveVgprRange(Alloc, ABase, 16); + reserveVgprRange(Alloc, BBase, 8); + if (CRange) + reserveVgprRange(Alloc, CBase, 16); + reserveVgprRange(Alloc, ScaleALo, 2); + reserveVgprRange(Alloc, ScaleBLo, 2); + + // Masking one eight-register FP4 A half overwrites exactly four registers. + // Preserve only those four values, not the four already-retained values. + // The first save slot doubles as the reversible scale-pair permutation + // temporary before the first A save and after the final A restore. Low-K + // runs for both M halves before high-K, so the same four slots serve all four + // WMMAs. Matrix B stays in place. + constexpr unsigned MatrixHalfWidth = 8; + constexpr unsigned SavedARegCount = MatrixHalfWidth / 2; + constexpr unsigned ScratchCount = SavedARegCount; + std::array SavedARegs; + for (unsigned &Reg : SavedARegs) { + std::optional Allocated = allocContiguousDeadOrAboveInBank( + Alloc, /*Count=*/1, /*Align=*/1, VgprBankSize, + ForwardDead.has_value()); + if (!Allocated) + return failClosed(Ctx, DI, + "fewer than four dead/above-KD VGPRs for exact M+K " + "split"); + Reg = *Allocated; + } + unsigned TmpReg = SavedARegs.front(); + + std::string ReplacementAsm; + raw_string_ostream OS(ReplacementAsm); + unsigned CurrentMode = *ActiveMode; + emitScalePairSplitInPlace(OS, ScaleALo, ScaleAHi, TmpReg, CurrentMode); + emitScalePairSplitInPlace(OS, ScaleBLo, ScaleBHi, TmpReg, CurrentMode); + + int HazardNops = classifyWmmaNops("v_wmma_scale_f32_16x16x128_f8f6f4").A0Nops; + auto EmitHazardNops = [&] { + for (int I = 0; I != HazardNops; ++I) + OS << "v_nop\n"; + }; + + for (bool HighK : {false, true}) { + for (unsigned MHalf = 0; MHalf != 2; ++MHalf) { + unsigned OriginalAHalf = ABase + MHalf * MatrixHalfWidth; + unsigned DstHalf = DBase + MHalf * MatrixHalfWidth; + unsigned ABank = OriginalAHalf / VgprBankSize; + unsigned BBank = BBase / VgprBankSize; + + unsigned SavedIndex = 0; + for (unsigned I = 0; I != MatrixHalfWidth; ++I) { + bool IsLow = ((I / 2) % 2) == 0; + if (IsLow == !HighK) + continue; + unsigned Saved = SavedARegs[SavedIndex++]; + emitVgprCopy(OS, Saved, OriginalAHalf + I, /*W=*/1, + Saved / VgprBankSize, CurrentMode); + } + assert(SavedIndex == SavedARegCount); + emitMaskVgprsInPlace(OS, /*KeepLow=*/!HighK, OriginalAHalf, + MatrixHalfWidth, /*SubW=*/2, CurrentMode); + + SmallVector WmmaMode = { + {VgprMsbOperand::Dst, DstHalf / VgprBankSize}, + {VgprMsbOperand::Src0, ABank}, + {VgprMsbOperand::Src1, BBank}}; + std::string Src2; + if (HighK) { + WmmaMode.push_back({VgprMsbOperand::Src2, DstHalf / VgprBankSize}); + Src2 = encodedVgprRange(DstHalf, MatrixHalfWidth); + } else if (CRange) { + unsigned CHalf = CBase + MHalf * MatrixHalfWidth; + WmmaMode.push_back({VgprMsbOperand::Src2, CHalf / VgprBankSize}); + Src2 = encodedVgprRange(CHalf, MatrixHalfWidth); + } else { + Src2 = Printed->Operands[3]; + } + emitModeForOperands(OS, CurrentMode, WmmaMode); + emitScale32Half(OS, DstHalf, OriginalAHalf, BBase, Src2, + HighK ? ScaleAHi : ScaleALo, HighK ? ScaleBHi : ScaleBLo, + HighK ? *HighSuffix : *LowSuffix); + + EmitHazardNops(); + SavedIndex = 0; + for (unsigned I = 0; I != MatrixHalfWidth; ++I) { + bool IsLow = ((I / 2) % 2) == 0; + if (IsLow == !HighK) + continue; + unsigned Saved = SavedARegs[SavedIndex++]; + emitVgprCopy(OS, OriginalAHalf + I, Saved, /*W=*/1, ABank, + CurrentMode); + } + assert(SavedIndex == SavedARegCount); + } + } + + emitScalePairRestoreInPlace(OS, ScaleALo, ScaleAHi, TmpReg, CurrentMode); + emitScalePairRestoreInPlace(OS, ScaleBLo, ScaleBHi, TmpReg, CurrentMode); + emitModeForOperands(OS, CurrentMode, + {{VgprMsbOperand::Src0, + getVgprMsbBank(*ActiveMode, VgprMsbOperand::Src0)}, + {VgprMsbOperand::Src1, + getVgprMsbBank(*ActiveMode, VgprMsbOperand::Src1)}, + {VgprMsbOperand::Src2, + getVgprMsbBank(*ActiveMode, VgprMsbOperand::Src2)}, + {VgprMsbOperand::Dst, + getVgprMsbBank(*ActiveMode, VgprMsbOperand::Dst)}}); + + SmallVector Replacement = + assembleInstructions(ReplacementAsm, Ctx.LS); + if (Replacement.empty()) + return failClosed(Ctx, DI, "M+K split assembly failed"); + + unsigned Extra = Alloc.extraVgprsNeeded(); + if (checkKernelVgprBump(Ctx, KernelName, Extra, PatchRequirement::Required) != + VgprBumpDecision::Apply) + return 0; + if (!emitToTrampoline(Ctx, DI.Offset, DI.Size, Replacement)) + return failClosed(Ctx, DI, "M+K split trampoline emission failed"); + + KernelPatchStats &Stats = Ctx.KernelStats[KernelName]; + Stats.ExtraVgprs = std::max(Stats.ExtraVgprs, Extra); + if (Extra == 0) + Stats.ScratchReused += ScratchCount; + Stats.ScratchAboveKd += Extra; + ScratchPatchInfo Info; + Info.Offset = DI.Offset; + Info.ScratchRegs.resize(Ctx.Config.MaxVgprs); + for (unsigned Reg : SavedARegs) + Info.ScratchRegs.set(Reg); + Ctx.OutScratchPatches.push_back(std::move(Info)); + + log() << "hotswap: wmma_scale16: exact M+K split at offset 0x" + << utohexstr(DI.Offset) << " (D=v" << DBase << ":" << (DBase + 15) + << ", A=v" << ABase << ":" << (ABase + 15) << ", B=v" << BBase << ":" + << (BBase + 7) << ", four saved-A VGPRs, tmp=v" << TmpReg << ", +" + << Extra << " vgpr, 4 WMMAs, " + << Replacement.size() << " bytes)\n"; + return 1; +} + // --------------------------------------------------------------------------- // patchWmmaScale16 -- dispatch // --------------------------------------------------------------------------- @@ -696,9 +1833,9 @@ static uint32_t applyWmmaScale16PatchesImpl(PatchContext &Ctx, size_t Idx) { if (Mnem == "v_wmma_scale16_f32_16x16x128_f8f6f4") return patchWmmaScale16_16x16(Ctx, Idx); + if (Mnem == "v_wmma_scale16_f32_32x16x128_f4") + return patchWmmaScale16_32x16(Ctx, Idx); - // The M=32 FP4 form needs an M-split in addition to the K-split; not yet - // lowered exactly, so fail closed rather than miscompile. if (Mnem.starts_with("v_wmma_scale16_f32_")) return failClosed(Ctx, Ctx.Decoded[Idx], "block-16 scaled variant has no exact lowering yet"); diff --git a/amd/comgr/src/comgr-hotswap-patch-wmma-split.cpp b/amd/comgr/src/comgr-hotswap-patch-wmma-split.cpp index 2429d6d992e5b..bd669e56096e6 100644 --- a/amd/comgr/src/comgr-hotswap-patch-wmma-split.cpp +++ b/amd/comgr/src/comgr-hotswap-patch-wmma-split.cpp @@ -72,6 +72,7 @@ #include "llvm/ADT/SmallVector.h" #include "llvm/ADT/StringExtras.h" #include "llvm/ADT/StringRef.h" +#include "llvm/Support/Endian.h" #include "llvm/Support/FormatVariadic.h" #include "llvm/Support/raw_ostream.h" @@ -652,6 +653,34 @@ VgprMsbState mergeVgprMsbState(VgprMsbState Old, VgprMsbState Incoming) { mergeVgprMsbValue(Old.Src2, Incoming.Src2)}; } +// The B0 f4gemm corpus uses one legacy VOP3 encoding that the A0 gfx1250 +// decoder splits into two unknown dwords. It cannot write MODE: recognize only +// the exact observed opcode/operand spellings, and only when no known entry +// reaches the decoder-created continuation dword. +bool isUndecodedB0Vop3Pair(const PatchContext &Ctx, size_t HeadIndex) { + if (HeadIndex + 1 >= Ctx.Decoded.size()) + return false; + const InternalDecodedInst &Head = Ctx.Decoded[HeadIndex]; + const InternalDecodedInst &Tail = Ctx.Decoded[HeadIndex + 1]; + if (Head.DecodeSucceeded || Head.Size != MinInstSize || + Tail.Offset != Head.Offset + MinInstSize || + Head.Offset > Ctx.TextSize || + 2 * MinInstSize > Ctx.TextSize - Head.Offset) + return false; + + uint32_t Word0 = support::endian::read32le(Ctx.Text + Head.Offset); + uint32_t Word1 = + support::endian::read32le(Ctx.Text + Head.Offset + MinInstSize); + constexpr uint32_t Bit14 = 1u << 14; + if ((Word0 & ~Bit14) != 0xd0310000u || + (Word1 != 0 && Word1 != 0x00100000u)) + return false; + if (Ctx.DirectControlFlow.Targets.contains(Tail.Offset) || + llvm::is_contained(Ctx.DeclaredEntries, Tail.Offset)) + return false; + return true; +} + // Populate Ctx.VgprMsbModeBefore with a per-instruction packed VGPR-MSB mode // via a forward CFG fixed point over each analyzable function. Fail-closed: // only exact, consistent modes are recorded; any conflict, unknown MODE write, @@ -680,7 +709,8 @@ void computeVgprMsbModes(PatchContext &Ctx) { // function's s_set_vgpr_msb prefix. The incoming mode is then unprovable // object-wide, so decline the whole analysis and let each required split fail // closed rather than risk seeding a wrong mode. - if (Ctx.DirectControlFlow.HasUnresolvedTargets) + if (Ctx.DirectControlFlow.HasUnresolvedTargets || + Ctx.DirectControlFlow.HasUnboundedIndirectEntries) return; // Functions entered at a non-start offset by a cross-function branch or call @@ -792,11 +822,21 @@ void computeVgprMsbModes(PatchContext &Ctx) { const size_t Count = static_cast(After - First); DenseMap OffsetToLocalIndex; OffsetToLocalIndex.reserve(Count); + BitVector UndecodedVop3Heads(Count); + BitVector UndecodedVop3Tails(Count); bool Valid = true; for (unsigned I = 0; I != Count; ++I) { OffsetToLocalIndex.try_emplace(First[I].Offset, I); - if (First[I].Mnemonic == "") + if (First[I].Mnemonic != "") + continue; + if (isUndecodedB0Vop3Pair(Ctx, GlobalFirst + I) && I + 1 < Count) { + UndecodedVop3Heads.set(I); + UndecodedVop3Tails.set(I + 1); + OffsetToLocalIndex.erase(First[I + 1].Offset); + ++I; + } else { Valid = false; + } } if (!Valid) continue; @@ -821,6 +861,13 @@ void computeVgprMsbModes(PatchContext &Ctx) { for (unsigned I = 0; I != Count && Valid; ++I) { const InternalDecodedInst &DI = First[I]; SmallVectorImpl &Out = Successors[I]; + if (UndecodedVop3Tails.test(I)) + continue; + if (UndecodedVop3Heads.test(I)) { + if (I + 2 < Count) + Out.push_back(I + 2); + continue; + } if (DI.Inst.getOpcode() == LS.SEndPgmOpcode || DI.Inst.getOpcode() == LS.SEndPgmSavedOpcode || LS.MIA->isReturn(DI.Inst) || isStandardLinkReturn(DI, LS)) @@ -936,7 +983,9 @@ void computeVgprMsbModes(PatchContext &Ctx) { for (size_t Next = 0; Next != Worklist.size(); ++Next) { unsigned I = Worklist[Next]; ModeBefore[GlobalFirst + I] = exactVgprMsbMode(In[I]); - VgprMsbState Out = transferVgprMsbState(In[I], First[I], LS); + VgprMsbState Out = UndecodedVop3Heads.test(I) + ? In[I] + : transferVgprMsbState(In[I], First[I], LS); for (unsigned Succ : Successors[I]) { VgprMsbState Merged = mergeVgprMsbState(In[Succ], Out); if (Merged.Dst != In[Succ].Dst || Merged.Src0 != In[Succ].Src0 || @@ -1136,6 +1185,12 @@ std::optional getActiveVgprMsbMode(PatchContext &Ctx, size_t Idx) { std::optional getLocallyEstablishedVgprMsbMode(PatchContext &Ctx, size_t Idx) { + // An unresolved transfer can enter at any instruction in the object, so an + // apparently adjacent setter is not a dominance proof in that case. + if (Ctx.DirectControlFlow.HasUnresolvedTargets || + Ctx.DirectControlFlow.HasUnboundedIndirectEntries) + return std::nullopt; + while (Idx > 0) { const InternalDecodedInst &Prev = Ctx.Decoded[Idx - 1]; const InternalDecodedInst &Current = Ctx.Decoded[Idx]; @@ -1148,6 +1203,15 @@ std::optional getLocallyEstablishedVgprMsbMode(PatchContext &Ctx, if (Entry == Current.Offset) return std::nullopt; + // The A0 decoder represents the exact B0 legacy-VOP3 pair as two unknown + // dwords. It cannot change MODE, so a straight-line local dominance scan + // may step over the complete pair. isUndecodedB0Vop3Pair also rejects a + // known entry into its continuation dword. + if (Idx >= 2 && isUndecodedB0Vop3Pair(Ctx, Idx - 2)) { + Idx -= 2; + continue; + } + if (std::optional Mode = getExactVgprMsbModeWritten(Prev, Ctx.LS)) return Mode; @@ -1155,14 +1219,22 @@ std::optional getLocallyEstablishedVgprMsbMode(PatchContext &Ctx, Prev.Inst.getOpcode() == Ctx.LS.SSetVgprMsbOpcode || instructionDefinesNamedRegister(Prev, "MODE", Ctx.LS) || (Ctx.LS.MIA && - (Ctx.LS.MIA->isBranch(Prev.Inst) || Ctx.LS.MIA->isCall(Prev.Inst) || - Ctx.LS.MIA->isReturn(Prev.Inst)))) + Ctx.LS.MIA->mayAffectControlFlow(Prev.Inst, *Ctx.LS.MRI))) return std::nullopt; --Idx; } return std::nullopt; } +int16_t transferExactVgprMsbMode(int16_t Incoming, + const InternalDecodedInst &DI, + const LLVMState &LS) { + VgprMsbState State = + Incoming >= 0 ? vgprMsbStateFromMode(static_cast(Incoming)) + : unknownVgprMsbState(); + return exactVgprMsbMode(transferVgprMsbState(State, DI, LS)); +} + unsigned getVgprMsbBank(unsigned Mode, VgprMsbOperand Operand) { return getVgprMsbs(Mode, Operand); } diff --git a/amd/comgr/test-lit/hotswap-wmma-scale16-32x16-refuse.s b/amd/comgr/test-lit/hotswap-wmma-scale16-32x16-refuse.s index 24a1476a8923f..fb096d32cdf15 100644 --- a/amd/comgr/test-lit/hotswap-wmma-scale16-32x16-refuse.s +++ b/amd/comgr/test-lit/hotswap-wmma-scale16-32x16-refuse.s @@ -1,40 +1,102 @@ -// COM: Fail-closed gate for the 32x16 FP4 block-16 scaled WMMA. The M=32 form -// COM: needs an M-split on top of the K-split and has no exact lowering yet, -// COM: so the rewrite refuses it: any unlowerable v_wmma_scale16_f32_* makes the -// COM: rewrite return an error instead of emitting a wrong code object. - -// RUN: %clang -target amdgcn-amd-amdhsa -mcpu=gfx1250 -nostdlib %s -o %t.elf - -// COM: Default mode: presence of the unlowerable form fails the rewrite. -// RUN: hotswap-rewrite %t.elf \ -// RUN: amdgcn-amd-amdhsa--gfx1250 amdgcn-amd-amdhsa--gfx1250 \ -// RUN: --expect-status ERROR \ -// RUN: | %FileCheck --check-prefix=REFUSE %s -// REFUSE: RESULT: ERROR - -// COM: Strict mode: same fail-closed result. -// RUN: hotswap-rewrite %t.elf \ -// RUN: amdgcn-amd-amdhsa--gfx1250 amdgcn-amd-amdhsa--gfx1250 \ -// RUN: --strict-mode --expect-status ERROR \ -// RUN: | %FileCheck --check-prefix=REFUSE %s +// COM: Fail-closed overlap gates for the staged M+K lowering. Each variant is +// COM: assembled into its own object so every rejection is exercised +// COM: independently. + +// RUN: %clang -target amdgcn-amd-amdhsa -mcpu=gfx1250 -nostdlib \ +// RUN: -Wa,-defsym,CASE=1 %s -o %t.da.elf +// RUN: env AMD_COMGR_EMIT_VERBOSE_LOGS=1 hotswap-rewrite %t.da.elf \ +// RUN: amdgcn-amd-amdhsa--gfx1250:gfx1250-b0-specific+ \ +// RUN: amdgcn-amd-amdhsa--gfx1250:gfx1250-b0-specific- \ +// RUN: --expect-status ERROR 2>&1 | %FileCheck --check-prefix=D-A %s +// D-A: destination overlaps matrix A +// D-A: RESULT: ERROR + +// RUN: %clang -target amdgcn-amd-amdhsa -mcpu=gfx1250 -nostdlib \ +// RUN: -Wa,-defsym,CASE=2 %s -o %t.db.elf +// RUN: env AMD_COMGR_EMIT_VERBOSE_LOGS=1 hotswap-rewrite %t.db.elf \ +// RUN: amdgcn-amd-amdhsa--gfx1250:gfx1250-b0-specific+ \ +// RUN: amdgcn-amd-amdhsa--gfx1250:gfx1250-b0-specific- \ +// RUN: --expect-status ERROR 2>&1 | %FileCheck --check-prefix=D-B %s +// D-B: destination overlaps matrix B +// D-B: RESULT: ERROR + +// RUN: %clang -target amdgcn-amd-amdhsa -mcpu=gfx1250 -nostdlib \ +// RUN: -Wa,-defsym,CASE=3 %s -o %t.ab.elf +// RUN: env AMD_COMGR_EMIT_VERBOSE_LOGS=1 hotswap-rewrite %t.ab.elf \ +// RUN: amdgcn-amd-amdhsa--gfx1250:gfx1250-b0-specific+ \ +// RUN: amdgcn-amd-amdhsa--gfx1250:gfx1250-b0-specific- \ +// RUN: --expect-status ERROR 2>&1 | %FileCheck --check-prefix=A-B %s +// A-B: matrix A overlaps matrix B +// A-B: RESULT: ERROR + +// RUN: %clang -target amdgcn-amd-amdhsa -mcpu=gfx1250 -nostdlib \ +// RUN: -Wa,-defsym,CASE=4 %s -o %t.dc.elf +// RUN: env AMD_COMGR_EMIT_VERBOSE_LOGS=1 hotswap-rewrite %t.dc.elf \ +// RUN: amdgcn-amd-amdhsa--gfx1250:gfx1250-b0-specific+ \ +// RUN: amdgcn-amd-amdhsa--gfx1250:gfx1250-b0-specific- \ +// RUN: --expect-status ERROR 2>&1 | %FileCheck --check-prefix=D-C %s +// D-C: partial destination/src2 overlap +// D-C: RESULT: ERROR + +// RUN: %clang -target amdgcn-amd-amdhsa -mcpu=gfx1250 -nostdlib \ +// RUN: -Wa,-defsym,CASE=5 %s -o %t.ac.elf +// RUN: env AMD_COMGR_EMIT_VERBOSE_LOGS=1 hotswap-rewrite %t.ac.elf \ +// RUN: amdgcn-amd-amdhsa--gfx1250:gfx1250-b0-specific+ \ +// RUN: amdgcn-amd-amdhsa--gfx1250:gfx1250-b0-specific- \ +// RUN: --expect-status ERROR 2>&1 | %FileCheck --check-prefix=A-C %s +// A-C: matrix A overlaps src2 +// A-C: RESULT: ERROR + +// RUN: %clang -target amdgcn-amd-amdhsa -mcpu=gfx1250 -nostdlib \ +// RUN: -Wa,-defsym,CASE=6 %s -o %t.scale-matrix.elf +// RUN: env AMD_COMGR_EMIT_VERBOSE_LOGS=1 hotswap-rewrite \ +// RUN: %t.scale-matrix.elf \ +// RUN: amdgcn-amd-amdhsa--gfx1250:gfx1250-b0-specific+ \ +// RUN: amdgcn-amd-amdhsa--gfx1250:gfx1250-b0-specific- \ +// RUN: --expect-status ERROR 2>&1 | %FileCheck --check-prefix=SCALE %s +// SCALE: scale pair overlaps a staged matrix operand +// SCALE: RESULT: ERROR + +// RUN: %clang -target amdgcn-amd-amdhsa -mcpu=gfx1250 -nostdlib \ +// RUN: -Wa,-defsym,CASE=7 %s -o %t.scale-scale.elf +// RUN: env AMD_COMGR_EMIT_VERBOSE_LOGS=1 hotswap-rewrite \ +// RUN: %t.scale-scale.elf \ +// RUN: amdgcn-amd-amdhsa--gfx1250:gfx1250-b0-specific+ \ +// RUN: amdgcn-amd-amdhsa--gfx1250:gfx1250-b0-specific- \ +// RUN: --expect-status ERROR 2>&1 | %FileCheck --check-prefix=SCALE %s .amdgcn_target "amdgcn-amd-amdhsa--gfx1250" .text -// --- Scale16 32x16 FP4 (block-16) -> refuse (fail closed) --- -.globl test_wmma_scale16_32x16 +.globl test_wmma_scale16_32x16_refuse .p2align 8 -.type test_wmma_scale16_32x16,@function -test_wmma_scale16_32x16: - v_wmma_scale16_f32_32x16x128_f4 v[0:15], v[16:31], v[32:39], v[0:15], v[40:41], v[42:43] +.type test_wmma_scale16_32x16_refuse,@function +test_wmma_scale16_32x16_refuse: +.if CASE == 1 + v_wmma_scale16_f32_32x16x128_f4 v[0:15], v[8:23], v[32:39], v[0:15], v[40:41], v[42:43] +.elseif CASE == 2 + v_wmma_scale16_f32_32x16x128_f4 v[0:15], v[16:31], v[8:15], v[0:15], v[40:41], v[42:43] +.elseif CASE == 3 + v_wmma_scale16_f32_32x16x128_f4 v[0:15], v[16:31], v[24:31], v[0:15], v[40:41], v[42:43] +.elseif CASE == 4 + v_wmma_scale16_f32_32x16x128_f4 v[0:15], v[24:39], v[40:47], v[8:23], v[48:49], v[50:51] +.elseif CASE == 5 + v_wmma_scale16_f32_32x16x128_f4 v[0:15], v[16:31], v[32:39], v[16:31], v[40:41], v[42:43] +.elseif CASE == 6 + v_wmma_scale16_f32_32x16x128_f4 v[0:15], v[16:31], v[32:39], v[0:15], v[16:17], v[42:43] +.elseif CASE == 7 + v_wmma_scale16_f32_32x16x128_f4 v[0:15], v[16:31], v[32:39], v[0:15], v[40:41], v[40:41] +.else + .error "unknown CASE" +.endif s_endpgm -.Ltest_wmma_scale16_32x16_end: -.size test_wmma_scale16_32x16, .Ltest_wmma_scale16_32x16_end-test_wmma_scale16_32x16 +.Ltest_wmma_scale16_32x16_refuse_end: +.size test_wmma_scale16_32x16_refuse, .Ltest_wmma_scale16_32x16_refuse_end-test_wmma_scale16_32x16_refuse .rodata .p2align 8 -.amdhsa_kernel test_wmma_scale16_32x16 - .amdhsa_next_free_vgpr 44 +.amdhsa_kernel test_wmma_scale16_32x16_refuse + .amdhsa_next_free_vgpr 52 .amdhsa_next_free_sgpr 2 .amdhsa_wavefront_size32 1 .end_amdhsa_kernel @@ -44,10 +106,10 @@ test_wmma_scale16_32x16: - 3 - 0 amdhsa.kernels: - - .name: test_wmma_scale16_32x16 - .symbol: test_wmma_scale16_32x16.kd + - .name: test_wmma_scale16_32x16_refuse + .symbol: test_wmma_scale16_32x16_refuse.kd .sgpr_count: 2 - .vgpr_count: 44 + .vgpr_count: 52 .kernarg_segment_size: 0 .group_segment_fixed_size: 0 .private_segment_fixed_size: 0 diff --git a/amd/comgr/test-lit/hotswap-wmma-scale16-32x16.s b/amd/comgr/test-lit/hotswap-wmma-scale16-32x16.s new file mode 100644 index 0000000000000..480096da18df8 --- /dev/null +++ b/amd/comgr/test-lit/hotswap-wmma-scale16-32x16.s @@ -0,0 +1,146 @@ +// COM: Exact decomposition of the corpus's M=32 FP4 block-16 scaled WMMA. +// COM: The schedule is low-M/low-K, high-M/low-K, low-M/high-K, +// COM: high-M/high-K. The low passes read the original C halves and even scale +// COM: bytes; the high passes accumulate from D and use odd scale bytes. Only +// COM: the four A registers overwritten by each FP4 mask are saved, and the +// COM: first save slot is reused as the reversible scale-pair permutation +// COM: temporary. + +// RUN: %clang -target amdgcn-amd-amdhsa -mcpu=gfx1250 -nostdlib %s -o %t.elf +// RUN: env AMD_COMGR_EMIT_VERBOSE_LOGS=1 hotswap-rewrite %t.elf \ +// RUN: amdgcn-amd-amdhsa--gfx1250:gfx1250-b0-specific+ \ +// RUN: amdgcn-amd-amdhsa--gfx1250:gfx1250-b0-specific- \ +// RUN: --output %t.out.elf 2>&1 | %FileCheck --check-prefix=API %s +// API: physical forward-dead proof at offset +// API-SAME: found 4 VGPRs +// API: wmma_scale16: exact M+K split +// API-SAME: four saved-A VGPRs, tmp=v55, +0 vgpr, 4 WMMAs +// API-NOT: error: +// API: liveness: kernel test_wmma_scale16_32x16: +// API-SAME: scratch_reused=4, scratch_above_kd=0 +// API: RESULT: SUCCESS + +// RUN: %llvm-objdump -d %t.out.elf | %FileCheck --check-prefix=DISASM %s +// DISASM-LABEL: : +// DISASM-NOT: v_wmma_scale16 +// DISASM: s_branch +// DISASM: s_endpgm +// +// COM: Forward byte split and its exact inverse. These selectors implement +// COM: [0,2,4,6] / [1,3,5,7], then restore the two original dwords exactly. +// DISASM: v_mov_b32_e32 v[[TMP:[0-9]+]], v40 +// DISASM-NEXT: v_perm_b32 v40, v[[TMP]], v41, 0x6040200 +// DISASM-NEXT: v_perm_b32 v41, v[[TMP]], v41, 0x7050301 +// DISASM: v_mov_b32_e32 v[[TMP]], v42 +// DISASM-NEXT: v_perm_b32 v42, v[[TMP]], v43, 0x6040200 +// DISASM-NEXT: v_perm_b32 v43, v[[TMP]], v43, 0x7050301 +// +// COM: D==C, low-low-high-high order. Reuse hints are stripped, and the +// COM: original C modifier appears only on the two low-K passes. +// DISASM-NOT: matrix_a_reuse +// DISASM-NOT: matrix_b_reuse +// DISASM: v_wmma_scale_f32_16x16x128_f8f6f4 v[0:7], v[16:23], v[32:39], v[0:7], v40, v42{{.*}}matrix_a_fmt:MATRIX_FMT_FP4{{.*}}matrix_b_fmt:MATRIX_FMT_FP4{{.*}}neg_lo:[0,0,1] +// DISASM: v_nop +// DISASM: v_wmma_scale_f32_16x16x128_f8f6f4 v[8:15], v[24:31], v[32:39], v[8:15], v40, v42{{.*}}matrix_a_fmt:MATRIX_FMT_FP4{{.*}}matrix_b_fmt:MATRIX_FMT_FP4{{.*}}neg_lo:[0,0,1] +// DISASM: v_nop +// DISASM: v_wmma_scale_f32_16x16x128_f8f6f4 v[0:7], v[16:23], v[32:39], v[0:7], v41, v43{{.*}}matrix_a_fmt:MATRIX_FMT_FP4{{.*}}matrix_b_fmt:MATRIX_FMT_FP4 +// DISASM-NOT: neg_lo +// DISASM: v_nop +// DISASM: v_wmma_scale_f32_16x16x128_f8f6f4 v[8:15], v[24:31], v[32:39], v[8:15], v41, v43{{.*}}matrix_a_fmt:MATRIX_FMT_FP4{{.*}}matrix_b_fmt:MATRIX_FMT_FP4 +// DISASM-NOT: neg_lo +// DISASM: v_nop +// DISASM: v_mov_b32_e32 v[[TMP]], v40 +// DISASM-NEXT: v_perm_b32 v40, v[[TMP]], v41, 0x5010400 +// DISASM-NEXT: v_perm_b32 v41, v[[TMP]], v41, 0x7030602 +// DISASM: v_mov_b32_e32 v[[TMP]], v42 +// DISASM-NEXT: v_perm_b32 v42, v[[TMP]], v43, 0x5010400 +// DISASM-NEXT: v_perm_b32 v43, v[[TMP]], v43, 0x7030602 +// +// COM: The second lowering permits matrix B to overlap C. Its first low pass +// COM: therefore has the same v[32:39] range in src1 and src2. +// DISASM: v_wmma_scale_f32_16x16x128_f8f6f4 v[0:7], v[16:23], v[32:39], v[32:39], v48, v50 +// DISASM-NOT: v_wmma_scale16 + +// RUN: hotswap-rewrite %t.out.elf \ +// RUN: amdgcn-amd-amdhsa--gfx1250:gfx1250-b0-specific- \ +// RUN: amdgcn-amd-amdhsa--gfx1250:gfx1250-b0-specific- \ +// RUN: --check-idempotent | %FileCheck --check-prefix=IDEM %s +// IDEM: IDEMPOTENT: YES + +.amdgcn_target "amdgcn-amd-amdhsa--gfx1250" +.text + +// Reduced opaque sequence from the f4gemm corpus. The A0 decoder splits this +// legacy B0 VOP3 into two unknown dwords. The sequence before the WMMA tests +// local MODE recovery; the sequence after it tests forward physical liveness. +.macro opaque_b0_vop3 + .long 0xd0310000 + .long 0x00100000 + v_cmp_ge_u16_e32 vcc_lo, s32, v18.l +.endm + +.globl test_wmma_scale16_32x16 +.p2align 8 +.type test_wmma_scale16_32x16,@function +test_wmma_scale16_32x16: + s_set_vgpr_msb 0 + opaque_b0_vop3 + v_wmma_scale16_f32_32x16x128_f4 v[0:15], v[16:31], v[32:39], v[0:15], v[40:41], v[42:43] matrix_a_scale:MATRIX_SCALE_ROW1 matrix_b_scale:MATRIX_SCALE_ROW1 matrix_a_scale_fmt:MATRIX_SCALE_FMT_E4M3 matrix_b_scale_fmt:MATRIX_SCALE_FMT_E4M3 matrix_a_reuse matrix_b_reuse neg_lo:[0,0,1] + opaque_b0_vop3 + v_mov_b32 v52, 0 + v_mov_b32 v53, 0 + v_mov_b32 v54, 0 + v_mov_b32 v55, 0 + s_branch test_wmma_scale16_32x16_bc_overlap +.Ltest_wmma_scale16_32x16_end: +.size test_wmma_scale16_32x16, .Ltest_wmma_scale16_32x16_end-test_wmma_scale16_32x16 + +.globl test_wmma_scale16_32x16_bc_overlap +.p2align 8 +.type test_wmma_scale16_32x16_bc_overlap,@function +test_wmma_scale16_32x16_bc_overlap: + v_wmma_scale16_f32_32x16x128_f4 v[0:15], v[16:31], v[32:39], v[32:47], v[48:49], v[50:51] matrix_a_scale_fmt:MATRIX_SCALE_FMT_E4M3 matrix_b_scale_fmt:MATRIX_SCALE_FMT_E4M3 + s_endpgm +.Ltest_wmma_scale16_32x16_bc_overlap_end: +.size test_wmma_scale16_32x16_bc_overlap, .Ltest_wmma_scale16_32x16_bc_overlap_end-test_wmma_scale16_32x16_bc_overlap + +.rodata +.p2align 8 +.amdhsa_kernel test_wmma_scale16_32x16 + .amdhsa_next_free_vgpr 1024 + .amdhsa_next_free_sgpr 34 + .amdhsa_wavefront_size32 1 +.end_amdhsa_kernel + +.amdhsa_kernel test_wmma_scale16_32x16_bc_overlap + .amdhsa_next_free_vgpr 52 + .amdhsa_next_free_sgpr 2 + .amdhsa_wavefront_size32 1 +.end_amdhsa_kernel + +.amdgpu_metadata + amdhsa.version: + - 3 + - 0 + amdhsa.kernels: + - .name: test_wmma_scale16_32x16 + .symbol: test_wmma_scale16_32x16.kd + .sgpr_count: 34 + .vgpr_count: 1024 + .kernarg_segment_size: 0 + .group_segment_fixed_size: 0 + .private_segment_fixed_size: 0 + .kernarg_segment_align: 8 + .wavefront_size: 32 + .max_flat_workgroup_size: 256 + - .name: test_wmma_scale16_32x16_bc_overlap + .symbol: test_wmma_scale16_32x16_bc_overlap.kd + .sgpr_count: 2 + .vgpr_count: 52 + .kernarg_segment_size: 0 + .group_segment_fixed_size: 0 + .private_segment_fixed_size: 0 + .kernarg_segment_align: 8 + .wavefront_size: 32 + .max_flat_workgroup_size: 256 +.end_amdgpu_metadata diff --git a/amd/comgr/test-unit/HotswapMCTest.cpp b/amd/comgr/test-unit/HotswapMCTest.cpp index ffbb482f912c8..7df6a9c22f505 100644 --- a/amd/comgr/test-unit/HotswapMCTest.cpp +++ b/amd/comgr/test-unit/HotswapMCTest.cpp @@ -2640,6 +2640,44 @@ TEST(CollectDirectBranchTargets, RejectsUndecodedMaterializationSlot) { EXPECT_TRUE(Info->HasUnresolvedTargets); } +TEST(CollectDirectBranchTargets, UndecodedVectorAluDoesNotCreateIndirectEntry) { + LLVMState S = initLLVM(makeGfx1250Ident()); + ASSERT_TRUE(S.Valid); + + // Reduced from a B0 f4gemm code object. The A0 MC decoder does not know this + // legacy-VOP3-major 0x34 vector instruction, but its encoding class cannot + // affect scalar control flow or MODE. + const uint8_t Bytes[] = {0x00, 0x00, 0x31, 0xd0, 0x00, 0x00, 0x10, 0x00}; + std::vector Decoded; + ASSERT_TRUE(decodeTextSection(Bytes, sizeof(Bytes), S, Decoded)); + ASSERT_FALSE(Decoded.empty()); + ASSERT_FALSE(Decoded.front().DecodeSucceeded); + + std::optional Info = collectDirectBranchTargets( + Decoded, S, /*TextAddr=*/0, sizeof(Bytes), /*DeclaredEntries=*/{0}, + /*FunctionRanges=*/{}, /*ExternalEntries=*/{}, Bytes); + ASSERT_TRUE(Info); + EXPECT_FALSE(Info->HasUnresolvedTargets); + EXPECT_FALSE(Info->HasUnboundedIndirectEntries); +} + +TEST(CollectDirectBranchTargets, UndecodedScalarClassRemainsUnbounded) { + LLVMState S = initLLVM(makeGfx1250Ident()); + ASSERT_TRUE(S.Valid); + + const uint8_t Bytes[] = {0xff, 0xff, 0xff, 0xff}; + std::vector Decoded; + ASSERT_TRUE(decodeTextSection(Bytes, sizeof(Bytes), S, Decoded)); + ASSERT_EQ(Decoded.size(), 1u); + ASSERT_FALSE(Decoded.front().DecodeSucceeded); + + std::optional Info = collectDirectBranchTargets( + Decoded, S, /*TextAddr=*/0, sizeof(Bytes), /*DeclaredEntries=*/{0}, + /*FunctionRanges=*/{}, /*ExternalEntries=*/{}, Bytes); + ASSERT_TRUE(Info); + EXPECT_TRUE(Info->HasUnboundedIndirectEntries); +} + TEST(CollectDirectBranchTargets, RejectsUnboundedIndirectEntry) { LLVMState S = initLLVM(makeGfx1250Ident()); ASSERT_TRUE(S.Valid); @@ -4350,6 +4388,225 @@ TEST(SafeSgprScratchBlock, OwnerlessCommitFailureIsAtomic) { EXPECT_EQ(Ctx.SgprDescriptorChargePasses, 0u); } +TEST(ForwardDeadVgprs, OpaqueBeforeKillRejects) { + constexpr unsigned MaxVgprs = 8; + constexpr unsigned Candidate = 3; + std::vector Nodes; + Nodes.emplace_back(MaxVgprs); + Nodes.emplace_back(MaxVgprs); + Nodes[0].Opaque = true; + Nodes[1].FullDefs.set(Candidate); + Nodes[1].SafeTerminal = true; + + std::optional Safe = + computeForwardDeadVgprs(Nodes, /*EntryNode=*/0, MaxVgprs); + ASSERT_TRUE(Safe); + EXPECT_FALSE(Safe->test(Candidate)); +} + +TEST(ForwardDeadVgprs, KillBeforeOpaqueAccepts) { + constexpr unsigned MaxVgprs = 8; + constexpr unsigned Candidate = 3; + std::vector Nodes; + Nodes.emplace_back(MaxVgprs); + Nodes.emplace_back(MaxVgprs); + Nodes[0].FullDefs.set(Candidate); + Nodes[0].Successors.push_back(1); + Nodes[1].Opaque = true; + + std::optional Safe = + computeForwardDeadVgprs(Nodes, /*EntryNode=*/0, MaxVgprs); + ASSERT_TRUE(Safe); + EXPECT_TRUE(Safe->test(Candidate)); +} + +TEST(ForwardDeadVgprs, ExternalExitBeforeKillRejects) { + constexpr unsigned MaxVgprs = 8; + constexpr unsigned Candidate = 3; + std::vector Nodes; + Nodes.emplace_back(MaxVgprs); + Nodes.emplace_back(MaxVgprs); + Nodes[0].HasUnsafeExit = true; + Nodes[0].Successors.push_back(1); + Nodes[1].FullDefs.set(Candidate); + Nodes[1].SafeTerminal = true; + + std::optional Safe = + computeForwardDeadVgprs(Nodes, /*EntryNode=*/0, MaxVgprs); + ASSERT_TRUE(Safe); + EXPECT_FALSE(Safe->test(Candidate)); +} + +TEST(ForwardDeadVgprs, KillBeforeExternalExitAccepts) { + constexpr unsigned MaxVgprs = 8; + constexpr unsigned Candidate = 3; + std::vector Nodes; + Nodes.emplace_back(MaxVgprs); + Nodes.emplace_back(MaxVgprs); + Nodes[0].FullDefs.set(Candidate); + Nodes[0].Successors.push_back(1); + Nodes[1].HasUnsafeExit = true; + Nodes[1].SafeTerminal = true; + + std::optional Safe = + computeForwardDeadVgprs(Nodes, /*EntryNode=*/0, MaxVgprs); + ASSERT_TRUE(Safe); + EXPECT_TRUE(Safe->test(Candidate)); +} + +TEST(ForwardDeadVgprs, OneUseBeforeKillBranchRejects) { + constexpr unsigned MaxVgprs = 8; + constexpr unsigned Candidate = 3; + std::vector Nodes; + for (unsigned I = 0; I != 3; ++I) + Nodes.emplace_back(MaxVgprs); + Nodes[0].Successors.push_back(1); + Nodes[0].Successors.push_back(2); + Nodes[1].Uses.set(Candidate); + Nodes[1].SafeTerminal = true; + Nodes[2].FullDefs.set(Candidate); + Nodes[2].SafeTerminal = true; + + std::optional Safe = + computeForwardDeadVgprs(Nodes, /*EntryNode=*/0, MaxVgprs); + ASSERT_TRUE(Safe); + EXPECT_FALSE(Safe->test(Candidate)); +} + +TEST(ForwardDeadVgprs, AllBranchesKillAccepts) { + constexpr unsigned MaxVgprs = 8; + constexpr unsigned Candidate = 3; + std::vector Nodes; + for (unsigned I = 0; I != 3; ++I) + Nodes.emplace_back(MaxVgprs); + Nodes[0].Successors.push_back(1); + Nodes[0].Successors.push_back(2); + Nodes[1].FullDefs.set(Candidate); + Nodes[1].SafeTerminal = true; + Nodes[2].FullDefs.set(Candidate); + Nodes[2].SafeTerminal = true; + + std::optional Safe = + computeForwardDeadVgprs(Nodes, /*EntryNode=*/0, MaxVgprs); + ASSERT_TRUE(Safe); + EXPECT_TRUE(Safe->test(Candidate)); +} + +TEST(ForwardDeadVgprs, LoopPathWithoutKillRejects) { + constexpr unsigned MaxVgprs = 8; + constexpr unsigned Candidate = 3; + std::vector Nodes; + Nodes.emplace_back(MaxVgprs); + Nodes.emplace_back(MaxVgprs); + Nodes[0].Successors.push_back(1); + Nodes[1].Successors.push_back(0); + + std::optional Safe = + computeForwardDeadVgprs(Nodes, /*EntryNode=*/0, MaxVgprs); + ASSERT_TRUE(Safe); + EXPECT_FALSE(Safe->test(Candidate)); +} + +TEST(ForwardDeadVgprs, LoopWithFullKillAccepts) { + constexpr unsigned MaxVgprs = 8; + constexpr unsigned Candidate = 3; + std::vector Nodes; + Nodes.emplace_back(MaxVgprs); + Nodes.emplace_back(MaxVgprs); + Nodes[0].FullDefs.set(Candidate); + Nodes[0].Successors.push_back(1); + Nodes[1].Successors.push_back(0); + + std::optional Safe = + computeForwardDeadVgprs(Nodes, /*EntryNode=*/0, MaxVgprs); + ASSERT_TRUE(Safe); + EXPECT_TRUE(Safe->test(Candidate)); +} + +TEST(ForwardDeadVgprs, BackedgeUseBeforePatchedSiteRejects) { + constexpr unsigned MaxVgprs = 8; + constexpr unsigned Candidate = 3; + std::vector Nodes; + for (unsigned I = 0; I != 3; ++I) + Nodes.emplace_back(MaxVgprs); + Nodes[0].Successors.push_back(1); + Nodes[1].Uses.set(Candidate); + Nodes[1].Successors.push_back(2); + Nodes[2].SafeTerminal = true; + + std::optional Safe = + computeForwardDeadVgprs(Nodes, /*EntryNode=*/0, MaxVgprs); + ASSERT_TRUE(Safe); + EXPECT_FALSE(Safe->test(Candidate)); +} + +TEST(ForwardDeadVgprs, BackedgeWithoutUseToPatchedSiteAccepts) { + constexpr unsigned MaxVgprs = 8; + constexpr unsigned Candidate = 3; + std::vector Nodes; + for (unsigned I = 0; I != 3; ++I) + Nodes.emplace_back(MaxVgprs); + Nodes[0].Successors.push_back(1); + Nodes[1].Successors.push_back(2); + Nodes[2].SafeTerminal = true; + + std::optional Safe = + computeForwardDeadVgprs(Nodes, /*EntryNode=*/0, MaxVgprs); + ASSERT_TRUE(Safe); + EXPECT_TRUE(Safe->test(Candidate)); +} + +TEST(ForwardDeadVgprs, PartialDefinitionIsUseNotKill) { + constexpr unsigned MaxVgprs = 8; + constexpr unsigned Candidate = 3; + std::vector Nodes; + Nodes.emplace_back(MaxVgprs); + // Partial/tied definitions consume the incoming full dword and therefore + // belong in Uses, never FullDefs. + Nodes[0].Uses.set(Candidate); + Nodes[0].SafeTerminal = true; + + std::optional Safe = + computeForwardDeadVgprs(Nodes, /*EntryNode=*/0, MaxVgprs); + ASSERT_TRUE(Safe); + EXPECT_FALSE(Safe->test(Candidate)); +} + +TEST(WmmaScale16, PhysicalVgprRangeMustFitOneBank) { + EXPECT_TRUE(physicalVgprRangeFitsOneBank(0, 16, 1024)); + EXPECT_TRUE(physicalVgprRangeFitsOneBank(248, 8, 1024)); + EXPECT_TRUE(physicalVgprRangeFitsOneBank(1016, 8, 1024)); + + EXPECT_FALSE(physicalVgprRangeFitsOneBank(0, 0, 1024)); + EXPECT_FALSE(physicalVgprRangeFitsOneBank(249, 8, 1024)); + EXPECT_FALSE(physicalVgprRangeFitsOneBank(255, 2, 1024)); + EXPECT_FALSE(physicalVgprRangeFitsOneBank(1017, 8, 1024)); + EXPECT_FALSE(physicalVgprRangeFitsOneBank(1024, 1, 1024)); +} + +TEST(WmmaScale16, UnrecognizedVectorRegisterCannotDisappearFromProof) { + LLVMState S = initLLVM(makeGfx1250Ident()); + ASSERT_TRUE(S.Valid); + ASSERT_NE(S.MRI, nullptr); + + auto FindRegister = [&](llvm::StringRef Name) { + for (unsigned Reg = 1; Reg != S.MRI->getNumRegs(); ++Reg) + if (Name == S.MRI->getName(Reg)) + return llvm::MCRegister(Reg); + return llvm::MCRegister(); + }; + + // AGPRs are vector registers but are deliberately not representable as an + // encoded v0..v255 range. Such an operand must invalidate the physical-VGPR + // proof; it cannot be ignored like a scalar register. + llvm::MCRegister Agpr0 = FindRegister("AGPR0"); + llvm::MCRegister Sgpr0 = FindRegister("SGPR0"); + ASSERT_TRUE(Agpr0); + ASSERT_TRUE(Sgpr0); + EXPECT_TRUE(isVectorRegisterOrAlias(Agpr0, *S.MRI)); + EXPECT_FALSE(isVectorRegisterOrAlias(Sgpr0, *S.MRI)); +} + TEST(FindNearestSled, RejectsOverflowingHeadroom) { std::vector Sleds = {{0, 64, 60, 0, 64}, {100, 128, 100, 100, 128}}; EXPECT_EQ(findNearestSled(Sleds, 0, std::numeric_limits::max()),