diff --git a/amd/comgr/src/comgr-hotswap-b0a0.cpp b/amd/comgr/src/comgr-hotswap-b0a0.cpp index 3c9b5c326230c..bdedc40fc941a 100755 --- a/amd/comgr/src/comgr-hotswap-b0a0.cpp +++ b/amd/comgr/src/comgr-hotswap-b0a0.cpp @@ -30,10 +30,12 @@ #include "comgr-env.h" #include "comgr-hotswap-internal.h" +#include "llvm/ADT/DenseMap.h" #include "llvm/ADT/DenseSet.h" #include "llvm/ADT/StringExtras.h" #include "llvm/ADT/Twine.h" #include "llvm/Support/Compiler.h" +#include "llvm/Support/Endian.h" #include "llvm/Support/MathExtras.h" #include @@ -42,6 +44,7 @@ #include #include #include +#include using namespace llvm; @@ -406,19 +409,19 @@ truncateNopSledsAtDirectTargets(std::vector &Sleds, return true; } -std::optional> encodeSetPCLongBranch(const LLVMState &LS, - uint64_t FromOffset, - uint64_t TargetOffset, - unsigned SgprBase) { - if ((SgprBase & 1u) != 0) { +std::optional> +encodeSetPCLongBranch(const LLVMState &LS, uint64_t FromOffset, + uint64_t TargetOffset, unsigned SgprBase, bool UseVcc) { + if (!UseVcc && (SgprBase & 1u) != 0) { log() << "hotswap: error: set-PC long branch requires an aligned " "SGPR pair, got s" << SgprBase << "\n"; return std::nullopt; } - const std::string Pair = "s[" + std::to_string(SgprBase) + ":" + - std::to_string(SgprBase + 1) + "]"; + const std::string Pair = UseVcc ? "vcc" + : "s[" + std::to_string(SgprBase) + ":" + + std::to_string(SgprBase + 1) + "]"; SmallVector GetPc = assembleSingleInst("s_get_pc_i64 " + Pair, LS); if (GetPc.empty()) return std::nullopt; @@ -489,10 +492,68 @@ getSetPcLongBranchLayoutSize(uint64_t FromOffset, uint64_t TargetOffset) { return SetPcReturnReserveBytes; } +static std::optional> +encodeSetPcGateway(const LLVMState &LS, uint64_t FromOffset, + uint64_t TargetOffset, unsigned SgprBase, bool UseVcc, + bool PreserveVcc) { + SmallVector Bytes; + uint64_t SetPcOffset = FromOffset; + if (PreserveVcc) { + if (!UseVcc) { + log() << "hotswap: error: VCC-preserving gateway does not use VCC\n"; + return std::nullopt; + } + Bytes = assembleSingleInst( + "s_mov_b32 s" + std::to_string(SgprBase) + ", vcc_lo", LS); + if (Bytes.size() != VccSaveRestoreBytes) + return std::nullopt; + std::optional Offset = checkedAddUint64( + FromOffset, Bytes.size(), "VCC-preserving set-PC gateway offset"); + if (!Offset) + return std::nullopt; + SetPcOffset = *Offset; + } + + std::optional> SetPc = + encodeSetPCLongBranch(LS, SetPcOffset, TargetOffset, SgprBase, UseVcc); + if (!SetPc) + return std::nullopt; + Bytes.append(SetPc->begin(), SetPc->end()); + return Bytes; +} + +static std::optional +getSetPcGatewayLayoutSize(uint64_t FromOffset, uint64_t TargetOffset, + unsigned SgprBase, bool UseVcc, + bool PreserveVcc) { + if (PreserveVcc && !UseVcc) + return std::nullopt; + if (!UseVcc && (SgprBase & 1u) != 0) + return std::nullopt; + + uint64_t SetPcOffset = FromOffset; + uint32_t PrefixBytes = 0; + if (PreserveVcc) { + std::optional Offset = checkedAddUint64( + FromOffset, VccSaveRestoreBytes, + "VCC-preserving set-PC gateway layout offset"); + if (!Offset) + return std::nullopt; + SetPcOffset = *Offset; + PrefixBytes = VccSaveRestoreBytes; + } + + std::optional SetPcBytes = + getSetPcLongBranchLayoutSize(SetPcOffset, TargetOffset); + if (!SetPcBytes) + return std::nullopt; + return PrefixBytes + *SetPcBytes; +} + Expected> findNearestSetPcGateway(std::vector &Gateways, const LLVMState &LS, uint64_t FromOffset, uint64_t TargetOffset, - unsigned SgprBase) { + unsigned SgprBase, bool UseVcc, bool PreserveVcc) { NopSled *Best = nullptr; uint32_t BestLayoutSize = 0; uint64_t BestUsableEnd = 0; @@ -508,10 +569,10 @@ findNearestSetPcGateway(std::vector &Gateways, const LLVMState &LS, if (Distance >= MaxSledDistance || Distance >= BestDistance || LS.encodeSBranch(FromOffset, Sled.WritePos).empty()) continue; - std::optional LayoutSize = - getSetPcLongBranchLayoutSize(Sled.WritePos, TargetOffset); - if ((SgprBase & 1u) != 0 || !LayoutSize) + getSetPcGatewayLayoutSize(Sled.WritePos, TargetOffset, SgprBase, + UseVcc, PreserveVcc); + if (!LayoutSize) return createStringError( Twine("failed to encode set-PC gateway at candidate offset 0x") + utohexstr(Sled.WritePos)); @@ -526,7 +587,8 @@ findNearestSetPcGateway(std::vector &Gateways, const LLVMState &LS, if (!Best) return std::nullopt; std::optional> BestBytes = - encodeSetPCLongBranch(LS, Best->WritePos, TargetOffset, SgprBase); + encodeSetPcGateway(LS, Best->WritePos, TargetOffset, SgprBase, UseVcc, + PreserveVcc); if (!BestBytes) return createStringError( Twine("failed to encode set-PC gateway at candidate offset 0x") + @@ -640,7 +702,8 @@ summarizeSafeSgprUsage(PatchContext &Ctx, std::optional findSafeSgprScratchBlock(PatchContext &Ctx, uint64_t TextOffset, unsigned Count, - unsigned Alignment, StringRef Context) { + unsigned Alignment, StringRef Context, + bool ReportNoSpace) { if (Count == 0 || Alignment == 0 || (Alignment & (Alignment - 1)) != 0) { log() << "hotswap: error: " << Context << ": invalid global SGPR block request (count=" << Count @@ -737,8 +800,10 @@ findSafeSgprScratchBlock(PatchContext &Ctx, uint64_t TextOffset, unsigned Count, } unsigned Base = (HighWatermark + Alignment - 1) & ~(Alignment - 1); if (Base > Ctx.Config.MaxSgprs || Count > Ctx.Config.MaxSgprs - Base) { - log() << "hotswap: error: " << Context << ": no aligned block of " << Count - << " safe SGPRs fits below s" << Ctx.Config.MaxSgprs << "\n"; + if (ReportNoSpace) + log() << "hotswap: error: " << Context << ": no aligned block of " + << Count << " safe SGPRs fits below s" << Ctx.Config.MaxSgprs + << "\n"; return std::nullopt; } return SafeSgprScratchBlock{Base, Count}; @@ -791,15 +856,449 @@ bool commitSafeSgprScratchBlock(PatchContext &Ctx, uint64_t TextOffset, return true; } -static std::optional -reserveSafeFarReturn(PatchContext &Ctx, uint64_t InstOffset) { - std::optional Scratch = findSafeSgprScratchBlock( - Ctx, InstOffset, /*Count=*/2, /*Alignment=*/2, "safe far return"); - if (!Scratch) +bool instructionReadsRegister(const InternalDecodedInst &DI, + const LLVMState &LS, MCRegister Register) { + const MCInstrDesc &Desc = LS.MCII->get(DI.Inst.getOpcode()); + unsigned DefCount = std::min(Desc.getNumDefs(), DI.Inst.getNumOperands()); + // A tied use makes its corresponding explicit def a read/modify/write + // operand. Some MCInst producers materialize a duplicate use operand while + // others only retain the destination operand, so consult the descriptor + // instead of relying on the decoded operand list to contain that duplicate. + for (unsigned Def = 0; Def != DefCount; ++Def) { + bool HasTiedUse = false; + for (unsigned Use = Desc.getNumDefs(); Use != Desc.getNumOperands(); + ++Use) { + if (Desc.getOperandConstraint(Use, MCOI::TIED_TO) == + static_cast(Def)) { + HasTiedUse = true; + break; + } + } + if (!HasTiedUse) + continue; + const MCOperand &Operand = DI.Inst.getOperand(Def); + if (Operand.isReg() && Operand.getReg() && + LS.MRI->regsOverlap(MCRegister(Operand.getReg()), Register)) + return true; + } + for (unsigned I = DefCount; I != DI.Inst.getNumOperands(); ++I) { + const MCOperand &Operand = DI.Inst.getOperand(I); + if (Operand.isReg() && Operand.getReg() && + LS.MRI->regsOverlap(MCRegister(Operand.getReg()), Register)) + return true; + } + for (MCPhysReg ImplicitUse : Desc.implicit_uses()) + if (LS.MRI->regsOverlap(MCRegister(ImplicitUse), Register)) + return true; + return false; +} + +static bool instructionWritesRegister(const InternalDecodedInst &DI, + const LLVMState &LS, + MCRegister Register) { + const MCInstrDesc &Desc = LS.MCII->get(DI.Inst.getOpcode()); + unsigned DefCount = std::min(Desc.getNumDefs(), DI.Inst.getNumOperands()); + for (unsigned I = 0; I != DefCount; ++I) { + const MCOperand &Operand = DI.Inst.getOperand(I); + if (Operand.isReg() && Operand.getReg() && + LS.MRI->regsOverlap(MCRegister(Operand.getReg()), Register)) + return true; + } + if (Desc.variadicOpsAreDefs()) { + unsigned VariadicBegin = + std::min(Desc.getNumOperands(), DI.Inst.getNumOperands()); + for (unsigned I = VariadicBegin; I != DI.Inst.getNumOperands(); ++I) { + const MCOperand &Operand = DI.Inst.getOperand(I); + if (Operand.isReg() && Operand.getReg() && + LS.MRI->regsOverlap(MCRegister(Operand.getReg()), Register)) + return true; + } + } + for (MCPhysReg ImplicitDef : Desc.implicit_defs()) + if (LS.MRI->regsOverlap(MCRegister(ImplicitDef), Register)) + return true; + return false; +} + +static bool replacementNeedsIncomingRegister(ArrayRef Replacement, + const LLVMState &LS, + MCRegister Register) { + std::vector Decoded; + if (!decodeTextSection(Replacement.data(), Replacement.size(), LS, Decoded)) + return true; + + for (const InternalDecodedInst &DI : Decoded) { + if (!DI.DecodeSucceeded || !LS.MIA || + LS.MIA->mayAffectControlFlow(DI.Inst, *LS.MRI)) + return true; + if (instructionReadsRegister(DI, LS, Register)) + return true; + if (instructionWritesRegister(DI, LS, Register)) + return false; + } + return false; +} + +static bool isRegisterDefinitelyDeadAtContinuation(PatchContext &Ctx, + uint64_t InstOffset, + uint32_t InstSize, + MCRegister Register) { + std::optional FunctionRange = + Ctx.Elf.findFunctionTextRangeAtOffset(InstOffset); + if (!FunctionRange) + return false; + + std::optional Continuation = checkedAddUint64( + InstOffset, InstSize, "far-return register liveness continuation"); + if (!Continuation) + return false; + std::vector::const_iterator It = + std::lower_bound(Ctx.Decoded.cbegin(), Ctx.Decoded.cend(), *Continuation, + [](const InternalDecodedInst &DI, uint64_t Offset) { + return DI.Offset < Offset; + }); + if (It == Ctx.Decoded.cend() || It->Offset != *Continuation) + return false; + + SmallVector Worklist; + DenseSet Visited; + Worklist.push_back(It - Ctx.Decoded.cbegin()); + while (!Worklist.empty()) { + size_t Index = Worklist.pop_back_val(); + if (!Visited.insert(Index).second) + continue; + const InternalDecodedInst &DI = Ctx.Decoded[Index]; + if (!DI.DecodeSucceeded || !Ctx.LS.MIA || + DI.Offset < FunctionRange->Begin || DI.Offset >= FunctionRange->End) + return false; + if (instructionReadsRegister(DI, Ctx.LS, Register)) + return false; + if (instructionWritesRegister(DI, Ctx.LS, Register) || + DI.Inst.getOpcode() == Ctx.LS.SEndPgmOpcode || + DI.Inst.getOpcode() == Ctx.LS.SEndPgmSavedOpcode) + continue; + + auto AddSuccessor = [&](uint64_t Offset) { + if (Offset < FunctionRange->Begin || Offset >= FunctionRange->End) + return false; + std::vector::const_iterator Successor = + std::lower_bound( + Ctx.Decoded.cbegin(), Ctx.Decoded.cend(), Offset, + [](const InternalDecodedInst &Candidate, uint64_t Target) { + return Candidate.Offset < Target; + }); + if (Successor == Ctx.Decoded.cend() || Successor->Offset != Offset) + return false; + Worklist.push_back(Successor - Ctx.Decoded.cbegin()); + return true; + }; + + if (Ctx.LS.MIA->isCall(DI.Inst) || Ctx.LS.MIA->isIndirectBranch(DI.Inst) || + Ctx.LS.MIA->isReturn(DI.Inst)) + return false; + if (Ctx.LS.MIA->isBranch(DI.Inst)) { + std::optional Target = + evaluateDirectControlFlowTarget(DI, Ctx.LS); + if (!Target || !AddSuccessor(*Target)) + return false; + if (Ctx.LS.MIA->isUnconditionalBranch(DI.Inst)) + continue; + } else if (Ctx.LS.MIA->mayAffectControlFlow(DI.Inst, *Ctx.LS.MRI) && + !Ctx.LS.MIA->isBarrier(DI.Inst)) { + return false; + } + + std::optional Fallthrough = checkedAddUint64( + DI.Offset, DI.Size, "far-return register liveness fallthrough"); + if (!Fallthrough || !AddSuccessor(*Fallthrough)) + return false; + } + return true; +} + +std::optional> +resolveNumberedSgprRegisters(const MCRegisterInfo &MRI, unsigned MaxSgprs) { + SmallVector Registers(MaxSgprs); + for (unsigned I = 1; I != MRI.getNumRegs(); ++I) { + MCRegister Register(I); + std::optional Index = numberedSgprIndex(MRI, Register); + if (Index && *Index < MaxSgprs) + Registers[*Index] = Register; + } + if (llvm::any_of(Registers, + [](MCRegister Register) { return !Register.isValid(); })) + return std::nullopt; + return Registers; +} + +void getNumberedSgprUsesAndDefs(const InternalDecodedInst &DI, + const LLVMState &LS, + ArrayRef NumberedSgprs, + BitVector &Uses, BitVector &Defs) { + assert(Uses.size() == NumberedSgprs.size() && + Defs.size() == NumberedSgprs.size()); + for (unsigned I = 0; I != NumberedSgprs.size(); ++I) { + if (instructionReadsRegister(DI, LS, NumberedSgprs[I])) + Uses.set(I); + if (instructionWritesRegister(DI, LS, NumberedSgprs[I])) + Defs.set(I); + } +} + +/// Return the numbered SGPRs whose incoming values can be observed by the +/// replacement. A malformed or control-flow-bearing replacement conservatively +/// keeps every value that has not already been overwritten. +BitVector +unsafeIncomingNumberedSgprsInReplacement(ArrayRef Replacement, + const LLVMState &LS, + ArrayRef NumberedSgprs) { + const unsigned MaxSgprs = NumberedSgprs.size(); + BitVector Unsafe(MaxSgprs); + BitVector Incoming(MaxSgprs, true); + std::vector Decoded; + if (!decodeTextSection(Replacement.data(), Replacement.size(), LS, Decoded)) { + Unsafe.set(); + return Unsafe; + } + + for (const InternalDecodedInst &DI : Decoded) { + if (!DI.DecodeSucceeded || !LS.MIA) { + Unsafe |= Incoming; + break; + } + BitVector Uses(MaxSgprs); + BitVector Defs(MaxSgprs); + getNumberedSgprUsesAndDefs(DI, LS, NumberedSgprs, Uses, Defs); + Uses &= Incoming; + Unsafe |= Uses; + Incoming.reset(Defs); + if (LS.MIA->mayAffectControlFlow(DI.Inst, *LS.MRI)) { + Unsafe |= Incoming; + break; + } + } + return Unsafe; +} + +/// Analyze all numbered SGPR incoming values in one monotone CFG walk. +/// Unsafe contains a register when some path reads its incoming value before +/// overwriting it, or reaches control flow that cannot be bounded precisely. +std::optional +unsafeIncomingNumberedSgprsInRange(ArrayRef Decoded, + const LLVMState &LS, uint64_t FunctionBegin, + uint64_t FunctionEnd, uint64_t Continuation, + ArrayRef NumberedSgprs) { + const unsigned MaxSgprs = NumberedSgprs.size(); + if (!LS.MIA) return std::nullopt; - if (!commitSafeSgprScratchBlock(Ctx, InstOffset, *Scratch, "safe far return")) + + auto FindInstruction = [&](uint64_t Offset) -> std::optional { + if (Offset < FunctionBegin || Offset >= FunctionEnd) + return std::nullopt; + auto It = + std::lower_bound(Decoded.begin(), Decoded.end(), Offset, + [](const InternalDecodedInst &DI, uint64_t Target) { + return DI.Offset < Target; + }); + if (It == Decoded.end() || It->Offset != Offset) + return std::nullopt; + return It - Decoded.begin(); + }; + std::optional ContinuationIndex = FindInstruction(Continuation); + if (!ContinuationIndex) return std::nullopt; - return Scratch; + + DenseMap IncomingAt; + IncomingAt.try_emplace(*ContinuationIndex, MaxSgprs, true); + SmallVector Worklist(1, *ContinuationIndex); + BitVector Queued(Decoded.size()); + Queued.set(*ContinuationIndex); + BitVector Unsafe(MaxSgprs); + + auto Propagate = [&](uint64_t Offset, const BitVector &Incoming) { + std::optional Successor = FindInstruction(Offset); + if (!Successor) { + Unsafe |= Incoming; + return; + } + auto It = IncomingAt.try_emplace(*Successor, MaxSgprs).first; + BitVector NewValues = Incoming; + NewValues.reset(It->second); + if (NewValues.none()) + return; + It->second |= Incoming; + if (!Queued.test(*Successor)) { + Queued.set(*Successor); + Worklist.push_back(*Successor); + } + }; + + while (!Worklist.empty()) { + size_t Index = Worklist.pop_back_val(); + Queued.reset(Index); + const InternalDecodedInst &DI = Decoded[Index]; + BitVector Incoming = IncomingAt.find(Index)->second; + if (!DI.DecodeSucceeded || DI.Offset < FunctionBegin || + DI.Offset >= FunctionEnd) { + Unsafe |= Incoming; + continue; + } + + BitVector Uses(MaxSgprs); + BitVector Defs(MaxSgprs); + getNumberedSgprUsesAndDefs(DI, LS, NumberedSgprs, Uses, Defs); + Uses &= Incoming; + Unsafe |= Uses; + Incoming.reset(Defs); + if (Incoming.none()) + continue; + + if (DI.Inst.getOpcode() == LS.SEndPgmOpcode || + DI.Inst.getOpcode() == LS.SEndPgmSavedOpcode) + continue; + if (LS.MIA->isCall(DI.Inst) || LS.MIA->isIndirectBranch(DI.Inst) || + LS.MIA->isReturn(DI.Inst)) { + Unsafe |= Incoming; + continue; + } + if (LS.MIA->isBranch(DI.Inst)) { + std::optional Target = evaluateDirectControlFlowTarget(DI, LS); + if (Target) + Propagate(*Target, Incoming); + else + Unsafe |= Incoming; + if (LS.MIA->isUnconditionalBranch(DI.Inst)) + continue; + } else if (LS.MIA->mayAffectControlFlow(DI.Inst, *LS.MRI) && + !LS.MIA->isBarrier(DI.Inst)) { + Unsafe |= Incoming; + continue; + } + + std::optional Fallthrough = checkedAddUint64( + DI.Offset, DI.Size, "far-return SGPR liveness fallthrough"); + if (Fallthrough) + Propagate(*Fallthrough, Incoming); + else + Unsafe |= Incoming; + } + return Unsafe; +} + +static std::optional unsafeIncomingNumberedSgprsAtContinuation( + PatchContext &Ctx, uint64_t InstOffset, uint32_t InstSize, + ArrayRef NumberedSgprs) { + std::optional FunctionRange = + Ctx.Elf.findFunctionTextRangeAtOffset(InstOffset); + if (!FunctionRange) + return std::nullopt; + std::optional Continuation = checkedAddUint64( + InstOffset, InstSize, "far-return SGPR liveness continuation"); + if (!Continuation) + return std::nullopt; + return unsafeIncomingNumberedSgprsInRange( + Ctx.Decoded, Ctx.LS, FunctionRange->Begin, FunctionRange->End, + *Continuation, NumberedSgprs); +} + +static std::optional +findLocallyDeadSgprPair(PatchContext &Ctx, uint64_t InstOffset, + uint32_t InstSize, ArrayRef Replacement) { + if (Ctx.Config.MaxSgprs < 2) + return std::nullopt; + std::optional> NumberedSgprs = + resolveNumberedSgprRegisters(*Ctx.LS.MRI, Ctx.Config.MaxSgprs); + if (!NumberedSgprs) + return std::nullopt; + std::optional ContinuationUnsafe = + unsafeIncomingNumberedSgprsAtContinuation(Ctx, InstOffset, InstSize, + *NumberedSgprs); + if (!ContinuationUnsafe) + return std::nullopt; + BitVector Unsafe = unsafeIncomingNumberedSgprsInReplacement( + Replacement, Ctx.LS, *NumberedSgprs); + Unsafe |= *ContinuationUnsafe; + + unsigned Base = (Ctx.Config.MaxSgprs - 2) & ~1u; + for (;;) { + if (!Unsafe.test(Base) && !Unsafe.test(Base + 1)) { + SafeSgprScratchBlock Scratch{Base, 2}; + if (commitSafeSgprScratchBlock(Ctx, InstOffset, Scratch, + "locally dead far-return SGPR pair")) + return Base; + } + if (Base == 0) + break; + Base -= 2; + } + return std::nullopt; +} + +struct FarReturnScratch { + bool Available = false; + unsigned SgprBase = 0; + bool UseVcc = false; + bool PreserveVcc = false; +}; + +static FarReturnScratch reserveSafeFarReturn(PatchContext &Ctx, + uint64_t InstOffset, + uint32_t InstSize, + ArrayRef Replacement) { + std::optional Scratch = findSafeSgprScratchBlock( + Ctx, InstOffset, /*Count=*/2, /*Alignment=*/2, "safe far return", + /*ReportNoSpace=*/false); + if (Scratch) { + if (!commitSafeSgprScratchBlock(Ctx, InstOffset, *Scratch, + "safe far return")) + return {}; + return FarReturnScratch{/*Available=*/true, Scratch->Base, + /*UseVcc=*/false, /*PreserveVcc=*/false}; + } + + if (Ctx.LS.VCCRegister.isValid() && + !replacementNeedsIncomingRegister(Replacement, Ctx.LS, + Ctx.LS.VCCRegister) && + isRegisterDefinitelyDeadAtContinuation(Ctx, InstOffset, InstSize, + Ctx.LS.VCCRegister)) { + log() << "hotswap: safe far return: reusing dead VCC at 0x" + << utohexstr(InstOffset) << "\n"; + return FarReturnScratch{/*Available=*/true, /*SgprBase=*/0, + /*UseVcc=*/true, /*PreserveVcc=*/false}; + } + + if (std::optional LocalPair = + findLocallyDeadSgprPair(Ctx, InstOffset, InstSize, Replacement)) { + log() << "hotswap: safe far return: reusing locally dead s[" << *LocalPair + << ":" << *LocalPair + 1 << "] at 0x" << utohexstr(InstOffset) + << "\n"; + return FarReturnScratch{/*Available=*/true, *LocalPair, + /*UseVcc=*/false, /*PreserveVcc=*/false}; + } + + if (InstSize >= VccPreservingSourceBytes) { + std::string Owner = + Ctx.Elf.findKernelAtAddress(InstOffset + Ctx.Elf.textAddr()); + std::optional WavefrontSize = + Owner.empty() ? std::nullopt : Ctx.Elf.getKernelWavefrontSize(Owner); + if (WavefrontSize == 32) { + std::optional Save = findSafeSgprScratchBlock( + Ctx, InstOffset, /*Count=*/1, /*Alignment=*/1, + "VCC-preserving far return", /*ReportNoSpace=*/false); + if (Save && commitSafeSgprScratchBlock(Ctx, InstOffset, *Save, + "VCC-preserving far return")) { + log() << "hotswap: safe far return: preserving live wave32 VCC_LO in s" + << Save->Base << " at 0x" << utohexstr(InstOffset) << "\n"; + return FarReturnScratch{/*Available=*/true, Save->Base, + /*UseVcc=*/true, /*PreserveVcc=*/true}; + } + } + } + + log() << "hotswap: safe far return: no register pair at 0x" + << utohexstr(InstOffset) + << "; deferring to the s_branch island planner\n"; + return {}; } bool isSBranchReachable(uint64_t From, uint64_t To) { @@ -819,11 +1318,13 @@ bool isSBranchReachable(uint64_t From, uint64_t To) { /// Queue a deferred trampoline for [\p InstOffset, +\p InstSize) with /// \p Replacement as its body; fixupTrampolineBranches fills in the edges once /// the pool layout is known. A site beyond s_branch reach of the appended pool -/// uses an SCC-neutral get-PC/add/set-PC sequence on the backward edge. +/// uses either an SCC-neutral get-PC/add/set-PC sequence or a chain of +/// registerless s_branch islands on the backward edge. /// Adjacent far sites are coalesced after patching to reduce gateway pressure. -/// Every far source edge then uses a short branch to nearby safe NOP padding; -/// that gateway uses the gfx12 SGPR-backed set-PC sequence. No source or return -/// edge executes gfx1250's broken s_add_pc_i64 instruction. +/// Every far source edge uses a short branch to nearby safe padding; that +/// gateway either continues through s_branch islands or uses the gfx12 +/// SGPR-backed set-PC sequence. No source or return edge executes gfx1250's +/// broken s_add_pc_i64 instruction. [[nodiscard]] bool emitToTrampoline(PatchContext &Ctx, uint64_t InstOffset, uint32_t InstSize, ArrayRef Replacement) { @@ -850,9 +1351,21 @@ bool isSBranchReachable(uint64_t From, uint64_t To) { const bool Far = !(isSBranchReachable(InstOffset, *PoolStart) && isSBranchReachable(*ShortBackFrom, *ReturnTo)); - uint64_t ReturnReserve = Far ? SetPcReturnReserveBytes : MinInstSize; + FarReturnScratch Scratch; + if (Far) + Scratch = reserveSafeFarReturn(Ctx, InstOffset, InstSize, Replacement); + uint64_t ReturnReserve = MinInstSize; + uint64_t BodyPrefix = 0; + if (Far && Scratch.Available) { + ReturnReserve = Scratch.PreserveVcc ? VccPreservingReturnReserveBytes + : SetPcReturnReserveBytes; + BodyPrefix = Scratch.PreserveVcc ? VccSaveRestoreBytes : 0; + } std::optional TrampolineSize = checkedAddUint64( - Replacement.size(), ReturnReserve, "queued trampoline size"); + Replacement.size(), BodyPrefix, "queued trampoline body size"); + if (TrampolineSize) + TrampolineSize = checkedAddUint64(*TrampolineSize, ReturnReserve, + "queued trampoline size"); if (!TrampolineSize) return false; std::optional QueuedBytes = @@ -864,6 +1377,13 @@ bool isSBranchReachable(uint64_t From, uint64_t To) { Trampoline T; T.OriginalOffset = InstOffset; T.OriginalSize = InstSize; + if (Scratch.PreserveVcc) { + SmallVector Restore = assembleSingleInst( + "s_mov_b32 vcc_lo, s" + std::to_string(Scratch.SgprBase), Ctx.LS); + if (Restore.size() != VccSaveRestoreBytes) + return false; + T.Bytes.append(Restore.begin(), Restore.end()); + } T.Bytes.insert(T.Bytes.end(), Replacement.begin(), Replacement.end()); if (std::optional Range = Ctx.Elf.findFunctionTextRangeAtOffset(InstOffset)) { @@ -885,14 +1405,12 @@ bool isSBranchReachable(uint64_t From, uint64_t To) { if (InstSize < MinInstSize) return declineFar(Twine(InstSize) + " B, smaller than " + Twine(MinInstSize) + " B forward branch"); - std::optional Scratch = - reserveSafeFarReturn(Ctx, InstOffset); - if (!Scratch) - return declineFar("no safe SGPR triple for set-PC return"); - T.Bytes.insert(T.Bytes.end(), SetPcReturnReserveBytes, uint8_t{0}); + T.Bytes.insert(T.Bytes.end(), ReturnReserve, uint8_t{0}); T.Long = true; - T.UsesSetPCBack = true; - T.LongBranchSgprBase = Scratch->Base; + T.UsesSetPCBack = Scratch.Available; + T.LongBranchSgprBase = Scratch.SgprBase; + T.LongBranchUsesVcc = Scratch.UseVcc; + T.LongBranchPreservesVcc = Scratch.PreserveVcc; Ctx.Profile.count(HotswapMetric::JumpLong); Ctx.OutTrampolines.emplace_back(std::move(T)); Ctx.QueuedTrampolineBytes = *QueuedBytes; @@ -1131,6 +1649,334 @@ matchPcMaterializedCall(ArrayRef Decoded, size_t CallIndex, MCRegister(Call.Inst.getOperand(0).getReg())}; } +using ReachingCallTargets = SmallVector; + +struct ReachingPcState { + bool Reached = false; + bool HasUnknown = false; + ReachingCallTargets Targets; + SmallVector ActiveMaterializations; +}; + +static bool mergeReachingPcState(ReachingPcState &Into, + const ReachingPcState &From) { + if (!From.Reached) + return false; + ReachingPcState Before = Into; + Into.Reached = true; + Into.HasUnknown |= From.HasUnknown; + for (uint64_t Target : From.Targets) + if (!llvm::is_contained(Into.Targets, Target)) + Into.Targets.push_back(Target); + for (size_t Completion : From.ActiveMaterializations) + if (!llvm::is_contained(Into.ActiveMaterializations, Completion)) + Into.ActiveMaterializations.push_back(Completion); + llvm::sort(Into.Targets); + llvm::sort(Into.ActiveMaterializations); + return Before.Reached != Into.Reached || + Before.HasUnknown != Into.HasUnknown || + Before.Targets != Into.Targets || + Before.ActiveMaterializations != Into.ActiveMaterializations; +} + +static bool isExactRegisterOperand(const MCInst &Inst, unsigned Index, + MCRegister Reg) { + return Index < Inst.getNumOperands() && Inst.getOperand(Index).isReg() && + Inst.getOperand(Index).getReg() == Reg; +} + +/// Recognize the two compiler-emitted ways a reusable register-call target is +/// materialized. The first is the canonical get-PC/add-nc pair. Tensile also +/// computes a 32-bit displacement in a temporary and propagates carry into the +/// high half: +/// +/// s_get_pc_i64 Pair +/// s_add_co_i32 Tmp, Imm0, Imm1 +/// s_add_co_u32 Pair.lo, Pair.lo, Tmp +/// s_add_co_ci_u32 Pair.hi, Pair.hi, 0 +/// +/// Return the completion instruction and absolute target. Intermediate +/// definitions deliberately remain "unknown" in the reaching-value solver. +static std::optional> +matchReusablePcMaterialization(ArrayRef Decoded, + size_t GetPcIndex, size_t FunctionEndIndex, + MCRegister Pair, const LLVMState &LS, + uint64_t TextAddr, ArrayRef Text) { + const InternalDecodedInst &GetPc = Decoded[GetPcIndex]; + if (!GetPc.DecodeSucceeded || GetPc.Inst.getOpcode() != LS.SGetPcI64Opcode || + !isExactRegisterOperand(GetPc.Inst, 0, Pair)) + return std::nullopt; + std::optional Pc = + checkedAddUint64(TextAddr, GetPc.Offset, "reusable get-PC address"); + if (!Pc) + return std::nullopt; + Pc = checkedAddUint64(*Pc, GetPc.Size, "reusable get-PC value"); + if (!Pc) + return std::nullopt; + + for (size_t I = GetPcIndex + 1; I < FunctionEndIndex; ++I) { + const InternalDecodedInst &DI = Decoded[I]; + if (!DI.DecodeSucceeded || isControlFlowBoundary(DI, LS)) + break; + if (!definesOverlappingRegister(DI, LS, Pair)) + continue; + if (DI.Inst.getOpcode() != LS.SAddNcU64Opcode || + DI.Inst.getNumOperands() != 3 || + !isExactRegisterOperand(DI.Inst, 0, Pair) || + !isExactRegisterOperand(DI.Inst, 1, Pair) || + !DI.Inst.getOperand(2).isImm()) + break; + uint64_t Delta = static_cast(DI.Inst.getOperand(2).getImm()); + return std::make_pair(I, *Pc + Delta); + } + + if (GetPcIndex + 3 >= FunctionEndIndex) + return std::nullopt; + const InternalDecodedInst &MakeDelta = Decoded[GetPcIndex + 1]; + const InternalDecodedInst &AddLow = Decoded[GetPcIndex + 2]; + const InternalDecodedInst &AddHigh = Decoded[GetPcIndex + 3]; + if (MakeDelta.Mnemonic != "s_add_co_i32" || + AddLow.Mnemonic != "s_add_co_u32" || + AddHigh.Mnemonic != "s_add_co_ci_u32" || + MakeDelta.Inst.getNumOperands() != 3 || + !MakeDelta.Inst.getOperand(0).isReg() || + !MakeDelta.Inst.getOperand(0).getReg() || + !MakeDelta.Inst.getOperand(2).isImm()) + return std::nullopt; + MCRegister DeltaReg(MakeDelta.Inst.getOperand(0).getReg()); + if (AddLow.Inst.getNumOperands() != 3 || !AddLow.Inst.getOperand(0).isReg() || + !AddLow.Inst.getOperand(1).isReg() || + !AddLow.Inst.getOperand(0).getReg() || + AddLow.Inst.getOperand(0).getReg() != + AddLow.Inst.getOperand(1).getReg() || + !isExactRegisterOperand(AddLow.Inst, 2, DeltaReg) || + AddHigh.Inst.getNumOperands() != 3 || + !AddHigh.Inst.getOperand(0).isReg() || + !AddHigh.Inst.getOperand(1).isReg() || + !AddHigh.Inst.getOperand(0).getReg() || + AddHigh.Inst.getOperand(0).getReg() != + AddHigh.Inst.getOperand(1).getReg() || + !AddHigh.Inst.getOperand(2).isImm() || + AddHigh.Inst.getOperand(2).getImm() != 0) + return std::nullopt; + MCRegister Low(AddLow.Inst.getOperand(0).getReg()); + MCRegister High(AddHigh.Inst.getOperand(0).getReg()); + std::optional LowIndex = numberedSgprIndex(*LS.MRI, Low); + std::optional HighIndex = numberedSgprIndex(*LS.MRI, High); + if (!LowIndex || !HighIndex || *HighIndex != *LowIndex + 1 || + !LS.MRI->regsOverlap(Low, Pair) || !LS.MRI->regsOverlap(High, Pair)) + return std::nullopt; + + std::optional FirstAddend; + if (MakeDelta.Inst.getOperand(1).isImm()) { + FirstAddend = static_cast(MakeDelta.Inst.getOperand(1).getImm()); + } else if (MakeDelta.Size >= 2 * MinInstSize && + MakeDelta.Offset <= Text.size() && + MakeDelta.Size <= Text.size() - MakeDelta.Offset) { + // The disassembler represents a SOP2 literal source as its literal + // register marker, not as an immediate operand. The literal dword is the + // final dword in this otherwise exact compiler-emitted instruction. + FirstAddend = support::endian::read32le(Text.data() + MakeDelta.Offset + + MakeDelta.Size - MinInstSize); + } + if (!FirstAddend) + return std::nullopt; + uint32_t Delta = *FirstAddend + + static_cast(MakeDelta.Inst.getOperand(2).getImm()); + return std::make_pair(GetPcIndex + 3, *Pc + Delta); +} + +struct ReachingCallGroup { + uint64_t Begin = 0; + uint64_t End = 0; + MCRegister TargetRegister; + SmallVector Calls; +}; + +/// Resolve register calls whose target pair is selected once and reused across +/// control flow. A monotone intraprocedural solver propagates finite target +/// sets from proven get-PC materializations. Any unrecognized pair definition +/// introduces Unknown, so a bypass around the selector remains fail-closed. +static std::vector resolveReusablePcCallTargets( + ArrayRef Decoded, const LLVMState &LS, + uint64_t TextAddr, uint64_t TextEnd, + ArrayRef FunctionRanges, + ArrayRef Text) { + std::vector Resolved(Decoded.size()); + SmallVector Groups; + + for (size_t I = 0; I != Decoded.size(); ++I) { + const InternalDecodedInst &Call = Decoded[I]; + if (!Call.DecodeSucceeded || Call.Inst.getOpcode() != LS.SSwapPcI64Opcode || + Call.Inst.getNumOperands() < 2) + continue; + const MCOperand &TargetOp = + Call.Inst.getOperand(Call.Inst.getNumOperands() - 1); + if (!TargetOp.isReg() || !TargetOp.getReg()) + continue; + MCRegister TargetRegister(TargetOp.getReg()); + + const ElfView::FunctionTextRange *Best = nullptr; + uint64_t Address = TextAddr + Call.Offset; + for (const ElfView::FunctionTextRange &Range : FunctionRanges) + if (Range.Begin <= Address && Address < Range.End && + (!Best || Range.Begin > Best->Begin)) + Best = &Range; + if (!Best || Best->Begin < TextAddr || Best->End > TextEnd) + continue; + uint64_t Begin = Best->Begin - TextAddr; + uint64_t End = Best->End - TextAddr; + ReachingCallGroup *Group = nullptr; + for (ReachingCallGroup &Candidate : Groups) + if (Candidate.Begin == Begin && Candidate.End == End && + Candidate.TargetRegister == TargetRegister) { + Group = &Candidate; + break; + } + if (!Group) { + Groups.push_back({Begin, End, TargetRegister, {}}); + Group = &Groups.back(); + } + Group->Calls.push_back(I); + } + + DenseMap OffsetToIndex; + for (size_t I = 0; I != Decoded.size(); ++I) + OffsetToIndex[Decoded[I].Offset] = I; + + for (const ReachingCallGroup &Group : Groups) { + ArrayRef::const_iterator Begin = + llvm::lower_bound(Decoded, Group.Begin, + [](const InternalDecodedInst &DI, uint64_t Offset) { + return DI.Offset < Offset; + }); + ArrayRef::const_iterator End = std::lower_bound( + Begin, Decoded.end(), Group.End, + [](const InternalDecodedInst &DI, uint64_t Offset) { + return DI.Offset < Offset; + }); + size_t BeginIndex = static_cast(Begin - Decoded.begin()); + size_t EndIndex = static_cast(End - Decoded.begin()); + if (BeginIndex == EndIndex) + continue; + + DenseMap Starters; + DenseMap Completions; + DenseMap> Intermediates; + for (size_t I = BeginIndex; I != EndIndex; ++I) { + std::optional> Match = + matchReusablePcMaterialization( + Decoded, I, EndIndex, Group.TargetRegister, LS, TextAddr, Text); + if (Match) { + Starters[I] = Match->first; + Completions[Match->first] = Match->second; + for (size_t J = I + 1; J != Match->first; ++J) + if (definesOverlappingRegister(Decoded[J], LS, Group.TargetRegister)) + Intermediates[J].push_back(Match->first); + } + } + if (Completions.empty()) + continue; + + std::vector Before(EndIndex - BeginIndex); + Before.front().Reached = true; + Before.front().HasUnknown = true; + SmallVector Worklist; + BitVector Queued(EndIndex - BeginIndex); + Worklist.push_back(BeginIndex); + Queued.set(0); + + while (!Worklist.empty()) { + size_t I = Worklist.pop_back_val(); + Queued.reset(I - BeginIndex); + ReachingPcState State = Before[I - BeginIndex]; + const InternalDecodedInst &DI = Decoded[I]; + + DenseMap::const_iterator Starter = Starters.find(I); + DenseMap::const_iterator Completion = + Completions.find(I); + if (Starter != Starters.end()) { + // The get-PC instruction overwrites the complete target pair. Record a + // token proving that this path entered the exact materialization; the + // completion may only produce a known target from that token. + State.HasUnknown = false; + State.Targets.clear(); + State.ActiveMaterializations.assign(1, Starter->second); + } else if (Completion != Completions.end()) { + bool HasMatchingToken = + llvm::is_contained(State.ActiveMaterializations, I); + bool HasBypassPath = + State.HasUnknown || !State.Targets.empty() || + llvm::any_of(State.ActiveMaterializations, + [I](size_t Active) { return Active != I; }); + State.HasUnknown = HasBypassPath || !HasMatchingToken; + State.Targets.clear(); + State.ActiveMaterializations.clear(); + if (HasMatchingToken) + State.Targets.push_back(Completion->second); + } else if (definesOverlappingRegister(DI, LS, Group.TargetRegister)) { + // Preserve only tokens for which this is a proven instruction inside + // the exact matched sequence. All other reaching values are clobbered. + SmallVector Preserved; + DenseMap>::const_iterator Intermediate = + Intermediates.find(I); + if (Intermediate != Intermediates.end()) + for (size_t Active : State.ActiveMaterializations) + if (llvm::is_contained(Intermediate->second, Active)) + Preserved.push_back(Active); + if (State.HasUnknown || !State.Targets.empty() || + Preserved.size() != State.ActiveMaterializations.size()) + State.HasUnknown = true; + State.Targets.clear(); + State.ActiveMaterializations = std::move(Preserved); + } + + if (llvm::is_contained(Group.Calls, I) && !State.HasUnknown && + State.ActiveMaterializations.empty() && !State.Targets.empty()) + Resolved[I] = State.Targets; + + SmallVector Successors; + auto appendFallthrough = [&]() { + if (I + 1 < EndIndex) + Successors.push_back(I + 1); + }; + if (DI.Inst.getOpcode() == LS.SEndPgmOpcode || + DI.Inst.getOpcode() == LS.SEndPgmSavedOpcode || + LS.MIA->isReturn(DI.Inst)) { + // No successor. + } else if (LS.MIA->isCall(DI.Inst)) { + appendFallthrough(); + } else if (LS.MIA->isIndirectBranch(DI.Inst)) { + // An indirect jump or bounded return leaves this intraprocedural path. + } else if (LS.MIA->isBranch(DI.Inst)) { + std::optional Target = + evaluateDirectControlFlowTarget(DI, LS); + if (Target) { + DenseMap::const_iterator TargetIndex = + OffsetToIndex.find(*Target); + if (TargetIndex != OffsetToIndex.end() && + TargetIndex->second >= BeginIndex && + TargetIndex->second < EndIndex) + Successors.push_back(TargetIndex->second); + } + if (!LS.MIA->isUnconditionalBranch(DI.Inst)) + appendFallthrough(); + } else { + appendFallthrough(); + } + + for (size_t Successor : Successors) { + if (mergeReachingPcState(Before[Successor - BeginIndex], State) && + !Queued.test(Successor - BeginIndex)) { + Worklist.push_back(Successor); + Queued.set(Successor - BeginIndex); + } + } + } + } + return Resolved; +} + struct KnownCallSite { size_t InstIndex = 0; uint64_t Target = 0; @@ -1143,6 +1989,26 @@ struct BoundedSetPcReturn { SmallVector Targets; }; +struct DirectTargetSource { + size_t InstIndex = 0; + uint64_t Target = 0; +}; + +struct FallthroughEntryInfo { + bool Proven = false; + uint64_t ChainBegin = 0; +}; + +struct ControlFlowScanIndex { + DenseMap MaterializedCalls; + DenseMap FallthroughEntries; + SmallVector Calls; + SmallVector SetPcIndices; + SmallVector BranchOrCallIndices; + SmallVector DirectTargetsByTarget; + bool HasUnboundedIndirectEntry = false; +}; + static bool hasPcRelativeOperand(const InternalDecodedInst &DI, const LLVMState &LS) { for (const MCOperandInfo &Operand : @@ -1183,78 +2049,122 @@ getDirectTextTarget(const InternalDecodedInst &DI, const LLVMState &LS, return AbsoluteTarget - TextAddr; } -static std::optional> collectKnownCallSites( - ArrayRef Decoded, const LLVMState &LS, - uint64_t TextAddr, uint64_t TextEnd, - ArrayRef> MaterializedCalls) { - SmallVector Calls; +static std::optional +buildControlFlowScanIndex(ArrayRef Decoded, + const LLVMState &LS, uint64_t TextAddr, + uint64_t TextEnd, + ArrayRef FunctionRanges) { + ControlFlowScanIndex Index; + DenseSet FunctionBegins; + for (const ElfView::FunctionTextRange &Range : FunctionRanges) + if (Range.Begin >= TextAddr && Range.Begin < TextEnd) + FunctionBegins.insert(Range.Begin - TextAddr); + + bool FallthroughProven = true; + uint64_t FallthroughChainBegin = 0; for (size_t I = 0; I != Decoded.size(); ++I) { const InternalDecodedInst &DI = Decoded[I]; - std::optional ReturnRegister = getCallReturnRegister(DI, LS); - if (!ReturnRegister) - continue; - - std::optional Target; - if (MaterializedCalls[I]) { - uint64_t AbsoluteTarget = MaterializedCalls[I]->Target; - if (AbsoluteTarget >= TextAddr && AbsoluteTarget < TextEnd) - Target = AbsoluteTarget - TextAddr; + if (I == 0) { + FallthroughChainBegin = DI.Offset; } else { - Target = getDirectTextTarget(DI, LS, TextAddr, TextEnd); + const InternalDecodedInst &Predecessor = Decoded[I - 1]; + bool EndOverflows = + Predecessor.Offset > + std::numeric_limits::max() - Predecessor.Size; + if (EndOverflows || Predecessor.Offset + Predecessor.Size != DI.Offset || + !Predecessor.DecodeSucceeded) { + FallthroughProven = false; + FallthroughChainBegin = DI.Offset; + } else if (LS.MIA->isBarrier(Predecessor.Inst)) { + FallthroughProven = true; + FallthroughChainBegin = DI.Offset; + } } - if (!Target) + if (FunctionBegins.contains(DI.Offset)) + Index.FallthroughEntries.try_emplace( + DI.Offset, + FallthroughEntryInfo{FallthroughProven, FallthroughChainBegin}); + + std::optional Materialized = + matchPcMaterializedCall(Decoded, I, LS, TextAddr); + if (Materialized) + Index.MaterializedCalls.try_emplace(I, *Materialized); + + std::optional ReturnRegister = getCallReturnRegister(DI, LS); + if (ReturnRegister) { + std::optional Target; + if (Materialized) { + uint64_t AbsoluteTarget = Materialized->Target; + if (AbsoluteTarget >= TextAddr && AbsoluteTarget < TextEnd) + Target = AbsoluteTarget - TextAddr; + } else { + Target = getDirectTextTarget(DI, LS, TextAddr, TextEnd); + } + if (Target) { + std::optional Continuation = checkedAddUint64( + DI.Offset, DI.Size, "known call continuation address"); + if (!Continuation) + return std::nullopt; + Index.Calls.push_back({I, *Target, *Continuation, *ReturnRegister}); + } + } + + if (DI.DecodeSucceeded && DI.Inst.getOpcode() == LS.SSetPcI64Opcode) + Index.SetPcIndices.push_back(I); + + // Set-PC returns are checked separately against BoundedReturnPositions. + // MC lowering erases their return pseudo identity, so including them in + // this generic bucket would make even a proven bounded return look like + // an arbitrary object-wide entry. + if (DI.DecodeSucceeded && DI.Inst.getOpcode() != LS.SSetPcI64Opcode && + !LS.MIA->isReturn(DI.Inst) && + (LS.MIA->isIndirectBranch(DI.Inst) || + DI.Inst.getOpcode() == LS.SAddPcI64Opcode)) + Index.HasUnboundedIndirectEntry = true; + + if ((!LS.MIA->isBranch(DI.Inst) && !LS.MIA->isCall(DI.Inst)) || + LS.MIA->isReturn(DI.Inst)) continue; + Index.BranchOrCallIndices.push_back(I); - std::optional Continuation = - checkedAddUint64(DI.Offset, DI.Size, "known call continuation address"); - if (!Continuation) - return std::nullopt; - Calls.push_back({I, *Target, *Continuation, *ReturnRegister}); + std::optional DirectTarget = + getDirectTextTarget(DI, LS, TextAddr, TextEnd); + if (DirectTarget) + Index.DirectTargetsByTarget.push_back({I, *DirectTarget}); } - return Calls; + llvm::sort(Index.DirectTargetsByTarget, + [](const DirectTargetSource &LHS, const DirectTargetSource &RHS) { + return std::tie(LHS.Target, LHS.InstIndex) < + std::tie(RHS.Target, RHS.InstIndex); + }); + return Index; } static bool hasUnprovenFallthroughEntry(ArrayRef Decoded, - const LLVMState &LS, uint64_t TextAddr, - uint64_t TextEnd, uint64_t FunctionBegin, uint64_t ReturnOffset, ArrayRef DeclaredEntries, - ArrayRef Calls) { + const ControlFlowScanIndex &Index) { if (FunctionBegin == 0) return false; - size_t BeginIndex = 0; - while (BeginIndex != Decoded.size() && - Decoded[BeginIndex].Offset < FunctionBegin) - ++BeginIndex; - if (BeginIndex == Decoded.size() || - Decoded[BeginIndex].Offset != FunctionBegin) { + DenseMap::const_iterator Fallthrough = + Index.FallthroughEntries.find(FunctionBegin); + if (Fallthrough == Index.FallthroughEntries.end()) { log() << "hotswap: s_set_pc_i64 at 0x" << utohexstr(ReturnOffset) << " is not a bounded return: function entry at 0x" << utohexstr(FunctionBegin) << " is not an instruction boundary\n"; return true; } - uint64_t ChainBegin = FunctionBegin; - size_t PredecessorIndex = BeginIndex; - while (PredecessorIndex != 0) { - const InternalDecodedInst &Predecessor = Decoded[PredecessorIndex - 1]; - std::optional PredecessorEnd = checkedAddUint64( - Predecessor.Offset, Predecessor.Size, "fallthrough predecessor end"); - if (!PredecessorEnd || *PredecessorEnd != ChainBegin || - !Predecessor.DecodeSucceeded) { - log() << "hotswap: s_set_pc_i64 at 0x" << utohexstr(ReturnOffset) - << " is not a bounded return: fallthrough into function entry " - "at 0x" - << utohexstr(FunctionBegin) << " is unprovable\n"; - return true; - } - if (LS.MIA->isBarrier(Predecessor.Inst)) - break; - ChainBegin = Predecessor.Offset; - --PredecessorIndex; + if (!Fallthrough->second.Proven) { + log() << "hotswap: s_set_pc_i64 at 0x" << utohexstr(ReturnOffset) + << " is not a bounded return: fallthrough into function entry " + "at 0x" + << utohexstr(FunctionBegin) << " is unprovable\n"; + return true; } + uint64_t ChainBegin = Fallthrough->second.ChainBegin; if (ChainBegin == FunctionBegin) return false; @@ -1268,7 +2178,7 @@ static bool hasUnprovenFallthroughEntry(ArrayRef Decoded, return true; } - for (const KnownCallSite &Call : Calls) { + for (const KnownCallSite &Call : Index.Calls) { uint64_t Source = Decoded[Call.InstIndex].Offset; if (Source >= ChainBegin && Source < FunctionBegin) continue; @@ -1287,16 +2197,30 @@ static bool hasUnprovenFallthroughEntry(ArrayRef Decoded, } } - for (const InternalDecodedInst &Source : Decoded) { - std::optional Target = - getDirectTextTarget(Source, LS, TextAddr, TextEnd); - if (!Target || *Target < ChainBegin || *Target >= FunctionBegin || - (Source.Offset >= ChainBegin && Source.Offset < FunctionBegin)) + SmallVector::const_iterator FirstTarget = + llvm::lower_bound(Index.DirectTargetsByTarget, ChainBegin, + [](const DirectTargetSource &Source, uint64_t Target) { + return Source.Target < Target; + }); + size_t FirstSourceIndex = Decoded.size(); + uint64_t FirstSourceTarget = 0; + for (SmallVector::const_iterator It = FirstTarget; + It != Index.DirectTargetsByTarget.end() && It->Target < FunctionBegin; + ++It) { + const InternalDecodedInst &Source = Decoded[It->InstIndex]; + if (Source.Offset >= ChainBegin && Source.Offset < FunctionBegin) continue; + if (It->InstIndex < FirstSourceIndex) { + FirstSourceIndex = It->InstIndex; + FirstSourceTarget = It->Target; + } + } + if (FirstSourceIndex != Decoded.size()) { log() << "hotswap: s_set_pc_i64 at 0x" << utohexstr(ReturnOffset) << " is not a bounded return: control flow at 0x" - << utohexstr(Source.Offset) << " enters the fallthrough chain at 0x" - << utohexstr(*Target) << "\n"; + << utohexstr(Decoded[FirstSourceIndex].Offset) + << " enters the fallthrough chain at 0x" + << utohexstr(FirstSourceTarget) << "\n"; return true; } return false; @@ -1307,10 +2231,36 @@ collectBoundedSetPcReturns(ArrayRef Decoded, const LLVMState &LS, uint64_t TextAddr, uint64_t TextEnd, ArrayRef DeclaredEntries, ArrayRef FunctionRanges, - ArrayRef Calls, - ArrayRef ExternalEntries) { + ArrayRef ExternalEntries, + const ControlFlowScanIndex &Index) { SmallVector Returns; - for (size_t ReturnIndex = 0; ReturnIndex != Decoded.size(); ++ReturnIndex) { + SmallVector, 16> CandidateRanges( + Index.SetPcIndices.size()); + for (size_t RangeIndex = 0; RangeIndex != FunctionRanges.size(); + ++RangeIndex) { + const ElfView::FunctionTextRange &Range = FunctionRanges[RangeIndex]; + if (Range.End <= Range.Begin) + continue; + SmallVector::const_iterator First = llvm::lower_bound( + Index.SetPcIndices, Range.Begin, + [&](size_t InstIndex, uint64_t Address) { + return TextAddr + Decoded[InstIndex].Offset < Address; + }); + SmallVector::const_iterator After = std::lower_bound( + First, Index.SetPcIndices.end(), Range.End, + [&](size_t InstIndex, uint64_t Address) { + return TextAddr + Decoded[InstIndex].Offset < Address; + }); + for (SmallVector::const_iterator It = First; It != After; + ++It) { + size_t Position = static_cast(It - Index.SetPcIndices.begin()); + CandidateRanges[Position].push_back(RangeIndex); + } + } + + for (size_t ReturnPosition = 0; ReturnPosition != Index.SetPcIndices.size(); + ++ReturnPosition) { + size_t ReturnIndex = Index.SetPcIndices[ReturnPosition]; const InternalDecodedInst &Return = Decoded[ReturnIndex]; // AMDGPUMCInstLower lowers S_SETPC_B64_return to S_SETPC_B64, so the // decoded instruction no longer carries MIA::isReturn identity. Recover @@ -1324,7 +2274,8 @@ collectBoundedSetPcReturns(ArrayRef Decoded, MCRegister ReturnRegister(Return.Inst.getOperand(0).getReg()); bool IsBounded = false; - for (const ElfView::FunctionTextRange &Range : FunctionRanges) { + for (size_t RangeIndex : CandidateRanges[ReturnPosition]) { + const ElfView::FunctionTextRange &Range = FunctionRanges[RangeIndex]; if (Range.Begin < TextAddr || Range.Begin >= TextEnd || Range.End <= Range.Begin || Range.End > TextEnd || (Range.Symbol && Range.Symbol->getBinding() != ELF::STB_LOCAL)) @@ -1347,9 +2298,9 @@ collectBoundedSetPcReturns(ArrayRef Decoded, if (!Safe) continue; - for (const ElfView::FunctionTextRange &Alias : FunctionRanges) { - if (Return.Offset + TextAddr < Alias.Begin || - Return.Offset + TextAddr >= Alias.End || Alias.Begin == Range.Begin) + for (size_t AliasIndex : CandidateRanges[ReturnPosition]) { + const ElfView::FunctionTextRange &Alias = FunctionRanges[AliasIndex]; + if (Alias.Begin == Range.Begin) continue; log() << "hotswap: s_set_pc_i64 at 0x" << utohexstr(Return.Offset) << " is not a bounded return: overlapping function entry at " @@ -1373,17 +2324,26 @@ collectBoundedSetPcReturns(ArrayRef Decoded, if (!Safe) continue; - if (hasUnprovenFallthroughEntry(Decoded, LS, TextAddr, TextEnd, - FunctionBegin, Return.Offset, - DeclaredEntries, Calls)) + if (hasUnprovenFallthroughEntry(Decoded, FunctionBegin, Return.Offset, + DeclaredEntries, Index)) continue; // The link pair must retain the value written by the incoming call // throughout the function. This includes blocks laid out after the // return that may branch back into its epilogue. - for (const InternalDecodedInst &DI : Decoded) { - if (DI.Offset < FunctionBegin || DI.Offset >= FunctionEnd) - continue; + ArrayRef::const_iterator FunctionFirst = + llvm::lower_bound(Decoded, FunctionBegin, + [](const InternalDecodedInst &DI, uint64_t Offset) { + return DI.Offset < Offset; + }); + ArrayRef::const_iterator FunctionAfter = + std::lower_bound(FunctionFirst, Decoded.end(), FunctionEnd, + [](const InternalDecodedInst &DI, uint64_t Offset) { + return DI.Offset < Offset; + }); + for (ArrayRef::const_iterator It = FunctionFirst; + It != FunctionAfter; ++It) { + const InternalDecodedInst &DI = *It; // MC call instructions carry no transitive callee-clobber information. // Without interprocedural proof, a nested callee may overwrite the // outer link pair even when the call defines a different return pair. @@ -1408,7 +2368,7 @@ collectBoundedSetPcReturns(ArrayRef Decoded, continue; SmallVector Targets; - for (const KnownCallSite &Call : Calls) { + for (const KnownCallSite &Call : Index.Calls) { if (Call.Target < FunctionBegin || Call.Target >= FunctionEnd) continue; if (Call.Target != FunctionBegin) { @@ -1436,29 +2396,39 @@ collectBoundedSetPcReturns(ArrayRef Decoded, // A branch from outside the function would bypass the call link // definition. Direct calls to the function entry are allowed only when // they were collected above with this exact return register. - for (size_t SourceIndex = 0; SourceIndex != Decoded.size(); - ++SourceIndex) { + SmallVector::const_iterator FirstTarget = + llvm::lower_bound( + Index.DirectTargetsByTarget, FunctionBegin, + [](const DirectTargetSource &Source, uint64_t Target) { + return Source.Target < Target; + }); + size_t FirstUnsafeSourceIndex = Decoded.size(); + uint64_t FirstUnsafeTarget = 0; + for (SmallVector::const_iterator It = FirstTarget; + It != Index.DirectTargetsByTarget.end() && It->Target < FunctionEnd; + ++It) { + size_t SourceIndex = It->InstIndex; const InternalDecodedInst &Source = Decoded[SourceIndex]; - std::optional Target = - getDirectTextTarget(Source, LS, TextAddr, TextEnd); - if (!Target || *Target < FunctionBegin || *Target >= FunctionEnd || - (Source.Offset >= FunctionBegin && Source.Offset < FunctionEnd)) + if (Source.Offset >= FunctionBegin && Source.Offset < FunctionEnd) continue; bool IsKnownEntryCall = false; - if (LS.MIA->isCall(Source.Inst) && *Target == FunctionBegin) - for (const KnownCallSite &Call : Calls) + if (LS.MIA->isCall(Source.Inst) && It->Target == FunctionBegin) + for (const KnownCallSite &Call : Index.Calls) IsKnownEntryCall |= Call.InstIndex == SourceIndex && Call.ReturnRegister == ReturnRegister; - if (!IsKnownEntryCall) { - log() << "hotswap: s_set_pc_i64 at 0x" << utohexstr(Return.Offset) - << " is not a bounded return: control flow at 0x" - << utohexstr(Source.Offset) << " enters at 0x" - << utohexstr(*Target) << "\n"; - Safe = false; - break; + if (!IsKnownEntryCall && SourceIndex < FirstUnsafeSourceIndex) { + FirstUnsafeSourceIndex = SourceIndex; + FirstUnsafeTarget = It->Target; } } + if (FirstUnsafeSourceIndex != Decoded.size()) { + log() << "hotswap: s_set_pc_i64 at 0x" << utohexstr(Return.Offset) + << " is not a bounded return: control flow at 0x" + << utohexstr(Decoded[FirstUnsafeSourceIndex].Offset) + << " enters at 0x" << utohexstr(FirstUnsafeTarget) << "\n"; + Safe = false; + } if (!Safe) continue; @@ -1475,54 +2445,41 @@ collectBoundedSetPcReturns(ArrayRef Decoded, return Returns; } -static const BoundedSetPcReturn * -findBoundedSetPcReturn(ArrayRef Returns, size_t InstIndex) { - for (const BoundedSetPcReturn &Return : Returns) - if (Return.InstIndex == InstIndex) - return &Return; - return nullptr; -} - static bool -hasKnownControlFlowEntry(ArrayRef Decoded, - const LLVMState &LS, uint64_t TextAddr, - uint64_t TextEnd, ArrayRef DeclaredEntries, +hasKnownControlFlowEntry(ArrayRef DeclaredEntries, ArrayRef BoundedReturns, + const DenseMap &BoundedReturnPositions, + const ControlFlowScanIndex &Index, uint64_t SequenceStart, uint64_t SequenceEnd) { for (uint64_t Entry : DeclaredEntries) if (Entry > SequenceStart && Entry <= SequenceEnd) return true; - for (size_t I = 0; I != Decoded.size(); ++I) { - const InternalDecodedInst &DI = Decoded[I]; - if (DI.DecodeSucceeded && DI.Inst.getOpcode() == LS.SSetPcI64Opcode) { - const BoundedSetPcReturn *Return = - findBoundedSetPcReturn(BoundedReturns, I); - if (!Return) + for (size_t InstIndex : Index.SetPcIndices) { + DenseMap::const_iterator It = + BoundedReturnPositions.find(InstIndex); + if (It == BoundedReturnPositions.end()) + return true; + const BoundedSetPcReturn &Return = BoundedReturns[It->second]; + for (uint64_t Target : Return.Targets) + if (Target > SequenceStart && Target <= SequenceEnd) return true; - for (uint64_t Target : Return->Targets) - if (Target > SequenceStart && Target <= SequenceEnd) - return true; - continue; - } + } - // Without bounding an indirect target, it may enter at any instruction in - // the materialization. Keep the call unresolved rather than relying on - // the indirect transfer's containing function alone. - if (DI.DecodeSucceeded && !LS.MIA->isReturn(DI.Inst) && - (LS.MIA->isIndirectBranch(DI.Inst) || - DI.Inst.getOpcode() == LS.SAddPcI64Opcode)) - return true; - if ((!LS.MIA->isBranch(DI.Inst) && !LS.MIA->isCall(DI.Inst)) || - LS.MIA->isReturn(DI.Inst)) - continue; + // Without bounding an indirect target, it may enter at any instruction in + // the materialization. Keep the call unresolved rather than relying on the + // indirect transfer's containing function alone. + if (Index.HasUnboundedIndirectEntry) + return true; - std::optional RelativeTarget = - getDirectTextTarget(DI, LS, TextAddr, TextEnd); - if (RelativeTarget && *RelativeTarget > SequenceStart && - *RelativeTarget <= SequenceEnd) - return true; - } + SmallVector::const_iterator First = + llvm::upper_bound(Index.DirectTargetsByTarget, SequenceStart, + [](uint64_t Target, const DirectTargetSource &Source) { + return Target < Source.Target; + }); + if (First != Index.DirectTargetsByTarget.end() && + First->Target <= SequenceEnd) + return true; return false; } @@ -1532,7 +2489,7 @@ std::optional collectDirectBranchTargets( ArrayRef Decoded, const LLVMState &LS, uint64_t TextAddr, uint64_t TextSize, ArrayRef DeclaredEntries, ArrayRef FunctionRanges, - ArrayRef ExternalEntries) { + ArrayRef ExternalEntries, ArrayRef Text) { if (!LS.MIA) { log() << "hotswap: MC branch analysis is unavailable; adjacent far " "trampolines will not be coalesced\n"; @@ -1544,28 +2501,69 @@ std::optional collectDirectBranchTargets( if (!TextEnd) return std::nullopt; - SmallVector, 16> MaterializedCalls( - Decoded.size()); - for (size_t I = 0; I != Decoded.size(); ++I) - MaterializedCalls[I] = matchPcMaterializedCall(Decoded, I, LS, TextAddr); + std::vector ReusableCalls = resolveReusablePcCallTargets( + Decoded, LS, TextAddr, *TextEnd, FunctionRanges, Text); - std::optional> Calls = - collectKnownCallSites(Decoded, LS, TextAddr, *TextEnd, MaterializedCalls); - if (!Calls) + std::optional Index = buildControlFlowScanIndex( + Decoded, LS, TextAddr, *TextEnd, FunctionRanges); + if (!Index) return std::nullopt; + + // Reusable targets are intentionally lower priority than an exact canonical + // one-shot materialization, matching the established call-site proof. Add + // only the non-canonical finite target sets to the sparse call index. + for (size_t I = 0; I != ReusableCalls.size(); ++I) { + if (ReusableCalls[I].empty() || Index->MaterializedCalls.contains(I)) + continue; + std::optional ReturnRegister = + getCallReturnRegister(Decoded[I], LS); + if (!ReturnRegister || + llvm::any_of(ReusableCalls[I], [TextAddr, TextEnd](uint64_t Target) { + return Target < TextAddr || Target >= *TextEnd; + })) + continue; + std::optional Continuation = checkedAddUint64( + Decoded[I].Offset, Decoded[I].Size, + "known reusable call continuation address"); + if (!Continuation) + return std::nullopt; + for (uint64_t Target : ReusableCalls[I]) + Index->Calls.push_back( + {I, Target - TextAddr, *Continuation, *ReturnRegister}); + } + std::optional> BoundedReturns = collectBoundedSetPcReturns(Decoded, LS, TextAddr, *TextEnd, - DeclaredEntries, FunctionRanges, *Calls, - ExternalEntries); + DeclaredEntries, FunctionRanges, + ExternalEntries, *Index); if (!BoundedReturns) return std::nullopt; + DenseMap BoundedReturnPositions; + for (size_t I = 0; I != BoundedReturns->size(); ++I) + BoundedReturnPositions.try_emplace((*BoundedReturns)[I].InstIndex, I); + + // Canonical one-shot materializations also participate in the reusable + // reaching-value solver so CFG joins can prove their exact path. Preserve + // the established fail-closed entry proof once bounded returns are known: + // an interior alias, fallthrough, or unbounded transfer may still bypass + // the materialization even when its local dataflow token is exact. + BitVector LocallyProvenMaterializedCalls(Decoded.size()); + for (const auto &Entry : Index->MaterializedCalls) { + size_t I = Entry.first; + if (ReusableCalls[I].empty()) + continue; + if (hasKnownControlFlowEntry( + DeclaredEntries, *BoundedReturns, BoundedReturnPositions, *Index, + Entry.second.SequenceStart, Entry.second.SequenceEnd)) { + ReusableCalls[I].clear(); + continue; + } + LocallyProvenMaterializedCalls.set(I); + } DirectControlFlowInfo Info; - for (size_t InstIndex = 0; InstIndex != Decoded.size(); ++InstIndex) { + for (size_t InstIndex : Index->BranchOrCallIndices) { const InternalDecodedInst &DI = Decoded[InstIndex]; - if ((!LS.MIA->isBranch(DI.Inst) && !LS.MIA->isCall(DI.Inst)) || - LS.MIA->isReturn(DI.Inst)) - continue; // Existing indirect branches are handled by // collectIndirectControlFlowFunctions(), which protects their containing // function from source relocation. Calls without a statically resolvable @@ -1588,13 +2586,41 @@ std::optional collectDirectBranchTargets( DI.Inst.getOperand(DI.Inst.getNumOperands() - 1).isImm()) { Target = static_cast( DI.Inst.getOperand(DI.Inst.getNumOperands() - 1).getImm()); - } else if (MaterializedCalls[InstIndex] && - !hasKnownControlFlowEntry( - Decoded, LS, TextAddr, *TextEnd, DeclaredEntries, - *BoundedReturns, - MaterializedCalls[InstIndex]->SequenceStart, - MaterializedCalls[InstIndex]->SequenceEnd)) { - Target = MaterializedCalls[InstIndex]->Target; + } else { + DenseMap::const_iterator Materialized = + Index->MaterializedCalls.find(InstIndex); + if (Materialized != Index->MaterializedCalls.end() && + !hasKnownControlFlowEntry(DeclaredEntries, *BoundedReturns, + BoundedReturnPositions, *Index, + Materialized->second.SequenceStart, + Materialized->second.SequenceEnd)) + Target = Materialized->second.Target; + } + if (!ReusableCalls[InstIndex].empty()) { + if (llvm::any_of(ReusableCalls[InstIndex], + [TextAddr, TextEnd](uint64_t ReusableTarget) { + return ReusableTarget < TextAddr || + ReusableTarget >= *TextEnd; + })) { + log() << "hotswap: unresolved call target at 0x" + << utohexstr(DI.Offset) << " (reusable target outside .text)\n"; + Info.HasUnresolvedTargets = true; + continue; + } + for (uint64_t ReusableTarget : ReusableCalls[InstIndex]) + Info.Targets.insert(ReusableTarget - TextAddr); + Info.BoundedIndirectTransfers.insert(DI.Offset); + if (LocallyProvenMaterializedCalls.test(InstIndex)) { + log() << "hotswap: resolved PC-materialized call at 0x" + << utohexstr(DI.Offset) << " to .text+0x" + << utohexstr(ReusableCalls[InstIndex].front() - TextAddr) + << "\n"; + } else { + log() << "hotswap: resolved reusable PC-materialized call at 0x" + << utohexstr(DI.Offset) << " to " + << ReusableCalls[InstIndex].size() << " target(s)\n"; + } + continue; } if (!Target) { log() << "hotswap: unresolved call target at 0x" << utohexstr(DI.Offset) @@ -1610,6 +2636,8 @@ std::optional collectDirectBranchTargets( log() << "hotswap: resolved PC-materialized call at 0x" << utohexstr(DI.Offset) << " to .text+0x" << utohexstr(RelativeTarget) << "\n"; + if (DI.Inst.getOperand(DI.Inst.getNumOperands() - 1).isReg()) + Info.BoundedIndirectTransfers.insert(DI.Offset); } continue; } @@ -1624,6 +2652,8 @@ std::optional collectDirectBranchTargets( } Info.Targets.insert(*Target); } + for (const BoundedSetPcReturn &Return : *BoundedReturns) + Info.BoundedIndirectTransfers.insert(Decoded[Return.InstIndex].Offset); return Info; } @@ -1644,15 +2674,22 @@ mergeAdjacentLongTrampolines(std::vector &Trampolines, Trampoline &Prev = Merged.back(); std::optional PrevEnd = checkedAddUint64( Prev.OriginalOffset, Prev.OriginalSize, "adjacent trampoline end"); + uint32_t BackReserve = Prev.LongBranchPreservesVcc + ? VccPreservingReturnReserveBytes + : SetPcReturnReserveBytes; + uint32_t BodyPrefix = + Prev.LongBranchPreservesVcc ? VccSaveRestoreBytes : 0; Adjacent = PrevEnd && *PrevEnd == T.OriginalOffset && Prev.Long && T.Long && Prev.UsesSetPCBack && T.UsesSetPCBack && + Prev.LongBranchPreservesVcc == T.LongBranchPreservesVcc && Prev.LongBranchSgprBase == T.LongBranchSgprBase && + Prev.LongBranchUsesVcc == T.LongBranchUsesVcc && Prev.HasFunctionRange && T.HasFunctionRange && Prev.FunctionStart == T.FunctionStart && Prev.FunctionEnd == T.FunctionEnd && !DirectBranchTargets.contains(T.OriginalOffset) && - Prev.Bytes.size() >= SetPcReturnReserveBytes && - T.Bytes.size() >= SetPcReturnReserveBytes; + Prev.Bytes.size() >= BackReserve && + T.Bytes.size() >= BackReserve + BodyPrefix; } if (!Adjacent) { @@ -1666,8 +2703,12 @@ mergeAdjacentLongTrampolines(std::vector &Trampolines, Merged.emplace_back(std::move(T)); continue; } - Prev.Bytes.resize(Prev.Bytes.size() - SetPcReturnReserveBytes); - Prev.Bytes.append(T.Bytes.begin(), T.Bytes.end()); + uint32_t BackReserve = Prev.LongBranchPreservesVcc + ? VccPreservingReturnReserveBytes + : SetPcReturnReserveBytes; + size_t BodyPrefix = Prev.LongBranchPreservesVcc ? VccSaveRestoreBytes : 0; + Prev.Bytes.resize(Prev.Bytes.size() - BackReserve); + Prev.Bytes.append(T.Bytes.begin() + BodyPrefix, T.Bytes.end()); Prev.OriginalSize += T.OriginalSize; ++MergeCount; } @@ -1768,12 +2809,15 @@ collectRelocationProtectedOffsets(ArrayRef Decoded, /// indirect destination, so leave the complete function in place. static DenseSet collectIndirectControlFlowFunctions(ArrayRef Decoded, - const LLVMState &LS, const ElfView &Elf) { + const LLVMState &LS, const ElfView &Elf, + const DenseSet &Bounded) { DenseSet Functions; if (!LS.MIA) return Functions; for (const InternalDecodedInst &DI : Decoded) { + if (Bounded.contains(DI.Offset)) + continue; if (LS.MIA->isBarrier(DI.Inst) || isEndProgram(DI, LS)) continue; if (!LS.MIA->isIndirectBranch(DI.Inst) && @@ -1793,7 +2837,8 @@ collectIndirectControlFlowFunctions(ArrayRef Decoded, /// Grow undersized far-site windows only through proven straight-line code. /// Patched neighbors are merged; ordinary instructions are copied verbatim /// into the trampoline body and retain their original order. This is bounded -/// to the 20 bytes required by the gfx12 SCC-neutral forward sequence. +/// to the source bytes required by the selected gfx12 set-PC sequence and, for +/// a live wave32 VCC, its restore landing pad. static void expandStraightLineTrampolines(PatchContext &Ctx, const DenseSet &DirectBranchTargets) { @@ -1803,15 +2848,20 @@ expandStraightLineTrampolines(PatchContext &Ctx, DenseSet Protected = collectRelocationProtectedOffsets( Ctx.Decoded, Ctx.LS, !Ctx.Config.RunB0A0Patches); DenseSet IndirectControlFlowFunctions = - collectIndirectControlFlowFunctions(Ctx.Decoded, Ctx.LS, Ctx.Elf); + collectIndirectControlFlowFunctions( + Ctx.Decoded, Ctx.LS, Ctx.Elf, + Ctx.DirectControlFlow.BoundedIndirectTransfers); for (size_t I = 0; I != Ctx.OutTrampolines.size(); ++I) { if (Ctx.OutTrampolines[I].HasFunctionRange && IndirectControlFlowFunctions.contains( Ctx.OutTrampolines[I].FunctionStart)) continue; - while (Ctx.OutTrampolines[I].Long && - Ctx.OutTrampolines[I].OriginalSize < SetPcForwardSequenceBytes) { + while (Ctx.OutTrampolines[I].Long && Ctx.OutTrampolines[I].UsesSetPCBack && + Ctx.OutTrampolines[I].OriginalSize < + (Ctx.OutTrampolines[I].LongBranchPreservesVcc + ? VccPreservingReturnReserveBytes + VccLandingPadBytes + : SetPcForwardSequenceBytes)) { Trampoline &T = Ctx.OutTrampolines[I]; std::optional End = checkedAddUint64( T.OriginalOffset, T.OriginalSize, "straight-line expansion end"); @@ -1820,9 +2870,13 @@ expandStraightLineTrampolines(PatchContext &Ctx, if (I + 1 < Ctx.OutTrampolines.size() && Ctx.OutTrampolines[I + 1].OriginalOffset == *End) { + if (T.LongBranchPreservesVcc) + break; Trampoline &Next = Ctx.OutTrampolines[I + 1]; if (!Next.Long || !Next.UsesSetPCBack || Next.LongBranchSgprBase != T.LongBranchSgprBase || + Next.LongBranchUsesVcc != T.LongBranchUsesVcc || + Next.LongBranchPreservesVcc != T.LongBranchPreservesVcc || !T.HasFunctionRange || !Next.HasFunctionRange || T.FunctionStart != Next.FunctionStart || T.FunctionEnd != Next.FunctionEnd || @@ -1844,21 +2898,27 @@ expandStraightLineTrampolines(PatchContext &Ctx, if (!Current) break; const InternalDecodedInst &DI = *Current; + uint32_t BackReserve = T.LongBranchPreservesVcc + ? VccPreservingReturnReserveBytes + : SetPcReturnReserveBytes; std::optional Range = Ctx.Elf.findFunctionTextRangeAtOffset(DI.Offset); if (!Range || !T.HasFunctionRange || Range->Begin != T.FunctionStart || Range->End != T.FunctionEnd || !isSafeStraightLineRelocation(DI, Ctx.LS, Protected) || - T.Bytes.size() < SetPcReturnReserveBytes) + T.Bytes.size() < BackReserve) break; - T.Bytes.insert(T.Bytes.end() - SetPcReturnReserveBytes, - Ctx.Text + DI.Offset, Ctx.Text + DI.Offset + DI.Size); + T.Bytes.insert(T.Bytes.end() - BackReserve, Ctx.Text + DI.Offset, + Ctx.Text + DI.Offset + DI.Size); T.OriginalSize += DI.Size; } - while (Ctx.OutTrampolines[I].Long && - Ctx.OutTrampolines[I].OriginalSize < SetPcForwardSequenceBytes) { + while (Ctx.OutTrampolines[I].Long && Ctx.OutTrampolines[I].UsesSetPCBack && + Ctx.OutTrampolines[I].OriginalSize < + (Ctx.OutTrampolines[I].LongBranchPreservesVcc + ? VccPreservingReturnReserveBytes + VccLandingPadBytes + : SetPcForwardSequenceBytes)) { Trampoline &T = Ctx.OutTrampolines[I]; if (DirectBranchTargets.contains(T.OriginalOffset)) break; @@ -1885,7 +2945,8 @@ expandStraightLineTrampolines(PatchContext &Ctx, if (!Range || !T.HasFunctionRange || Range->Begin != T.FunctionStart || Range->End != T.FunctionEnd) break; - T.Bytes.insert(T.Bytes.begin(), Ctx.Text + DI.Offset, + size_t BodyPrefix = T.LongBranchPreservesVcc ? VccSaveRestoreBytes : 0; + T.Bytes.insert(T.Bytes.begin() + BodyPrefix, Ctx.Text + DI.Offset, Ctx.Text + DI.Offset + DI.Size); T.OriginalOffset = DI.Offset; T.OriginalSize += DI.Size; @@ -1962,7 +3023,8 @@ buildExternalGatewaySleds(ArrayRef Decoded, Expected countReachableSetPcGatewaySlots(ArrayRef Gateways, const LLVMState &LS, uint64_t FromOffset, uint64_t TargetOffset, - unsigned SgprBase, uint64_t MaxSlots) { + unsigned SgprBase, uint64_t MaxSlots, + bool UseVcc, bool PreserveVcc) { uint64_t Slots = 0; for (const NopSled &Sled : Gateways) { if (FromOffset < Sled.FunctionStart || FromOffset >= Sled.FunctionEnd) @@ -1975,10 +3037,10 @@ countReachableSetPcGatewaySlots(ArrayRef Gateways, const LLVMState &LS, if (Distance >= MaxSledDistance || LS.encodeSBranch(FromOffset, Candidate).empty()) break; - std::optional LayoutSize = - getSetPcLongBranchLayoutSize(Candidate, TargetOffset); - if ((SgprBase & 1u) != 0 || !LayoutSize) + getSetPcGatewayLayoutSize(Candidate, TargetOffset, SgprBase, UseVcc, + PreserveVcc); + if (!LayoutSize) return createStringError( Twine("invalid set-PC gateway while counting candidate " "offset 0x") + @@ -2007,6 +3069,7 @@ allocateForwardBranchIslands(std::vector &Gateways, uint64_t Current = FromOffset; while (!isSBranchReachable(Current, TargetOffset)) { + bool Forward = TargetOffset > Current; size_t BestIndex = Gateways.size(); uint64_t BestOffset = 0; for (size_t I = 0; I != Gateways.size(); ++I) { @@ -2015,12 +3078,516 @@ allocateForwardBranchIslands(std::vector &Gateways, FromOffset >= Sled.FunctionEnd) continue; uint64_t UsableEnd = std::min(Sled.End, Sled.FunctionEnd); - if (Sled.WritePos >= TargetOffset || Sled.WritePos <= Current || + bool MakesProgress = + Forward ? Sled.WritePos > Current && Sled.WritePos < TargetOffset + : Sled.WritePos < Current && Sled.WritePos > TargetOffset; + if (!MakesProgress || Sled.WritePos > UsableEnd || + MinInstSize > UsableEnd - Sled.WritePos || + !isSBranchReachable(Current, Sled.WritePos)) + continue; + if (BestIndex == Gateways.size() || + (Forward ? Sled.WritePos > BestOffset : Sled.WritePos < BestOffset)) { + BestIndex = I; + BestOffset = Sled.WritePos; + } + } + + if (BestIndex == Gateways.size()) { + for (size_t I = Allocations.size(); I != 0; --I) { + const Allocation &A = Allocations[I - 1]; + Gateways[A.SledIndex].WritePos = A.PreviousWritePos; + } + return std::nullopt; + } + + NopSled &Best = Gateways[BestIndex]; + Allocations.push_back({BestIndex, Best.WritePos}); + Islands.push_back(Best.WritePos); + Current = Best.WritePos; + Best.WritePos += MinInstSize; + UsedSleds.insert(BestIndex); + } + return Islands; +} + +static SmallVector encodeScc1Branch(const LLVMState &LS, + uint64_t FromOffset, + uint64_t TargetOffset) { + std::optional PcBase = + checkedAddUint64(FromOffset, MinInstSize, "conditional branch PC base"); + if (!PcBase || ((TargetOffset - *PcBase) & (MinInstSize - 1)) != 0) + return {}; + int64_t DwordDelta = + static_cast(TargetOffset - *PcBase) / MinInstSize; + if (DwordDelta < BranchOffsetMin || DwordDelta > BranchOffsetMax) + return {}; + return assembleSingleInst("s_cbranch_scc1 " + std::to_string(DwordDelta), LS); +} + +/// Plan common far gateways before the final pool layout. An 8-byte source +/// cannot hold the 20-byte SCC-neutral set-PC sequence, but it can preserve +/// its identity without touching SCC: +/// +/// s_get_pc_i64 ScratchSource +/// s_branch CommonGateway +/// +/// The common gateway reaches a dispatcher in the pool through a second +/// scratch pair. The dispatcher saves SCC, compares the recorded source PCs, +/// restores SCC in the selected stub, and branches to the matching trampoline +/// body. One 20-byte .text gateway can therefore serve hundreds of otherwise +/// independent 8-byte patch sites. +static bool planSharedDispatchGateways(PatchContext &Ctx, + std::vector &TextGateways) { + struct Candidate { + size_t Index = 0; + unsigned ScratchBase = 0; + }; + SmallVector Candidates; + uint64_t MissingScratchCandidates = 0; + uint64_t FirstMissingScratch = 0; + uint64_t TP = Ctx.PoolBaseOffset; + for (size_t I = 0; I != Ctx.OutTrampolines.size(); ++I) { + Trampoline &T = Ctx.OutTrampolines[I]; + uint64_t ThisTP = TP; + std::optional Next = + checkedAddUint64(TP, T.Bytes.size(), "shared dispatcher pool layout"); + if (!Next) + return false; + TP = *Next; + if (!T.Long || T.OriginalSize < 2 * MinInstSize || + isSBranchReachable(T.OriginalOffset, ThisTP)) + continue; + std::optional> Direct = encodeSetPCLongBranch( + Ctx.LS, T.OriginalOffset, ThisTP, T.LongBranchSgprBase); + if (Direct && Direct->size() <= T.OriginalSize) + continue; + std::optional Scratch = findSafeSgprScratchBlock( + Ctx, T.OriginalOffset, /*Count=*/4, + /*Alignment=*/2, "shared far-dispatch gateway"); + if (Scratch) { + Candidates.push_back({I, Scratch->Base}); + } else { + if (MissingScratchCandidates == 0) + FirstMissingScratch = T.OriginalOffset; + ++MissingScratchCandidates; + } + } + if (MissingScratchCandidates != 0) + log() << "hotswap: shared far-dispatch skipped " << MissingScratchCandidates + << " site(s) without four safe SGPRs" + << " (first at 0x" << utohexstr(FirstMissingScratch) << ")\n"; + + // The ordinary planner is simpler and uses fewer scratch registers for + // small objects. Shared dispatch is a capacity mechanism for dense far-site + // workloads, not a replacement for individual gateways. + if (Candidates.size() < 8) + return true; + + BitVector Assigned(Ctx.OutTrampolines.size()); + SmallVector, 8> Groups; + constexpr size_t MaxGroupSites = 1024; + for (const Candidate &Seed : Candidates) { + if (Assigned.test(Seed.Index)) + continue; + const Trampoline &SeedT = Ctx.OutTrampolines[Seed.Index]; + + size_t SledIndex = TextGateways.size(); + uint64_t BestDistance = std::numeric_limits::max(); + for (size_t I = 0; I != TextGateways.size(); ++I) { + const NopSled &Sled = TextGateways[I]; + uint64_t From = SeedT.OriginalOffset + MinInstSize; + if (SeedT.OriginalOffset < Sled.FunctionStart || + SeedT.OriginalOffset >= Sled.FunctionEnd) + continue; + uint64_t UsableEnd = std::min(Sled.End, Sled.FunctionEnd); + if (Sled.WritePos > UsableEnd || + SetPcForwardSequenceBytes > UsableEnd - Sled.WritePos || + !isSBranchReachable(From, Sled.WritePos)) + continue; + uint64_t Distance = + From > Sled.WritePos ? From - Sled.WritePos : Sled.WritePos - From; + if (Distance < BestDistance) { + BestDistance = Distance; + SledIndex = I; + } + } + uint64_t GatewayOffset = 0; + uint64_t SecondaryGatewayOffset = 0; + std::vector WorkingGateways; + SmallVector SeedIslands; + if (SledIndex != TextGateways.size()) { + GatewayOffset = TextGateways[SledIndex].WritePos; + WorkingGateways = TextGateways; + WorkingGateways[SledIndex].WritePos += SetPcForwardSequenceBytes; + } else { + size_t BestIslandCount = std::numeric_limits::max(); + for (size_t I = 0; I != TextGateways.size(); ++I) { + const NopSled &Sled = TextGateways[I]; + uint64_t UsableEnd = std::min(Sled.End, Sled.FunctionEnd); + if (Sled.WritePos > UsableEnd || + SetPcForwardSequenceBytes > UsableEnd - Sled.WritePos) + continue; + std::vector Trial = TextGateways; + uint64_t TrialGateway = Trial[I].WritePos; + Trial[I].WritePos += SetPcForwardSequenceBytes; + std::optional> Islands = + allocateForwardBranchIslands( + Trial, SeedT.OriginalOffset + MinInstSize, TrialGateway); + if (!Islands || Islands->empty() || Islands->size() >= BestIslandCount) + continue; + BestIslandCount = Islands->size(); + GatewayOffset = TrialGateway; + WorkingGateways = std::move(Trial); + SeedIslands = std::move(*Islands); + } + if (WorkingGateways.empty()) { + // Split the SCC-neutral 20-byte sequence across an 8-byte get-PC + // segment and a 16-byte add/set-PC segment. This admits functions + // that have no single 20-byte padding window. + for (size_t I = 0; I != TextGateways.size() && WorkingGateways.empty(); + ++I) { + const NopSled &First = TextGateways[I]; + uint64_t FirstEnd = std::min(First.End, First.FunctionEnd); + uint64_t SourceBranch = SeedT.OriginalOffset + MinInstSize; + if (SeedT.OriginalOffset < First.FunctionStart || + SeedT.OriginalOffset >= First.FunctionEnd || + First.WritePos > FirstEnd || + 2 * MinInstSize > FirstEnd - First.WritePos || + !isSBranchReachable(SourceBranch, First.WritePos)) + continue; + std::vector FirstReserved = TextGateways; + GatewayOffset = FirstReserved[I].WritePos; + FirstReserved[I].WritePos += 2 * MinInstSize; + for (size_t J = 0; J != FirstReserved.size(); ++J) { + const NopSled &Second = FirstReserved[J]; + uint64_t SecondEnd = std::min(Second.End, Second.FunctionEnd); + if (SeedT.OriginalOffset < Second.FunctionStart || + SeedT.OriginalOffset >= Second.FunctionEnd || + Second.WritePos > SecondEnd || + 4 * MinInstSize > SecondEnd - Second.WritePos || + !isSBranchReachable(GatewayOffset + MinInstSize, + Second.WritePos)) + continue; + WorkingGateways = FirstReserved; + SecondaryGatewayOffset = WorkingGateways[J].WritePos; + WorkingGateways[J].WritePos += 4 * MinInstSize; + break; + } + } + } + if (WorkingGateways.empty()) + continue; + } + + SmallVector Members; + DenseMap> MemberIslands; + DenseMap MemberRelays; + uint64_t GroupBodyBytes = 0; + constexpr uint64_t MaxDispatcherSpan = 120 * 1024; + for (const Candidate &C : Candidates) { + if (Members.size() == MaxGroupSites || Assigned.test(C.Index) || + C.ScratchBase != Seed.ScratchBase) + continue; + const Trampoline &T = Ctx.OutTrampolines[C.Index]; + uint64_t ProposedSpan = + 8 + 28 * (Members.size() + 1) + GroupBodyBytes + T.Bytes.size(); + if (ProposedSpan > MaxDispatcherSpan) + continue; + uint64_t From = T.OriginalOffset + MinInstSize; + SmallVector Islands; + if (C.Index == Seed.Index && !SeedIslands.empty()) { + Islands = SeedIslands; + } else if (!isSBranchReachable(From, GatewayOffset)) { + continue; + } + Members.push_back(C.Index); + GroupBodyBytes += T.Bytes.size(); + if (!Islands.empty()) + MemberIslands[C.Index] = std::move(Islands); + } + DenseSet LocalMembers; + SmallVector, 32> RelayAnchors; + for (size_t Index : Members) { + LocalMembers.insert(Index); + RelayAnchors.push_back( + {Ctx.OutTrampolines[Index].OriginalOffset + MinInstSize, Index}); + } + llvm::sort(RelayAnchors); + for (const Candidate &C : Candidates) { + if (Members.size() == MaxGroupSites || Assigned.test(C.Index) || + LocalMembers.contains(C.Index) || C.ScratchBase != Seed.ScratchBase) + continue; + const Trampoline &T = Ctx.OutTrampolines[C.Index]; + uint64_t ProposedSpan = + 8 + 28 * (Members.size() + 1) + GroupBodyBytes + T.Bytes.size(); + if (ProposedSpan > MaxDispatcherSpan) + continue; + uint64_t From = T.OriginalOffset + MinInstSize; + auto It = + llvm::lower_bound(RelayAnchors, std::make_pair(From, size_t{0})); + std::optional Relay; + if (It != RelayAnchors.end() && isSBranchReachable(From, It->first)) + Relay = It->first; + if (It != RelayAnchors.begin()) { + --It; + if (isSBranchReachable(From, It->first)) + Relay = It->first; + } + if (!Relay) + continue; + Members.push_back(C.Index); + LocalMembers.insert(C.Index); + MemberRelays[C.Index] = *Relay; + GroupBodyBytes += T.Bytes.size(); + RelayAnchors.insert( + llvm::lower_bound(RelayAnchors, std::make_pair(From, C.Index)), + {From, C.Index}); + } + // A single site with a normal 20-byte gateway gains nothing from the + // dispatcher and would unnecessarily consume two extra SGPRs. Leave it to + // the established direct planner. A split 8+16-byte gateway is retained + // because the established planner cannot represent it. + if (Members.size() == 1 && SecondaryGatewayOffset == 0) + continue; + if (Members.empty()) + continue; + TextGateways = std::move(WorkingGateways); + + uint32_t Group = Groups.size() + 1; + for (size_t Index : Members) { + Trampoline &T = Ctx.OutTrampolines[Index]; + SafeSgprScratchBlock Scratch{Seed.ScratchBase, 4}; + if (!commitSafeSgprScratchBlock(Ctx, T.OriginalOffset, Scratch, + "shared far-dispatch gateway")) + return false; + T.UsesSharedDispatcherForward = true; + T.SharedDispatcherGroup = Group; + T.SharedDispatcherSgprBase = Seed.ScratchBase; + T.SharedDispatcherGatewayOffset = GatewayOffset; + DenseMap::const_iterator Relay = + MemberRelays.find(Index); + if (Relay != MemberRelays.end()) + T.SharedDispatcherRelayOffset = Relay->second; + T.SharedDispatcherSecondaryGatewayOffset = SecondaryGatewayOffset; + DenseMap>::iterator Islands = + MemberIslands.find(Index); + if (Islands != MemberIslands.end()) { + T.ForwardBranchIslands = std::move(Islands->second); + T.ForwardBranchTargetOffset = GatewayOffset; + } + Assigned.set(Index); + } + Groups.push_back(std::move(Members)); + } + + if (Groups.empty()) + return true; + + std::vector Reordered; + Reordered.reserve(Ctx.OutTrampolines.size()); + for (size_t I = 0; I != Ctx.OutTrampolines.size(); ++I) + if (!Assigned.test(I)) + Reordered.push_back(std::move(Ctx.OutTrampolines[I])); + for (const SmallVector &Group : Groups) + for (size_t Index : Group) + Reordered.push_back(std::move(Ctx.OutTrampolines[Index])); + Ctx.OutTrampolines = std::move(Reordered); + + for (size_t I = 0; I != Ctx.OutTrampolines.size();) { + Trampoline &First = Ctx.OutTrampolines[I]; + if (!First.UsesSharedDispatcherForward) { + ++I; + continue; + } + uint32_t Group = First.SharedDispatcherGroup; + size_t Count = 0; + while (I + Count != Ctx.OutTrampolines.size() && + Ctx.OutTrampolines[I + Count].SharedDispatcherGroup == Group) + ++Count; + uint64_t Prefix = 8 + 28 * Count; + if (Prefix > std::numeric_limits::max()) + return false; + First.PoolEntryPrefixBytes = static_cast(Prefix); + First.Bytes.insert(First.Bytes.begin(), Prefix, uint8_t{0}); + I += Count; + } + + log() << "hotswap: planned " << Groups.size() + << " shared far-dispatch gateway group(s) for " << Assigned.count() + << " source site(s)\n"; + return true; +} + +static bool emitSharedDispatchers(PatchContext &Ctx) { + DenseMap> Groups; + SmallVector PoolOffsets; + uint64_t TP = Ctx.PoolBaseOffset; + for (size_t I = 0; I != Ctx.OutTrampolines.size(); ++I) { + PoolOffsets.push_back(TP); + Trampoline &T = Ctx.OutTrampolines[I]; + if (T.UsesSharedDispatcherForward) + Groups[T.SharedDispatcherGroup].push_back(I); + std::optional Next = + checkedAddUint64(TP, T.Bytes.size(), "shared dispatcher final layout"); + if (!Next) + return false; + TP = *Next; + } + + for (auto &KV : Groups) { + ArrayRef Members = KV.second; + if (Members.empty()) + continue; + Trampoline &Owner = Ctx.OutTrampolines[Members.front()]; + uint64_t DispatcherOffset = PoolOffsets[Members.front()]; + auto fail = [&](const Twine &Reason) { + log() << "hotswap: error: shared dispatcher group " << KV.first + << " at 0x" << utohexstr(DispatcherOffset) << ": " << Reason + << "\n"; + return false; + }; + unsigned Base = Owner.SharedDispatcherSgprBase; + const std::string SourceLow = "s" + std::to_string(Base); + const std::string CursorLow = "s" + std::to_string(Base + 2); + // After the SCC-neutral gateway has transferred control, only the low + // cursor half is needed: every source and pool address differs by less + // than 4 GiB, so modulo-2^32 deltas remain exact across a load-address + // wrap. Reuse the cursor high half to preserve SCC. + const std::string Save = "s" + std::to_string(Base + 3); + + Owner.HasForwardGateway = true; + Owner.ForwardGatewayOffset = Owner.SharedDispatcherGatewayOffset; + if (Owner.SharedDispatcherSecondaryGatewayOffset == 0) { + std::optional> Gateway = + encodeSetPCLongBranch(Ctx.LS, Owner.SharedDispatcherGatewayOffset, + DispatcherOffset, Base + 2); + if (!Gateway || Gateway->size() > SetPcForwardSequenceBytes) + return fail("single-segment gateway encoding failed"); + Owner.ForwardGatewayBytes = std::move(*Gateway); + } else { + const std::string GatewayPair = "s[" + std::to_string(Base + 2) + ":" + + std::to_string(Base + 3) + "]"; + Owner.ForwardGatewayBytes = + assembleSingleInst("s_get_pc_i64 " + GatewayPair, Ctx.LS); + SmallVector ToSecond = + Ctx.LS.encodeSBranch(Owner.SharedDispatcherGatewayOffset + + Owner.ForwardGatewayBytes.size(), + Owner.SharedDispatcherSecondaryGatewayOffset); + if (Owner.ForwardGatewayBytes.size() != MinInstSize || + ToSecond.size() != MinInstSize) + return fail("split gateway first segment encoding failed"); + Owner.ForwardGatewayBytes.append(ToSecond); + + uint64_t PcBase = Owner.SharedDispatcherGatewayOffset + MinInstSize; + uint64_t Delta = DispatcherOffset - PcBase; + SmallVector Lines; + Lines.push_back("s_add_nc_u64 " + GatewayPair + ", " + GatewayPair + + ", 0x" + utohexstr(Delta)); + Lines.push_back("s_set_pc_i64 " + GatewayPair); + Owner.SecondaryForwardGatewayBytes = + assembleInstructions(joinAsmLines(Lines), Ctx.LS); + if (Owner.SecondaryForwardGatewayBytes.empty() || + Owner.SecondaryForwardGatewayBytes.size() > 4 * MinInstSize) + return fail("split gateway second segment encoding failed"); + while (Owner.SecondaryForwardGatewayBytes.size() < 4 * MinInstSize) + Owner.SecondaryForwardGatewayBytes.append(Ctx.LS.SNopBytes); + } + + SmallVector BodyOffsets; + for (size_t Member : Members) + BodyOffsets.push_back(PoolOffsets[Member] + + Ctx.OutTrampolines[Member].PoolEntryPrefixBytes); + + SmallVector Bytes; + auto appendInst = [&](StringRef Asm) { + SmallVector Encoded = assembleSingleInst(Asm, Ctx.LS); + if (Encoded.empty()) + return false; + Bytes.append(Encoded); + return true; + }; + if (!appendInst("s_cselect_b32 " + Save + ", 1, 0")) + return fail("SCC save encoding failed"); + + uint64_t CursorValue = DispatcherOffset; + uint64_t StubBase = + DispatcherOffset + 4 + 20 * Members.size() + MinInstSize; + for (size_t J = 0; J != Members.size(); ++J) { + const Trampoline &T = Ctx.OutTrampolines[Members[J]]; + uint64_t SourcePc = T.OriginalOffset + MinInstSize; + uint64_t Distance = SourcePc > CursorValue ? SourcePc - CursorValue + : CursorValue - SourcePc; + if (Distance >= (uint64_t{1} << 32)) + return fail("source-to-dispatcher span exceeds 32 bits"); + uint64_t Delta = SourcePc - CursorValue; + SmallVector Add = assembleSingleInst( + "s_add_co_u32 " + CursorLow + ", " + CursorLow + ", 0x" + + utohexstr(static_cast(Delta)), + Ctx.LS); + if (Add.empty() || Add.size() > 3 * MinInstSize) + return fail("source-PC cursor add encoding failed"); + Bytes.append(Add); + while (Add.size() < 3 * MinInstSize) { + Bytes.append(Ctx.LS.SNopBytes); + Add.append(Ctx.LS.SNopBytes); + } + if (!appendInst("s_cmp_eq_u32 " + SourceLow + ", " + CursorLow)) + return fail("source-PC compare encoding failed"); + uint64_t BranchFrom = DispatcherOffset + Bytes.size(); + SmallVector Branch = + encodeScc1Branch(Ctx.LS, BranchFrom, StubBase + J * 2 * MinInstSize); + if (Branch.size() != MinInstSize) + return fail("source-PC conditional branch is out of range"); + Bytes.append(Branch); + CursorValue = SourcePc; + } + if (!appendInst("s_trap 2")) + return fail("unmatched-source trap encoding failed"); + for (size_t J = 0; J != Members.size(); ++J) { + if (!appendInst("s_cmp_lg_u32 " + Save + ", 0")) + return fail("SCC restore encoding failed"); + uint64_t BranchFrom = DispatcherOffset + Bytes.size(); + SmallVector Branch = + Ctx.LS.encodeSBranch(BranchFrom, BodyOffsets[J]); + if (Branch.size() != MinInstSize) + return fail("selected trampoline body is out of branch range"); + Bytes.append(Branch); + } + if (Bytes.size() != Owner.PoolEntryPrefixBytes) + return fail("dispatcher size differs from reserved prefix"); + std::memcpy(Owner.Bytes.data(), Bytes.data(), Bytes.size()); + } + return true; +} + +static std::optional> +allocateBackwardBranchIslands(std::vector &Gateways, + uint64_t OwnerOffset, uint64_t FromOffset, + uint64_t TargetOffset) { + struct Allocation { + size_t SledIndex = 0; + uint64_t PreviousWritePos = 0; + }; + SmallVector Allocations; + SmallVector Islands; + DenseSet UsedSleds; + uint64_t Current = FromOffset; + + while (!isSBranchReachable(Current, TargetOffset)) { + size_t BestIndex = Gateways.size(); + uint64_t BestOffset = std::numeric_limits::max(); + for (size_t I = 0; I != Gateways.size(); ++I) { + NopSled &Sled = Gateways[I]; + if (UsedSleds.contains(I) || OwnerOffset < Sled.FunctionStart || + OwnerOffset >= Sled.FunctionEnd) + continue; + uint64_t UsableEnd = std::min(Sled.End, Sled.FunctionEnd); + if (Sled.WritePos <= TargetOffset || Sled.WritePos >= Current || Sled.WritePos > UsableEnd || MinInstSize > UsableEnd - Sled.WritePos || !isSBranchReachable(Current, Sled.WritePos)) continue; - if (BestIndex == Gateways.size() || Sled.WritePos > BestOffset) { + if (BestIndex == Gateways.size() || Sled.WritePos < BestOffset) { BestIndex = I; BestOffset = Sled.WritePos; } @@ -2055,9 +3622,13 @@ assignLongBranchGateways(PatchContext &Ctx, DirectBranchTargets); for (const NopSled &Sled : Ctx.NopSleds) Gateways.push_back(Sled); + if (!planSharedDispatchGateways(Ctx, Gateways) || + !emitSharedDispatchers(Ctx)) + return false; } DenseMap PoolIslandOwners; + DenseMap SourceTailIslandOwners; uint64_t IslandLayoutOffset = Ctx.PoolBaseOffset; for (size_t I = 0; I != Ctx.OutTrampolines.size(); ++I) { Trampoline &T = Ctx.OutTrampolines[I]; @@ -2082,6 +3653,7 @@ assignLongBranchGateways(PatchContext &Ctx, uint64_t InitialCandidateSlots = 0; }; std::vector Pending; + uint64_t ReturnBranchIslandChains = 0; uint64_t TrampOffset = Ctx.PoolBaseOffset; for (size_t I = 0; I != Ctx.OutTrampolines.size(); ++I) { Trampoline &T = Ctx.OutTrampolines[I]; @@ -2093,25 +3665,91 @@ assignLongBranchGateways(PatchContext &Ctx, TrampOffset = *Next; if (!T.Long) continue; + if (T.UsesSharedDispatcherForward) + continue; + + if (!T.UsesSetPCBack) { + const uint64_t TrailingIsland = + T.HasPoolBranchIsland ? PoolBranchIslandBytes : 0; + if (T.Bytes.size() < TrailingIsland + MinInstSize) { + log() << "hotswap: error: registerless return reservation is " + "truncated at 0x" + << utohexstr(T.OriginalOffset) << "\n"; + return false; + } + uint64_t BackSlot = *Next - TrailingIsland - MinInstSize; + std::optional ReturnTo = + checkedAddUint64(T.OriginalOffset, T.OriginalSize, + "registerless trampoline return target"); + if (!ReturnTo) + return false; + std::optional> ReturnIslands = + allocateBackwardBranchIslands(Gateways, T.OriginalOffset, BackSlot, + *ReturnTo); + if (!ReturnIslands) { + log() << "hotswap: error: no safe return s_branch island chain for " + "far site 0x" + << utohexstr(T.OriginalOffset) << "\n"; + return false; + } + T.ReturnBranchIslands = std::move(*ReturnIslands); + T.ReturnBranchTargetOffset = *ReturnTo; + ReturnBranchIslandChains += !T.ReturnBranchIslands.empty(); + } if (isSBranchReachable(T.OriginalOffset, TP)) { T.UsesShortBranchForward = true; + if (T.LongBranchPreservesVcc) + std::memcpy(T.Bytes.data(), Ctx.LS.SNopBytes.data(), MinInstSize); continue; } - std::optional> Direct = encodeSetPCLongBranch( - Ctx.LS, T.OriginalOffset, TP, T.LongBranchSgprBase); - if (Direct && Direct->size() <= T.OriginalSize) { - T.UsesDirectSetPCForward = true; - T.DirectSetPCForwardBytes = std::move(*Direct); - continue; + if (T.UsesSetPCBack) { + std::optional> Direct = + T.LongBranchPreservesVcc + ? encodeSetPcGateway(Ctx.LS, T.OriginalOffset, TP, + T.LongBranchSgprBase, T.LongBranchUsesVcc, + /*PreserveVcc=*/true) + : encodeSetPCLongBranch(Ctx.LS, T.OriginalOffset, TP, + T.LongBranchSgprBase, + T.LongBranchUsesVcc); + uint64_t RequiredSourceBytes = + Direct ? Direct->size() + + (T.LongBranchPreservesVcc ? VccLandingPadBytes : 0) + : 0; + if (Direct && RequiredSourceBytes <= T.OriginalSize) { + T.UsesDirectSetPCForward = true; + T.DirectSetPCForwardBytes = std::move(*Direct); + continue; + } } Pending.push_back({I, TP, 0}); } + + // Once a source is replaced by a one-dword branch, the remainder of its + // original instruction window is unreachable and can provide a safe relay. + // Add these only after selecting direct set-PC sources, whose longer forward + // sequence consumes the tail. Shared dispatch and VCC preservation likewise + // reserve the second dword. Relays are object-wide: unlike an arbitrary NOP + // sled they cannot be reached by the owning function's original fallthrough. + for (size_t I = 0; I != Ctx.OutTrampolines.size(); ++I) { + const Trampoline &T = Ctx.OutTrampolines[I]; + if (T.OriginalSize < 2 * MinInstSize || T.UsesDirectSetPCForward || + T.UsesSharedDispatcherForward || T.LongBranchPreservesVcc) + continue; + uint64_t Tail = T.OriginalOffset + MinInstSize; + SourceTailIslandOwners[Tail] = I; + Gateways.push_back({Tail, Tail + MinInstSize, Tail, 0, + std::numeric_limits::max()}); + } + for (PendingGateway &P : Pending) { const Trampoline &T = Ctx.OutTrampolines[P.TrampolineIndex]; + if (!T.UsesSetPCBack) + continue; Expected CandidateSlots = countReachableSetPcGatewaySlots( Gateways, Ctx.LS, T.OriginalOffset, P.TargetOffset, - T.LongBranchSgprBase, Pending.size()); + T.LongBranchSgprBase, Pending.size(), T.LongBranchUsesVcc, + T.LongBranchPreservesVcc); if (!CandidateSlots) { log() << "hotswap: error: failed to count gateways for far site 0x" << utohexstr(T.OriginalOffset) << ": " @@ -2121,36 +3759,25 @@ assignLongBranchGateways(PatchContext &Ctx, P.InitialCandidateSlots = *CandidateSlots; } - std::vector StillPending; - StillPending.reserve(Pending.size()); - uint64_t BranchIslandChains = 0; - for (const PendingGateway &P : Pending) { - Trampoline &T = Ctx.OutTrampolines[P.TrampolineIndex]; - std::optional> Islands = - allocateForwardBranchIslands(Gateways, T.OriginalOffset, - P.TargetOffset); - if (!Islands || Islands->empty()) { - StillPending.push_back(P); - continue; - } - T.ForwardBranchIslands = std::move(*Islands); - T.ForwardBranchTargetOffset = P.TargetOffset; - ++BranchIslandChains; - } - Pending = std::move(StillPending); - std::stable_sort(Pending.begin(), Pending.end(), [](const PendingGateway &LHS, const PendingGateway &RHS) { return LHS.InitialCandidateSlots < RHS.InitialCandidateSlots; }); + std::vector StillPending; + StillPending.reserve(Pending.size()); uint64_t AssignedGateways = 0; for (const PendingGateway &P : Pending) { Trampoline &T = Ctx.OutTrampolines[P.TrampolineIndex]; + if (!T.UsesSetPCBack) { + StillPending.push_back(P); + continue; + } Expected> GatewayOrErr = findNearestSetPcGateway(Gateways, Ctx.LS, T.OriginalOffset, - P.TargetOffset, T.LongBranchSgprBase); + P.TargetOffset, T.LongBranchSgprBase, + T.LongBranchUsesVcc, T.LongBranchPreservesVcc); if (!GatewayOrErr) { log() << "hotswap: error: failed to plan gateway for far site 0x" << utohexstr(T.OriginalOffset) << ": " @@ -2159,10 +3786,8 @@ assignLongBranchGateways(PatchContext &Ctx, } std::optional Gateway = std::move(*GatewayOrErr); if (!Gateway) { - log() << "hotswap: error: no safe short-branch gateway for far site 0x" - << utohexstr(T.OriginalOffset) << " (" << P.InitialCandidateSlots - << " initial candidate slot(s))\n"; - return false; + StillPending.push_back(P); + continue; } T.HasForwardGateway = true; T.ForwardGatewayOffset = Gateway->Sled->WritePos; @@ -2170,12 +3795,50 @@ assignLongBranchGateways(PatchContext &Ctx, Gateway->Sled->WritePos += T.ForwardGatewayBytes.size(); ++AssignedGateways; } - if (!Pending.empty()) + Pending = std::move(StillPending); + + uint64_t BranchIslandChains = 0; + StillPending.clear(); + StillPending.reserve(Pending.size()); + for (const PendingGateway &P : Pending) { + Trampoline &T = Ctx.OutTrampolines[P.TrampolineIndex]; + std::optional> Islands = + allocateForwardBranchIslands(Gateways, T.OriginalOffset, + P.TargetOffset); + if (!Islands || Islands->empty()) { + StillPending.push_back(P); + continue; + } + T.ForwardBranchIslands = std::move(*Islands); + T.ForwardBranchTargetOffset = P.TargetOffset; + if (T.LongBranchPreservesVcc) + std::memcpy(T.Bytes.data(), Ctx.LS.SNopBytes.data(), MinInstSize); + ++BranchIslandChains; + } + Pending = std::move(StillPending); + + if (!Pending.empty()) { + const PendingGateway &P = Pending.front(); + const Trampoline &T = Ctx.OutTrampolines[P.TrampolineIndex]; + if (!T.UsesSetPCBack) + log() << "hotswap: error: no safe forward s_branch island chain for " + "registerless far site 0x" + << utohexstr(T.OriginalOffset) << "\n"; + else + log() << "hotswap: error: no safe short-branch gateway for far site 0x" + << utohexstr(T.OriginalOffset) << " (" << P.InitialCandidateSlots + << " initial candidate slot(s))\n"; + return false; + } + if (AssignedGateways != 0) log() << "hotswap: assigned " << AssignedGateways << " SCC-neutral forward gateway(s)\n"; if (BranchIslandChains != 0) log() << "hotswap: assigned " << BranchIslandChains << " forward s_branch island chain(s)\n"; + if (ReturnBranchIslandChains != 0) + log() << "hotswap: assigned " << ReturnBranchIslandChains + << " return s_branch island chain(s)\n"; for (Trampoline &T : Ctx.OutTrampolines) { if (T.HasForwardGateway) { @@ -2188,6 +3851,17 @@ assignLongBranchGateways(PatchContext &Ctx, } std::memcpy(Ctx.Text + T.ForwardGatewayOffset, T.ForwardGatewayBytes.data(), T.ForwardGatewayBytes.size()); + if (!T.SecondaryForwardGatewayBytes.empty()) { + uint64_t Offset = T.SharedDispatcherSecondaryGatewayOffset; + if (Offset > Ctx.TextSize || + T.SecondaryForwardGatewayBytes.size() > Ctx.TextSize - Offset) { + log() << "hotswap: error: secondary forward gateway at 0x" + << utohexstr(Offset) << " extends past .text.\n"; + return false; + } + std::memcpy(Ctx.Text + Offset, T.SecondaryForwardGatewayBytes.data(), + T.SecondaryForwardGatewayBytes.size()); + } } for (size_t I = 0; I != T.ForwardBranchIslands.size(); ++I) { uint64_t From = T.ForwardBranchIslands[I]; @@ -2203,7 +3877,14 @@ assignLongBranchGateways(PatchContext &Ctx, } DenseMap::const_iterator Owner = PoolIslandOwners.find(From); - if (Owner != PoolIslandOwners.end()) { + DenseMap::const_iterator SourceOwner = + SourceTailIslandOwners.find(From); + if (SourceOwner != SourceTailIslandOwners.end()) { + Trampoline &OwnerT = Ctx.OutTrampolines[SourceOwner->second]; + OwnerT.HasSourceTailBranchIsland = true; + OwnerT.SourceTailBranchIslandOffset = From; + OwnerT.SourceTailBranchTargetOffset = To; + } else if (Owner != PoolIslandOwners.end()) { Trampoline &OwnerT = Ctx.OutTrampolines[Owner->second]; std::memcpy(OwnerT.Bytes.data() + OwnerT.Bytes.size() - PoolBranchIslandBytes, @@ -2217,6 +3898,40 @@ assignLongBranchGateways(PatchContext &Ctx, std::memcpy(Ctx.Text + From, Branch.data(), Branch.size()); } } + for (size_t I = 0; I != T.ReturnBranchIslands.size(); ++I) { + uint64_t From = T.ReturnBranchIslands[I]; + uint64_t To = I + 1 == T.ReturnBranchIslands.size() + ? T.ReturnBranchTargetOffset + : T.ReturnBranchIslands[I + 1]; + SmallVector Branch = Ctx.LS.encodeSBranch(From, To); + if (Branch.size() != MinInstSize) { + log() << "hotswap: error: failed to encode return branch island at 0x" + << utohexstr(From) << "\n"; + return false; + } + DenseMap::const_iterator Owner = + PoolIslandOwners.find(From); + DenseMap::const_iterator SourceOwner = + SourceTailIslandOwners.find(From); + if (SourceOwner != SourceTailIslandOwners.end()) { + Trampoline &OwnerT = Ctx.OutTrampolines[SourceOwner->second]; + OwnerT.HasSourceTailBranchIsland = true; + OwnerT.SourceTailBranchIslandOffset = From; + OwnerT.SourceTailBranchTargetOffset = To; + } else if (Owner != PoolIslandOwners.end()) { + Trampoline &OwnerT = Ctx.OutTrampolines[Owner->second]; + std::memcpy(OwnerT.Bytes.data() + OwnerT.Bytes.size() - + PoolBranchIslandBytes, + Branch.data(), Branch.size()); + } else { + if (From > Ctx.TextSize || Branch.size() > Ctx.TextSize - From) { + log() << "hotswap: error: return branch island at 0x" + << utohexstr(From) << " is outside .text and trampoline pool\n"; + return false; + } + std::memcpy(Ctx.Text + From, Branch.data(), Branch.size()); + } + } } return true; } @@ -2305,7 +4020,8 @@ static std::optional applyGfx1250B0toA0Rules( Elf.functionTextRanges(); std::optional ControlFlow = collectDirectBranchTargets( Decoded, LS, Elf.textAddr(), Elf.textSize(), DeclaredEntries->Entries, - FunctionRanges, DeclaredEntries->ExternalEntries); + FunctionRanges, DeclaredEntries->ExternalEntries, + ArrayRef(Text, TextSize)); if (!ControlFlow) return std::nullopt; if (ControlFlow->HasUnresolvedTargets) { @@ -2583,13 +4299,10 @@ fixupTrampolineBranches(std::vector &Trampolines, uint8_t *Text, return false; TrampOffset = *NextTrampOffset; - if (T.Long && !T.UsesSetPCBack) { - log() << "hotswap: error: far trampoline lacks safe set-PC return at 0x" - << utohexstr(T.OriginalOffset) << "\n"; - return false; - } const uint32_t BackReserve = - T.UsesSetPCBack ? SetPcReturnReserveBytes : MinInstSize; + T.LongBranchPreservesVcc + ? VccPreservingReturnReserveBytes + : (T.UsesSetPCBack ? SetPcReturnReserveBytes : MinInstSize); const uint32_t TrailingIsland = T.HasPoolBranchIsland ? PoolBranchIslandBytes : 0; if (T.Bytes.size() < BackReserve + TrailingIsland) { @@ -2606,11 +4319,35 @@ fixupTrampolineBranches(std::vector &Trampolines, uint8_t *Text, return false; std::optional> BrBack; - if (T.UsesSetPCBack) { - BrBack = - encodeSetPCLongBranch(LS, BackSlot, *ReturnTo, T.LongBranchSgprBase); + if (T.LongBranchPreservesVcc) { + SmallVector Save = assembleSingleInst( + "s_mov_b32 s" + std::to_string(T.LongBranchSgprBase) + ", vcc_lo", + LS); + std::optional SetPcOffset = checkedAddUint64( + BackSlot, Save.size(), "VCC-preserving return set-PC offset"); + uint64_t LandingDisplacement = T.UsesDirectSetPCForward + ? T.DirectSetPCForwardBytes.size() + : MinInstSize; + std::optional Landing = + checkedAddUint64(T.OriginalOffset, LandingDisplacement, + "VCC-preserving return landing offset"); + if (Save.size() != VccSaveRestoreBytes || !SetPcOffset || !Landing) + return false; + std::optional> SetPc = encodeSetPCLongBranch( + LS, *SetPcOffset, *Landing, T.LongBranchSgprBase, /*UseVcc=*/true); + if (SetPc) { + Save.append(SetPc->begin(), SetPc->end()); + BrBack = std::move(Save); + } + } else if (T.UsesSetPCBack) { + BrBack = encodeSetPCLongBranch(LS, BackSlot, *ReturnTo, + T.LongBranchSgprBase, T.LongBranchUsesVcc); } else { - SmallVector ShortBranch = LS.encodeSBranch(BackSlot, *ReturnTo); + uint64_t BranchTarget = T.ReturnBranchIslands.empty() + ? *ReturnTo + : T.ReturnBranchIslands.front(); + SmallVector ShortBranch = + LS.encodeSBranch(BackSlot, BranchTarget); if (!ShortBranch.empty()) BrBack = std::move(ShortBranch); } @@ -2627,7 +4364,23 @@ fixupTrampolineBranches(std::vector &Trampolines, uint8_t *Text, SmallVector BrFwd; if (T.Long) { - if (T.UsesShortBranchForward) { + if (T.UsesSharedDispatcherForward) { + const std::string Pair = + "s[" + std::to_string(T.SharedDispatcherSgprBase) + ":" + + std::to_string(T.SharedDispatcherSgprBase + 1) + "]"; + BrFwd = assembleSingleInst("s_get_pc_i64 " + Pair, LS); + if (BrFwd.size() != MinInstSize) + return false; + uint64_t BranchTarget = + T.SharedDispatcherRelayOffset ? T.SharedDispatcherRelayOffset + : T.ForwardBranchIslands.empty() ? T.SharedDispatcherGatewayOffset + : T.ForwardBranchIslands.front(); + SmallVector Branch = + LS.encodeSBranch(T.OriginalOffset + BrFwd.size(), BranchTarget); + if (Branch.size() != MinInstSize) + return false; + BrFwd.append(Branch); + } else if (T.UsesShortBranchForward) { BrFwd = LS.encodeSBranch(T.OriginalOffset, TP); } else if (!T.ForwardBranchIslands.empty()) { BrFwd = @@ -2650,11 +4403,55 @@ fixupTrampolineBranches(std::vector &Trampolines, uint8_t *Text, return false; } std::memcpy(Text + T.OriginalOffset, BrFwd.data(), BrFwd.size()); + uint32_t PadStart = BrFwd.size(); + if (T.LongBranchPreservesVcc) { + uint64_t LandingDisplacement = + T.UsesDirectSetPCForward ? BrFwd.size() : MinInstSize; + if ((!T.UsesDirectSetPCForward && BrFwd.size() != MinInstSize) || + LandingDisplacement > T.OriginalSize || + VccLandingPadBytes > T.OriginalSize - LandingDisplacement) { + log() << "hotswap: error: VCC-preserving source window is invalid at " + "0x" + << utohexstr(T.OriginalOffset) << "\n"; + return false; + } + SmallVector Restore = assembleSingleInst( + "s_mov_b32 vcc_lo, s" + std::to_string(T.LongBranchSgprBase), LS); + if (Restore.size() != VccSaveRestoreBytes) { + log() << "hotswap: error: failed to encode VCC restore landing at 0x" + << utohexstr(T.OriginalOffset + LandingDisplacement) << "\n"; + return false; + } + std::memcpy(Text + T.OriginalOffset + LandingDisplacement, Restore.data(), + Restore.size()); + PadStart = LandingDisplacement + VccLandingPadBytes; + } // Pad the tail of the replaced slot with cached s_nop bytes. - for (uint32_t I = BrFwd.size(); I + MinInstSize <= T.OriginalSize; + for (uint32_t I = PadStart; I + MinInstSize <= T.OriginalSize; I += MinInstSize) std::memcpy(Text + T.OriginalOffset + I, LS.SNopBytes.data(), MinInstSize); + if (T.HasSourceTailBranchIsland) { + if (T.SourceTailBranchIslandOffset < T.OriginalOffset || + T.SourceTailBranchIslandOffset - T.OriginalOffset < PadStart || + T.SourceTailBranchIslandOffset - T.OriginalOffset > + T.OriginalSize - MinInstSize) { + log() << "hotswap: error: source-tail branch island overlaps the " + "forward sequence at 0x" + << utohexstr(T.OriginalOffset) << "\n"; + return false; + } + SmallVector Relay = LS.encodeSBranch( + T.SourceTailBranchIslandOffset, T.SourceTailBranchTargetOffset); + if (Relay.size() != MinInstSize) { + log() << "hotswap: error: source-tail branch island encoding failed " + "at 0x" + << utohexstr(T.SourceTailBranchIslandOffset) << "\n"; + return false; + } + std::memcpy(Text + T.SourceTailBranchIslandOffset, Relay.data(), + Relay.size()); + } } return true; } @@ -2754,12 +4551,22 @@ static amd_comgr_status_t retargetCodeObjectImpl( << "parseable ELF64 (" << toString(ViewOrErr.takeError()) << ").\n"; return AMD_COMGR_STATUS_ERROR_INVALID_ARGUMENT; } + ElfView &Elf = *ViewOrErr; if (ViewOrErr->textSize() == 0) { - log() << "hotswap: error: retargetCodeObject: input ELF has empty " - << ".text section; nothing to rewrite.\n"; - return AMD_COMGR_STATUS_ERROR_INVALID_ARGUMENT; + if (!Elf.isValidDataOnlyObject()) { + log() << "hotswap: error: retargetCodeObject: empty .text does not " + "describe a valid data-only code object.\n"; + return AMD_COMGR_STATUS_ERROR_INVALID_ARGUMENT; + } + std::unique_ptr Result = + copyOutputBuffer(ElfData, ElfSize, "data-only"); + if (!Result) + return AMD_COMGR_STATUS_ERROR_OUT_OF_RESOURCES; + Out = std::move(Result); + log() << "hotswap: accepted data-only code object with empty .text; " + "returning a byte-identical copy.\n"; + return AMD_COMGR_STATUS_SUCCESS; } - ElfView &Elf = *ViewOrErr; if (Prof) Profile.add(HotswapMetric::ElfParse, profNowNs() - ParseT0, 0); diff --git a/amd/comgr/src/comgr-hotswap-elf.cpp b/amd/comgr/src/comgr-hotswap-elf.cpp index a14e9aef10b4a..692d4b69a9e36 100644 --- a/amd/comgr/src/comgr-hotswap-elf.cpp +++ b/amd/comgr/src/comgr-hotswap-elf.cpp @@ -203,7 +203,6 @@ static bool rewriteMetadataNotes(uint8_t *Elf, const ELFFileT &File, } PendingWrites.push_back({DescOffset, std::move(NewBlob)}); } - if (Err) { log() << "hotswap: error: " << Context << ": failed to iterate AMDGPU notes: " << toString(std::move(Err)) @@ -506,6 +505,163 @@ std::vector ElfView::functionTextRanges() const { return std::vector(Ranges.begin(), Ranges.end()); } +bool ElfView::isValidDataOnlyObject() const { + if (textSize() != 0) { + log() << "hotswap: error: data-only validation requires an empty .text " + "section.\n"; + return false; + } + + for (const ELFT::Shdr &Shdr : Sections) { + if ((Shdr.sh_flags & ELF::SHF_EXECINSTR) != 0 && Shdr.sh_size != 0) { + Expected NameOrErr = File.getSectionName(Shdr); + if (!NameOrErr) { + log() << "hotswap: error: data-only validation found a non-empty " + "executable section with an unreadable name: " + << toString(NameOrErr.takeError()) << "\n"; + return false; + } + log() << "hotswap: error: data-only object has non-empty executable " + "section '" + << *NameOrErr << "'.\n"; + return false; + } + + if (Shdr.sh_type != ELF::SHT_SYMTAB && Shdr.sh_type != ELF::SHT_DYNSYM) + continue; + + Expected SymsOrErr = File.symbols(&Shdr); + if (!SymsOrErr) { + log() << "hotswap: error: data-only validation failed to read symbols: " + << toString(SymsOrErr.takeError()) << "\n"; + return false; + } + Expected StrTabOrErr = + File.getStringTableForSymtab(Shdr, Sections); + if (!StrTabOrErr) { + log() << "hotswap: error: data-only validation failed to read the " + "symbol string table: " + << toString(StrTabOrErr.takeError()) << "\n"; + return false; + } + + for (const ELFT::Sym &Sym : *SymsOrErr) { + Expected NameOrErr = Sym.getName(*StrTabOrErr); + if (!NameOrErr) { + log() << "hotswap: error: data-only validation found an unreadable " + "symbol name: " + << toString(NameOrErr.takeError()) << "\n"; + return false; + } + if ((Sym.getType() == ELF::STT_FUNC || + Sym.getType() == ELF::STT_GNU_IFUNC) && + Sym.st_shndx != ELF::SHN_UNDEF) { + if (Sym.st_shndx == TextSectionIndex) + log() << "hotswap: error: data-only object has a function/ifunc " + "symbol in empty .text.\n"; + else + log() << "hotswap: error: data-only object has defined " + "function/ifunc symbol '" + << *NameOrErr << "'.\n"; + return false; + } + if (NameOrErr->ends_with(".kd")) { + log() << "hotswap: error: data-only object has kernel descriptor " + "symbol '" + << *NameOrErr << "'.\n"; + return false; + } + } + } + + bool SawMetadataNote = false; + auto validateMetadataNote = [&](ELFT::Note Note) { + if (Note.getName() != "AMDGPU" || Note.getType() != ELF::NT_AMDGPU_METADATA) + return true; + SawMetadataNote = true; + + ArrayRef Desc = Note.getDesc(4); + if (Desc.empty()) { + log() << "hotswap: error: data-only AMDGPU metadata note has an " + "empty descriptor.\n"; + return false; + } + StringRef Blob(reinterpret_cast(Desc.data()), Desc.size()); + msgpack::Document Doc; + if (!Doc.readFromBlob(Blob, false)) { + log() << "hotswap: error: failed to parse data-only AMDGPU metadata " + "note.\n"; + return false; + } + + msgpack::DocNode Root = Doc.getRoot(); + if (!Root.isMap()) { + log() << "hotswap: error: data-only AMDGPU metadata root is not a map.\n"; + return false; + } + msgpack::MapDocNode &RootMap = Root.getMap(); + msgpack::DocNode::MapTy::iterator KernelsIt = + RootMap.find("amdhsa.kernels"); + if (KernelsIt == RootMap.end() || !KernelsIt->second.isArray()) { + log() << "hotswap: error: data-only AMDGPU metadata has no valid " + "amdhsa.kernels array.\n"; + return false; + } + if (!KernelsIt->second.getArray().empty()) { + log() << "hotswap: error: data-only AMDGPU metadata claims " + << KernelsIt->second.getArray().size() << " kernel(s).\n"; + return false; + } + return true; + }; + + Expected PhdrsOrErr = File.program_headers(); + if (!PhdrsOrErr) { + log() << "hotswap: error: data-only validation failed to read program " + "headers: " + << toString(PhdrsOrErr.takeError()) << "\n"; + return false; + } + for (const ELFT::Phdr &Phdr : *PhdrsOrErr) { + if (Phdr.p_type != ELF::PT_NOTE) + continue; + + Error Err = Error::success(); + for (ELFT::Note Note : File.notes(Phdr, Err)) + if (!validateMetadataNote(Note)) + return false; + if (Err) { + log() << "hotswap: error: data-only validation failed to iterate AMDGPU " + "notes: " + << toString(std::move(Err)) << "\n"; + return false; + } + } + + // Linked code objects normally expose metadata through PT_NOTE. Match + // COMGR's metadata lookup fallback for relocatable/unusual objects whose + // note exists only in the section table. + if (!SawMetadataNote) { + for (const ELFT::Shdr &Shdr : Sections) { + if (Shdr.sh_type != ELF::SHT_NOTE) + continue; + + Error Err = Error::success(); + for (ELFT::Note Note : File.notes(Shdr, Err)) + if (!validateMetadataNote(Note)) + return false; + if (Err) { + log() + << "hotswap: error: data-only validation failed to iterate AMDGPU " + "note sections: " + << toString(std::move(Err)) << "\n"; + return false; + } + } + } + return true; +} + // -- ElfView::findKernelAtAddress --------------------------------------------- const ElfView::FunctionTextRange * diff --git a/amd/comgr/src/comgr-hotswap-internal.h b/amd/comgr/src/comgr-hotswap-internal.h index 35daca85023f2..e7068fed6abce 100644 --- a/amd/comgr/src/comgr-hotswap-internal.h +++ b/amd/comgr/src/comgr-hotswap-internal.h @@ -313,12 +313,19 @@ struct Trampoline { uint64_t OriginalOffset = 0; uint32_t OriginalSize = 0; llvm::SmallVector Bytes; - // When set, the pool is beyond s_branch reach. The source site branches to a - // nearby NOP gateway, which uses the scratch-backed gfx12 set-PC sequence - // to reach the pool without executing s_add_pc_i64. + // When set, the pool is beyond s_branch reach. The source and return edges + // use safe branch islands and, when a dead register pair is available, the + // scratch-backed gfx12 set-PC sequence. Neither executes s_add_pc_i64. bool Long = false; bool UsesSetPCBack = false; unsigned LongBranchSgprBase = 0; + // When numbered SGPRs are exhausted, a far edge may use VCC after proving + // that the replacement does not consume its incoming value and that VCC is + // dead at the continuation. + bool LongBranchUsesVcc = false; + // A wave32 far edge may preserve live VCC_LO in one safe numbered SGPR. + // Its source tail is a restore-and-fallthrough landing pad. + bool LongBranchPreservesVcc = false; bool HasPoolBranchIsland = false; uint64_t PoolBranchIslandOffset = 0; bool UsesShortBranchForward = false; @@ -326,9 +333,28 @@ struct Trampoline { llvm::SmallVector DirectSetPCForwardBytes; llvm::SmallVector ForwardBranchIslands; uint64_t ForwardBranchTargetOffset = 0; + // Dwords after the source's forward sequence are unreachable. A spare tail + // dword in an eight-byte or larger/coalesced source window can therefore + // serve as one safe, registerless relay for another far edge. + bool HasSourceTailBranchIsland = false; + uint64_t SourceTailBranchIslandOffset = 0; + uint64_t SourceTailBranchTargetOffset = 0; + llvm::SmallVector ReturnBranchIslands; + uint64_t ReturnBranchTargetOffset = 0; bool HasForwardGateway = false; uint64_t ForwardGatewayOffset = 0; llvm::SmallVector ForwardGatewayBytes; + // Multiple 8-byte far sites can share one SCC-neutral gateway. Each source + // records its PC and branches to that gateway; a dispatcher prefixed to the + // first group trampoline maps the source PC to the corresponding body. + bool UsesSharedDispatcherForward = false; + uint32_t SharedDispatcherGroup = 0; + unsigned SharedDispatcherSgprBase = 0; + uint64_t SharedDispatcherGatewayOffset = 0; + uint64_t SharedDispatcherRelayOffset = 0; + uint64_t SharedDispatcherSecondaryGatewayOffset = 0; + llvm::SmallVector SecondaryForwardGatewayBytes; + uint32_t PoolEntryPrefixBytes = 0; // A far-site run may only be coalesced within one known function. Unknown // ranges stay unmerged because adjacent symbols are independent entries. bool HasFunctionRange = false; @@ -477,6 +503,12 @@ static constexpr uint32_t MinInstSize = 4; static constexpr uint32_t SetPcReturnReserveBytes = 20; static constexpr uint32_t SetPcForwardSequenceBytes = SetPcReturnReserveBytes; +static constexpr uint32_t VccSaveRestoreBytes = MinInstSize; +static constexpr uint32_t VccPreservingReturnReserveBytes = + VccSaveRestoreBytes + SetPcReturnReserveBytes; +static constexpr uint32_t VccLandingPadBytes = MinInstSize; +static constexpr uint32_t VccPreservingSourceBytes = + MinInstSize + VccLandingPadBytes; static constexpr uint32_t PoolBranchIslandBytes = MinInstSize; // s_branch encoding: 16-bit signed dword offset field bounds. Used by @@ -559,6 +591,12 @@ class ElfView { /// Zero-size symbols extend to the next function symbol or `.text` end. std::vector functionTextRanges() const; + /// Validate that an object with a present, empty `.text` section contains + /// data only: no executable section contents, no defined function/ifunc + /// symbols, no kernel descriptor symbols, and no AMDGPU metadata kernel + /// entries. Malformed symbol tables, notes, or metadata fail closed. + bool isValidDataOnlyObject() const; + /// Find the kernel function symbol whose range includes \p TextAddress. /// Returns "" if no matching function symbol exists. std::string findKernelAtAddress(uint64_t TextAddress) const; @@ -1114,6 +1152,31 @@ struct VgprAllocator { return Base; } + /// Allocate \p N contiguous VGPRs above the kernel count without crossing a + /// \p BankSize-register boundary. Textual AMDGPU assembly only names + /// v0-v255; keeping a generated operand in one physical bank lets a caller + /// encode its low bits under one s_set_vgpr_msb mode. + std::optional + allocContiguousAboveKdInBank(unsigned N, unsigned Align = 2, + unsigned BankSize = 256) { + if (N == 0 || N > BankSize) + return std::nullopt; + unsigned OldNext = NextAboveKd; + unsigned Base = NextAboveKd; + if (Align > 1 && (Base % Align) != 0) + Base += Align - (Base % Align); + if (Base / BankSize != (Base + N - 1) / BankSize) + Base = ((Base / BankSize) + 1) * BankSize; + if (Align > 1 && (Base % Align) != 0) + Base += Align - (Base % Align); + if (Base + N > MaxVgprs) + return std::nullopt; + ExtraAllocated += (Base + N) - OldNext; + LiveAtPoint.set(Base, Base + N); + NextAboveKd = Base + N; + return Base; + } + unsigned extraVgprsNeeded() const { return ExtraAllocated; } }; @@ -1203,6 +1266,10 @@ struct SafeSgprUsageSummary { struct DirectControlFlowInfo { llvm::DenseSet Targets; + // Register-based transfers whose complete finite target set was proven. + // These do not make every instruction in their containing function a + // potential indirect destination. + llvm::DenseSet BoundedIndirectTransfers; bool HasUnresolvedTargets = false; }; @@ -1273,6 +1340,32 @@ struct PatchContext { llvm::StringMap KernelVgprGranuleCache; }; +enum class VgprMsbOperand : unsigned { + Src0 = 0, + Src1 = 2, + Src2 = 4, + Dst = 6, +}; + +/// Populate PatchContext::VgprMsbModeBefore if it has not been computed yet. +void ensureVgprMsbModes(PatchContext &Ctx); + +/// Return the exact VGPR-MSB mode before Decoded[Idx] proven by whole-function +/// CFG analysis. +[[nodiscard]] std::optional getActiveVgprMsbMode(PatchContext &Ctx, + size_t Idx); + +/// Recover an exact mode by scanning backward through the local straight-line +/// instruction sequence containing \p Idx. This is intentionally separate +/// from CFG mode recovery: only a lowering whose original operands already +/// depend on the local setter may use it when unrelated opaque control flow +/// prevents object-wide analysis. +[[nodiscard]] std::optional +getLocallyEstablishedVgprMsbMode(PatchContext &Ctx, size_t Idx); + +unsigned getVgprMsbBank(unsigned Mode, VgprMsbOperand Operand); +void setVgprMsbBank(unsigned &Mode, VgprMsbOperand Operand, unsigned Bank); + /// Return occupancy limits for \p Processor from COMGR's ISA metadata table. std::optional getSubtargetOccupancyLimits(llvm::StringRef Processor); @@ -1314,7 +1407,8 @@ struct SafeSgprScratchBlock { /// nullopt after logging when no block fits below RewriteConfig::MaxSgprs. std::optional findSafeSgprScratchBlock(PatchContext &Ctx, uint64_t TextOffset, unsigned Count, - unsigned Alignment, llvm::StringRef Context); + unsigned Alignment, llvm::StringRef Context, + bool ReportNoSpace = true); /// Charge a previously selected global block to the kernel owning \p /// TextOffset. If the site is in an ordinary device function, conservatively @@ -1333,12 +1427,15 @@ bool commitSafeSgprScratchBlock(PatchContext &Ctx, uint64_t TextOffset, uint32_t InstSize, llvm::ArrayRef Replacement); -/// Encode an SCC-neutral indirect long branch using the aligned SGPR pair at -/// \p SgprBase. The displacement uses gfx12's s_add_nc_u64; no s_add_pc_i64 +/// Encode an SCC-neutral indirect long branch using either the aligned +/// numbered pair at \p SgprBase or VCC when \p UseVcc is true. The caller must +/// prove VCC dead across the edge or preserve its wave32 low half before +/// selecting it. The displacement uses gfx12's s_add_nc_u64; no s_add_pc_i64 /// is emitted. std::optional> encodeSetPCLongBranch(const LLVMState &LS, uint64_t FromOffset, - uint64_t TargetOffset, unsigned SgprBase); + uint64_t TargetOffset, unsigned SgprBase, + bool UseVcc = false); struct EncodedSetPcGateway { NopSled *Sled = nullptr; @@ -1347,19 +1444,22 @@ struct EncodedSetPcGateway { /// Find the nearest short-branch-reachable gateway whose remaining space fits /// the set-PC sequence. Candidate widths are computed from the displacement; -/// only the selected candidate is encoded. The returned plan does not advance -/// the sled or modify text. +/// only the selected candidate is encoded. When \p PreserveVcc is true, +/// prepend a VCC_LO save to \p SgprBase. The returned plan does not advance the +/// sled or modify text. llvm::Expected> findNearestSetPcGateway(std::vector &Gateways, const LLVMState &LS, uint64_t FromOffset, uint64_t TargetOffset, - unsigned SgprBase); + unsigned SgprBase, bool UseVcc = false, + bool PreserveVcc = false); /// Count set-PC gateway slots reachable from \p FromOffset, up to \p MaxSlots. /// Candidate widths are computed without assembly. Zero means that no /// candidate fits; an Error means that a reachable candidate is invalid. llvm::Expected countReachableSetPcGatewaySlots( llvm::ArrayRef Gateways, const LLVMState &LS, uint64_t FromOffset, - uint64_t TargetOffset, unsigned SgprBase, uint64_t MaxSlots); + uint64_t TargetOffset, unsigned SgprBase, uint64_t MaxSlots, + bool UseVcc = false, bool PreserveVcc = false); /// Return whether an s_branch at \p From can encode \p To, including the /// instruction-relative PC base, alignment, signed range, and overflow checks. @@ -1413,7 +1513,41 @@ std::optional collectDirectBranchTargets( uint64_t TextAddr, uint64_t TextSize, llvm::ArrayRef DeclaredEntries, llvm::ArrayRef FunctionRanges = {}, - llvm::ArrayRef ExternalEntries = {}); + llvm::ArrayRef ExternalEntries = {}, + llvm::ArrayRef Text = {}); + +/// Return whether \p DI consumes the incoming value of \p Register, including +/// explicit read/modify/write destinations represented by MC tied-operand +/// constraints. This conservative query underpins scratch-register liveness. +[[nodiscard]] bool instructionReadsRegister(const InternalDecodedInst &DI, + const LLVMState &LS, + llvm::MCRegister Register); + +/// Resolve s0 through s(MaxSgprs - 1) to their physical MC registers. +std::optional> +resolveNumberedSgprRegisters(const llvm::MCRegisterInfo &MRI, + unsigned MaxSgprs); + +/// Collect numbered SGPR uses and definitions using MC register overlap, so +/// tuple and subregister operands conservatively affect their numbered SGPRs. +void getNumberedSgprUsesAndDefs(const InternalDecodedInst &DI, + const LLVMState &LS, + llvm::ArrayRef NumberedSgprs, + llvm::BitVector &Uses, llvm::BitVector &Defs); + +/// Return numbered SGPR incoming values observed by a replacement before a +/// definition, conservatively retaining values at malformed or opaque control +/// flow. +llvm::BitVector unsafeIncomingNumberedSgprsInReplacement( + llvm::ArrayRef Replacement, const LLVMState &LS, + llvm::ArrayRef NumberedSgprs); + +/// Return numbered SGPR incoming values that may be read before a definition, +/// or remain live at an opaque or invalid control-flow boundary. +std::optional unsafeIncomingNumberedSgprsInRange( + llvm::ArrayRef Decoded, const LLVMState &LS, + uint64_t FunctionBegin, uint64_t FunctionEnd, uint64_t Continuation, + llvm::ArrayRef NumberedSgprs); [[nodiscard]] bool emitReplacementCode(PatchContext &Ctx, uint64_t InstOffset, uint32_t InstSize, diff --git a/amd/comgr/src/comgr-hotswap-occupancy.cpp b/amd/comgr/src/comgr-hotswap-occupancy.cpp index 08c5ffc8f4039..f27bb2cd37a6a 100644 --- a/amd/comgr/src/comgr-hotswap-occupancy.cpp +++ b/amd/comgr/src/comgr-hotswap-occupancy.cpp @@ -225,6 +225,19 @@ VgprBumpDecision checkKernelVgprBump(PatchContext &Ctx, StringRef KernelName, if (!Capacity) return failOrDeclineVgprBump(Ctx, Requirement); + // Some input kernels already reserve enough VGPRs that their declared + // maximum workgroup cannot meet the generic waves/EU target. Do not reject + // a bump merely because it cannot repair that pre-existing condition: the + // patch is occupancy-neutral when the proposed allocation admits at least + // as many waves as the input allocation. + std::optional CurrentCapacity = computeWorkgroupCapacity( + *CurrentVgprs, Cached->second->MaxFlatWorkgroupSize, + Cached->second->WavefrontSize, *Limits); + if (!CurrentCapacity) + return failOrDeclineVgprBump(Ctx, Requirement); + if (Capacity->AchievableWavesPerEU >= CurrentCapacity->AchievableWavesPerEU) + return VgprBumpDecision::Apply; + VgprBumpDecision Decision = decideVgprBump(Requirement, *Capacity); if (Decision == VgprBumpDecision::Apply) return Decision; diff --git a/amd/comgr/src/comgr-hotswap-patch-trampoline.cpp b/amd/comgr/src/comgr-hotswap-patch-trampoline.cpp index f9beef347793a..ef3c2cf05802f 100755 --- a/amd/comgr/src/comgr-hotswap-patch-trampoline.cpp +++ b/amd/comgr/src/comgr-hotswap-patch-trampoline.cpp @@ -42,6 +42,7 @@ #include #include #include +#include #include using namespace llvm; @@ -402,6 +403,95 @@ std::optional> expandDs2AddrImpl(const MCInst &Inst, return std::nullopt; } +bool hasUnencodableVgprName(StringRef Asm) { + for (size_t Pos = Asm.find('v'); Pos != StringRef::npos; + Pos = Asm.find('v', Pos + 1)) { + StringRef Tail = Asm.substr(Pos + 1); + Tail.consume_front("["); + unsigned Index = 0; + if (!Tail.consumeInteger(10, Index) && Index > 255) + return true; + } + return false; +} + +bool normalizeVgprOperand(StringRef Input, VgprMsbOperand Role, unsigned &Mode, + std::string &Output) { + StringRef Operand = Input.trim(); + StringRef Suffix; + size_t Space = Operand.find(' '); + if (Space != StringRef::npos) { + Suffix = Operand.substr(Space); + Operand = Operand.take_front(Space); + } + if (!Operand.consume_front("v")) { + Output = Input.trim().str(); + return true; + } + + bool IsRange = Operand.consume_front("["); + if (IsRange && !Operand.consume_back("]")) + return false; + StringRef LoText; + StringRef HiText; + std::tie(LoText, HiText) = Operand.split(':'); + if (!IsRange) + LoText = HiText = Operand; + if (LoText.empty() || HiText.empty()) + return false; + + unsigned Lo = 0; + unsigned Hi = 0; + if (LoText.getAsInteger(10, Lo) || HiText.getAsInteger(10, Hi) || Hi < Lo || + Lo / 256 != Hi / 256) + return false; + unsigned Bank = Lo / 256; + if (Bank > 3) + return false; + setVgprMsbBank(Mode, Role, Bank); + if (IsRange) + Output = + ("v[" + Twine(Lo & 255) + ":" + Twine(Hi & 255) + "]" + Suffix).str(); + else + Output = ("v" + Twine(Lo & 255) + Suffix).str(); + return true; +} + +std::optional> +normalizeDsVgprBanks(StringRef Asm, StringRef FromMnem, unsigned OldMode) { + size_t MnemEnd = Asm.find(' '); + if (MnemEnd == StringRef::npos) + return std::nullopt; + StringRef Mnem = Asm.take_front(MnemEnd); + SmallVector Operands; + Asm.substr(MnemEnd + 1) + .split(Operands, ',', /*MaxSplit=*/-1, + /*KeepEmpty=*/false); + + SmallVector Roles; + if (FromMnem.starts_with("ds_load_")) + Roles = {VgprMsbOperand::Dst, VgprMsbOperand::Src0}; + else if (FromMnem.starts_with("ds_storexchg_")) + Roles = {VgprMsbOperand::Dst, VgprMsbOperand::Src0, VgprMsbOperand::Src1}; + else if (FromMnem.starts_with("ds_store_")) + Roles = {VgprMsbOperand::Src0, VgprMsbOperand::Src1}; + else + return std::nullopt; + if (Operands.size() != Roles.size()) + return std::nullopt; + + unsigned NewMode = OldMode; + std::string Normalized = Mnem.str(); + for (unsigned I = 0; I != Operands.size(); ++I) { + std::string Operand; + if (!normalizeVgprOperand(Operands[I], Roles[I], NewMode, Operand)) + return std::nullopt; + Normalized += I == 0 ? " " : ", "; + Normalized += Operand; + } + return std::pair{std::move(Normalized), NewMode}; +} + // -- patchDs2Addr ----------------------------------------------------------- // // Expand one ds_*_2addr_* instruction (stride64 or non-stride64) into two @@ -421,9 +511,44 @@ bool patchDs2Addr(PatchContext &Ctx, size_t Idx) { if (!Expanded) return failRequiredPatch(Ctx); + bool NeedsBankNormalization = + llvm::any_of(*Expanded, [](const std::string &Asm) { + return hasUnencodableVgprName(Asm); + }); + std::optional ActiveMode; + if (NeedsBankNormalization) { + ActiveMode = getActiveVgprMsbMode(Ctx, Idx); + if (!ActiveMode) { + log() << "hotswap: error: ds_2addr at 0x" << utohexstr(DI.Offset) + << " crosses v255 but the active VGPR-MSB mode is unknown\n"; + return failRequiredPatch(Ctx); + } + } + std::string Combined; - for (const std::string &Line : *Expanded) - Combined += Line + "\n"; + for (const std::string &Line : *Expanded) { + if (!NeedsBankNormalization) { + Combined += Line + "\n"; + continue; + } + std::optional> Normalized = + normalizeDsVgprBanks(Line, DI.Mnemonic, *ActiveMode); + if (!Normalized) { + log() << "hotswap: error: ds_2addr at 0x" << utohexstr(DI.Offset) + << " has a VGPR operand that crosses a 256-register bank\n"; + return failRequiredPatch(Ctx); + } + unsigned NewMode = Normalized->second; + if (NewMode != *ActiveMode) + Combined += + ("s_set_vgpr_msb " + Twine(NewMode | (*ActiveMode << 8)) + "\n") + .str(); + Combined += Normalized->first + "\n"; + if (NewMode != *ActiveMode) + Combined += + ("s_set_vgpr_msb " + Twine(*ActiveMode | (NewMode << 8)) + "\n") + .str(); + } // Drain the DS counter right after the split pair so both halves are // guaranteed complete before any downstream consumer. The original code // tracked completion of the single 2-addr instruction via a later diff --git a/amd/comgr/src/comgr-hotswap-patch-wmma-scale16.cpp b/amd/comgr/src/comgr-hotswap-patch-wmma-scale16.cpp index 9a05d05aa1019..0b676fc5bee7c 100644 --- a/amd/comgr/src/comgr-hotswap-patch-wmma-scale16.cpp +++ b/amd/comgr/src/comgr-hotswap-patch-wmma-scale16.cpp @@ -34,9 +34,12 @@ /// Each pass's block-32 scale is a byte-gather of the block-16 scale bytes: /// even bytes feed the low subblocks, odd bytes the high ones. /// -/// The masked A copy lands in a contiguous VGPR block allocated above the -/// kernel's VGPR count and below MaxVgprs (256 on GFX1250), so it stays in VGPR -/// bank 0 and needs no s_set_vgpr_msb switch. +/// The replacement is assembled from textual register names, for which the +/// AMDGPU parser accepts v0-v255. Scratch may nevertheless live above v255: +/// the generated assembly encodes each scratch VGPR's low byte and brackets +/// mixed-bank instructions with role-specific s_set_vgpr_msb transitions. +/// Matrix B is copied into the scratch bank so the lowered WMMA can use one +/// src1 bank for both its gathered scale and matrix-B operands. /// /// Fail-closed fallback: when the scratch budget (A-width VGPRs plus a few /// scale/temp VGPRs and one scratch SGPR) is unavailable, the pass marks the @@ -56,6 +59,8 @@ #include "llvm/ADT/Twine.h" #include "llvm/Support/raw_ostream.h" +#include + using namespace llvm; namespace COMGR { @@ -68,9 +73,14 @@ static constexpr unsigned VOP3PXSize = 16; // AMDGPU SRC operand encoding: VGPRs are 256 + N. static constexpr unsigned VgprEncBase = 256; +static constexpr unsigned VgprBankSize = 256; static std::string vgprName(unsigned N) { return ("v" + Twine(N)).str(); } +static std::string encodedVgprName(unsigned Physical) { + return vgprName(Physical % VgprBankSize); +} + static bool isVgprEncoding(unsigned Enc) { return Enc >= VgprEncBase; } static std::optional decodeVgprEncoding(unsigned Enc) { @@ -105,6 +115,7 @@ static void writeScaleSrc1(uint8_t *Raw, unsigned Enc) { // -- Base WMMA uop field accessors (bytes 8-15) ------------------------------ // VDST: byte[8] (8-bit raw VGPR number, no +256) // SRC0: byte[12] + byte[13] bit[0] (9-bit; matrix A) +// SRC1: byte[13] bits[7:1] + byte[14] bits[1:0] (9-bit; matrix B) // SRC2: byte[14] bits[7:2] + byte[15] bits[2:0] (9-bit; accumulator C) static unsigned extractVdst(const uint8_t *Raw) { return Raw[8]; } @@ -114,6 +125,11 @@ static void writeSrc0(uint8_t *Raw, unsigned Enc) { Raw[13] = (Raw[13] & 0xFE) | ((Enc >> 8) & 0x01); } +static void writeSrc1(uint8_t *Raw, unsigned Enc) { + Raw[13] = (Raw[13] & 0x01) | ((Enc & 0x7F) << 1); + Raw[14] = (Raw[14] & 0xFC) | ((Enc >> 7) & 0x03); +} + static void writeSrc2(uint8_t *Raw, unsigned Enc) { Raw[14] = (Raw[14] & 0x03) | ((Enc & 0x3F) << 2); Raw[15] = (Raw[15] & 0xF8) | ((Enc >> 6) & 0x07); @@ -162,28 +178,105 @@ static SmallVector rewriteScale16ToScale(const uint8_t *OrigRaw, // (odd byte 2j+1) for pass-high, packed into one VGPR as // [byte0..3] = k-block 0..3. -static void emitGatherEven(raw_string_ostream &OS, StringRef Lo, StringRef Hi, - StringRef Dst, StringRef T) { +using VgprBankRequirement = std::pair; + +static void +emitModeForOperands(raw_string_ostream &OS, unsigned &CurrentMode, + std::initializer_list Requirements) { + unsigned NewMode = CurrentMode; + for (const VgprBankRequirement &Requirement : Requirements) + setVgprMsbBank(NewMode, Requirement.first, Requirement.second); + if (NewMode == CurrentMode) + return; + OS << "s_set_vgpr_msb " << (NewMode | (CurrentMode << 8)) << "\n"; + CurrentMode = NewMode; +} + +static void emitGatherEven(raw_string_ostream &OS, unsigned Lo, unsigned Hi, + unsigned Dst, unsigned T, unsigned ScratchBank, + unsigned &CurrentMode) { + std::string LoName = encodedVgprName(Lo); + std::string HiName = encodedVgprName(Hi); + std::string DstName = encodedVgprName(Dst); + std::string TName = encodedVgprName(T); + // Dst = { Lo[7:0], Lo[23:16], Hi[7:0], Hi[23:16] } (bytes 0,2,4,6) - OS << "v_and_b32 " << Dst << ", 0xff, " << Lo << "\n"; - OS << "v_bfe_u32 " << T << ", " << Lo << ", 16, 8\n"; - OS << "v_lshl_or_b32 " << Dst << ", " << T << ", 8, " << Dst << "\n"; - OS << "v_and_b32 " << T << ", 0xff, " << Hi << "\n"; - OS << "v_lshl_or_b32 " << Dst << ", " << T << ", 16, " << Dst << "\n"; - OS << "v_bfe_u32 " << T << ", " << Hi << ", 16, 8\n"; - OS << "v_lshl_or_b32 " << Dst << ", " << T << ", 24, " << Dst << "\n"; + emitModeForOperands(OS, CurrentMode, + {{VgprMsbOperand::Dst, ScratchBank}, + {VgprMsbOperand::Src1, Lo / VgprBankSize}}); + OS << "v_and_b32 " << DstName << ", 0xff, " << LoName << "\n"; + emitModeForOperands(OS, CurrentMode, + {{VgprMsbOperand::Dst, ScratchBank}, + {VgprMsbOperand::Src0, Lo / VgprBankSize}}); + OS << "v_bfe_u32 " << TName << ", " << LoName << ", 16, 8\n"; + emitModeForOperands(OS, CurrentMode, + {{VgprMsbOperand::Dst, ScratchBank}, + {VgprMsbOperand::Src0, ScratchBank}, + {VgprMsbOperand::Src2, ScratchBank}}); + OS << "v_lshl_or_b32 " << DstName << ", " << TName << ", 8, " << DstName + << "\n"; + emitModeForOperands(OS, CurrentMode, + {{VgprMsbOperand::Dst, ScratchBank}, + {VgprMsbOperand::Src1, Hi / VgprBankSize}}); + OS << "v_and_b32 " << TName << ", 0xff, " << HiName << "\n"; + emitModeForOperands(OS, CurrentMode, + {{VgprMsbOperand::Dst, ScratchBank}, + {VgprMsbOperand::Src0, ScratchBank}, + {VgprMsbOperand::Src2, ScratchBank}}); + OS << "v_lshl_or_b32 " << DstName << ", " << TName << ", 16, " << DstName + << "\n"; + emitModeForOperands(OS, CurrentMode, + {{VgprMsbOperand::Dst, ScratchBank}, + {VgprMsbOperand::Src0, Hi / VgprBankSize}}); + OS << "v_bfe_u32 " << TName << ", " << HiName << ", 16, 8\n"; + emitModeForOperands(OS, CurrentMode, + {{VgprMsbOperand::Dst, ScratchBank}, + {VgprMsbOperand::Src0, ScratchBank}, + {VgprMsbOperand::Src2, ScratchBank}}); + OS << "v_lshl_or_b32 " << DstName << ", " << TName << ", 24, " << DstName + << "\n"; } -static void emitGatherOdd(raw_string_ostream &OS, StringRef Lo, StringRef Hi, - StringRef Dst, StringRef T) { +static void emitGatherOdd(raw_string_ostream &OS, unsigned Lo, unsigned Hi, + unsigned Dst, unsigned T, unsigned ScratchBank, + unsigned &CurrentMode) { + std::string LoName = encodedVgprName(Lo); + std::string HiName = encodedVgprName(Hi); + std::string DstName = encodedVgprName(Dst); + std::string TName = encodedVgprName(T); + // Dst = { Lo[15:8], Lo[31:24], Hi[15:8], Hi[31:24] } (bytes 1,3,5,7) - OS << "v_bfe_u32 " << Dst << ", " << Lo << ", 8, 8\n"; - OS << "v_bfe_u32 " << T << ", " << Lo << ", 24, 8\n"; - OS << "v_lshl_or_b32 " << Dst << ", " << T << ", 8, " << Dst << "\n"; - OS << "v_bfe_u32 " << T << ", " << Hi << ", 8, 8\n"; - OS << "v_lshl_or_b32 " << Dst << ", " << T << ", 16, " << Dst << "\n"; - OS << "v_lshrrev_b32 " << T << ", 24, " << Hi << "\n"; - OS << "v_lshl_or_b32 " << Dst << ", " << T << ", 24, " << Dst << "\n"; + emitModeForOperands(OS, CurrentMode, + {{VgprMsbOperand::Dst, ScratchBank}, + {VgprMsbOperand::Src0, Lo / VgprBankSize}}); + OS << "v_bfe_u32 " << DstName << ", " << LoName << ", 8, 8\n"; + OS << "v_bfe_u32 " << TName << ", " << LoName << ", 24, 8\n"; + emitModeForOperands(OS, CurrentMode, + {{VgprMsbOperand::Dst, ScratchBank}, + {VgprMsbOperand::Src0, ScratchBank}, + {VgprMsbOperand::Src2, ScratchBank}}); + OS << "v_lshl_or_b32 " << DstName << ", " << TName << ", 8, " << DstName + << "\n"; + emitModeForOperands(OS, CurrentMode, + {{VgprMsbOperand::Dst, ScratchBank}, + {VgprMsbOperand::Src0, Hi / VgprBankSize}}); + OS << "v_bfe_u32 " << TName << ", " << HiName << ", 8, 8\n"; + emitModeForOperands(OS, CurrentMode, + {{VgprMsbOperand::Dst, ScratchBank}, + {VgprMsbOperand::Src0, ScratchBank}, + {VgprMsbOperand::Src2, ScratchBank}}); + OS << "v_lshl_or_b32 " << DstName << ", " << TName << ", 16, " << DstName + << "\n"; + emitModeForOperands(OS, CurrentMode, + {{VgprMsbOperand::Dst, ScratchBank}, + {VgprMsbOperand::Src1, Hi / VgprBankSize}}); + OS << "v_lshrrev_b32 " << TName << ", 24, " << HiName << "\n"; + emitModeForOperands(OS, CurrentMode, + {{VgprMsbOperand::Dst, ScratchBank}, + {VgprMsbOperand::Src0, ScratchBank}, + {VgprMsbOperand::Src2, ScratchBank}}); + OS << "v_lshl_or_b32 " << DstName << ", " << TName << ", 24, " << DstName + << "\n"; } // A' = mask ? A : 0, per lane, for W consecutive VGPRs from ABase into SBase. @@ -193,11 +286,16 @@ static void emitGatherOdd(raw_string_ostream &OS, StringRef Lo, StringRef Hi, // high-16 in lanes 16-31, so a lane mask isolates a subblock. static void emitLaneMaskCopy(raw_string_ostream &OS, StringRef MaskSgpr, uint32_t MaskImm, unsigned SBase, unsigned ABase, - unsigned W) { + unsigned W, unsigned ScratchBank, + unsigned &CurrentMode) { OS << "s_mov_b32 " << MaskSgpr << ", 0x" << utohexstr(MaskImm) << "\n"; - for (unsigned I = 0; I < W; ++I) - OS << "v_cndmask_b32_e64 " << vgprName(SBase + I) << ", 0, " - << vgprName(ABase + I) << ", " << MaskSgpr << "\n"; + for (unsigned I = 0; I < W; ++I) { + emitModeForOperands(OS, CurrentMode, + {{VgprMsbOperand::Dst, ScratchBank}, + {VgprMsbOperand::Src1, (ABase + I) / VgprBankSize}}); + OS << "v_cndmask_b32_e64 " << encodedVgprName(SBase + I) << ", 0, " + << encodedVgprName(ABase + I) << ", " << MaskSgpr << "\n"; + } } // A' keeps the VGPRs of the low (KeepLow=true) or high 16-K subblocks and zeros @@ -210,25 +308,45 @@ static void emitLaneMaskCopy(raw_string_ostream &OS, StringRef MaskSgpr, // subblock's VGPRs instead. static void emitVgprSelectCopy(raw_string_ostream &OS, bool KeepLow, unsigned SBase, unsigned ABase, unsigned W, - unsigned SubW) { + unsigned SubW, unsigned ScratchBank, + unsigned &CurrentMode) { for (unsigned I = 0; I < W; ++I) { bool IsLow = ((I / SubW) % 2) == 0; - if (IsLow == KeepLow) - OS << "v_mov_b32 " << vgprName(SBase + I) << ", " << vgprName(ABase + I) - << "\n"; - else - OS << "v_mov_b32 " << vgprName(SBase + I) << ", 0\n"; + if (IsLow == KeepLow) { + emitModeForOperands(OS, CurrentMode, + {{VgprMsbOperand::Dst, ScratchBank}, + {VgprMsbOperand::Src0, (ABase + I) / VgprBankSize}}); + OS << "v_mov_b32 " << encodedVgprName(SBase + I) << ", " + << encodedVgprName(ABase + I) << "\n"; + } else { + emitModeForOperands(OS, CurrentMode, + {{VgprMsbOperand::Dst, ScratchBank}}); + OS << "v_mov_b32 " << encodedVgprName(SBase + I) << ", 0\n"; + } + } +} + +static void emitVgprCopy(raw_string_ostream &OS, unsigned DstBase, + unsigned SrcBase, unsigned W, unsigned ScratchBank, + unsigned &CurrentMode) { + for (unsigned I = 0; I < W; ++I) { + emitModeForOperands(OS, CurrentMode, + {{VgprMsbOperand::Dst, ScratchBank}, + {VgprMsbOperand::Src0, (SrcBase + I) / VgprBankSize}}); + OS << "v_mov_b32 " << encodedVgprName(DstBase + I) << ", " + << encodedVgprName(SrcBase + I) << "\n"; } } -// Parse the matrix-A (src0) VGPR range from the printer's canonical form. +// Parse a matrix VGPR range from the printer's canonical form. struct VgprRange { unsigned Base; unsigned Width; }; static std::optional -matrixAOperandRange(PatchContext &Ctx, const InternalDecodedInst &DI) { +matrixOperandRange(PatchContext &Ctx, const InternalDecodedInst &DI, + unsigned OperandIndex) { SmallString<256> Buf; raw_svector_ostream OS(Buf); Ctx.LS.MCIP->printInst(&DI.Inst, /*Address=*/0, /*Annot=*/"", *Ctx.LS.STI, @@ -238,17 +356,18 @@ matrixAOperandRange(PatchContext &Ctx, const InternalDecodedInst &DI) { if (MnemEnd == StringRef::npos) return std::nullopt; StringRef Rest = S.substr(MnemEnd).ltrim(); - // Operand 0 = vdst, operand 1 = src0 (matrix A). - size_t Comma0 = Rest.find(','); - if (Comma0 == StringRef::npos) - return std::nullopt; - Rest = Rest.substr(Comma0 + 1).ltrim(); - size_t Comma1 = Rest.find(','); - StringRef A = (Comma1 == StringRef::npos) ? Rest : Rest.substr(0, Comma1); - A = A.trim(); - if (!A.starts_with("v[") || !A.ends_with("]")) + for (unsigned I = 0; I < OperandIndex; ++I) { + size_t Comma = Rest.find(','); + if (Comma == StringRef::npos) + return std::nullopt; + Rest = Rest.substr(Comma + 1).ltrim(); + } + size_t End = Rest.find(','); + StringRef Operand = (End == StringRef::npos) ? Rest : Rest.substr(0, End); + Operand = Operand.trim(); + if (!Operand.starts_with("v[") || !Operand.ends_with("]")) return std::nullopt; - StringRef Inside = A.drop_front(2).drop_back(1); + StringRef Inside = Operand.drop_front(2).drop_back(1); StringRef LoS, HiS; std::tie(LoS, HiS) = Inside.split(':'); unsigned Lo = 0, Hi = 0; @@ -329,14 +448,40 @@ static uint32_t patchWmmaScale16_16x16(PatchContext &Ctx, size_t Idx) { if (!ScaleABase || !ScaleBBase) return failClosed(Ctx, DI, "non-VGPR block-16 scale operand"); - unsigned ScaleALo = *ScaleABase, ScaleAHi = ScaleALo + 1; - unsigned ScaleBLo = *ScaleBBase, ScaleBHi = ScaleBLo + 1; - - std::optional ARange = matrixAOperandRange(Ctx, DI); - if (!ARange) - return failClosed(Ctx, DI, "could not determine matrix-A VGPR range"); - unsigned ABase = ARange->Base; + std::optional ActiveMode = getActiveVgprMsbMode(Ctx, Idx); + // A compiler-emitted scale16 whose immediately preceding instruction sets + // the mode already depends on that setter for the original fused operands. + // Preserve that local contract when unrelated opaque control flow prevents + // object-wide mode recovery. + if (!ActiveMode) + ActiveMode = getLocallyEstablishedVgprMsbMode(Ctx, Idx); + if (!ActiveMode) + return failClosed(Ctx, DI, "cannot determine active VGPR-MSB mode"); + + unsigned OrigSrc0Bank = getVgprMsbBank(*ActiveMode, VgprMsbOperand::Src0); + unsigned OrigSrc1Bank = getVgprMsbBank(*ActiveMode, VgprMsbOperand::Src1); + unsigned OrigDstBank = getVgprMsbBank(*ActiveMode, VgprMsbOperand::Dst); + + unsigned ScaleALo = *ScaleABase + OrigSrc0Bank * VgprBankSize; + unsigned ScaleAHi = ScaleALo + 1; + unsigned ScaleBLo = *ScaleBBase + OrigSrc1Bank * VgprBankSize; + unsigned ScaleBHi = ScaleBLo + 1; + if (ScaleAHi >= Ctx.Config.MaxVgprs || ScaleBHi >= Ctx.Config.MaxVgprs) + return failClosed(Ctx, DI, "block-16 scale tuple exceeds VGPR capacity"); + + std::optional ARange = + matrixOperandRange(Ctx, DI, /*OperandIndex=*/1); + std::optional BRange = + matrixOperandRange(Ctx, DI, /*OperandIndex=*/2); + if (!ARange || !BRange) + return failClosed(Ctx, DI, "could not determine matrix-A/B VGPR ranges"); + unsigned ABase = ARange->Base + OrigSrc0Bank * VgprBankSize; unsigned AWidth = ARange->Width; + unsigned BBase = BRange->Base + OrigSrc1Bank * VgprBankSize; + unsigned BWidth = BRange->Width; + if (ABase + AWidth > Ctx.Config.MaxVgprs || + BBase + BWidth > Ctx.Config.MaxVgprs) + return failClosed(Ctx, DI, "matrix operand exceeds VGPR capacity"); // The masking scheme depends on the matrix-A data format. std::optional Plan = matrixAMaskPlan(Ctx, DI); @@ -362,18 +507,28 @@ static uint32_t patchWmmaScale16_16x16(PatchContext &Ctx, size_t Idx) { VgprAllocator Alloc(Ctx.Liveness.liveBefore(Idx), KdCount, Ctx.Config.MaxVgprs); - // Four block-32 scale VGPRs (A/B x low/high) plus one byte-extraction temp, - // then a contiguous, even-aligned block for the masked A copy. - std::optional ScaleAloReg = Alloc.alloc(); - std::optional ScaleBloReg = Alloc.alloc(); - std::optional ScaleAhiReg = Alloc.alloc(); - std::optional ScaleBhiReg = Alloc.alloc(); - std::optional TmpReg = Alloc.alloc(); - std::optional SBase = - Alloc.allocContiguousAboveKd(AWidth, /*Align=*/2); - if (!ScaleAloReg || !ScaleBloReg || !ScaleAhiReg || !ScaleBhiReg || !TmpReg || - !SBase) - return failClosed(Ctx, DI, "insufficient scratch VGPRs for exact K-split"); + // Keep every generated operand in one physical VGPR bank. Five scalar + // temporaries precede an even-aligned masked-A block and a matrix-B copy. + // Copying B is necessary because both scale-B and matrix-B use the WMMA + // src1 VGPR-MSB field. + constexpr unsigned ScalarScratchCount = 5; + unsigned AOffset = (ScalarScratchCount + 1) & ~1u; + unsigned BOffset = AOffset + AWidth; + unsigned ScratchCount = BOffset + BWidth; + std::optional ScratchBase = Alloc.allocContiguousAboveKdInBank( + ScratchCount, /*Align=*/2, VgprBankSize); + if (!ScratchBase) + return failClosed(Ctx, DI, + "no single-bank above-KD VGPR block for exact K-split"); + + unsigned ScaleAloReg = *ScratchBase; + unsigned ScaleBloReg = *ScratchBase + 1; + unsigned ScaleAhiReg = *ScratchBase + 2; + unsigned ScaleBhiReg = *ScratchBase + 3; + unsigned TmpReg = *ScratchBase + 4; + unsigned SBase = *ScratchBase + AOffset; + unsigned BCopyBase = *ScratchBase + BOffset; + unsigned ScratchBank = *ScratchBase / VgprBankSize; // The lane-mask scheme (FP8/BF8) needs one scratch SGPR for the wave-lane // bitmask; the VGPR-select scheme (FP4/FP6) uses plain v_mov and needs none. @@ -389,53 +544,97 @@ static uint32_t patchWmmaScale16_16x16(PatchContext &Ctx, size_t Idx) { } // Preamble + pass-low masked copy (assembled together), then pass-high copy. - std::string PreAsm, HiAsm; - raw_string_ostream PreOS(PreAsm), HiOS(HiAsm); - - emitGatherEven(PreOS, vgprName(ScaleALo), vgprName(ScaleAHi), - vgprName(*ScaleAloReg), vgprName(*TmpReg)); - emitGatherEven(PreOS, vgprName(ScaleBLo), vgprName(ScaleBHi), - vgprName(*ScaleBloReg), vgprName(*TmpReg)); - emitGatherOdd(PreOS, vgprName(ScaleALo), vgprName(ScaleAHi), - vgprName(*ScaleAhiReg), vgprName(*TmpReg)); - emitGatherOdd(PreOS, vgprName(ScaleBLo), vgprName(ScaleBHi), - vgprName(*ScaleBhiReg), vgprName(*TmpReg)); + std::string PreAsm, HiAsm, PostAsm; + raw_string_ostream PreOS(PreAsm), HiOS(HiAsm), PostOS(PostAsm); + unsigned PreMode = *ActiveMode; + + emitGatherEven(PreOS, ScaleALo, ScaleAHi, ScaleAloReg, TmpReg, ScratchBank, + PreMode); + emitGatherEven(PreOS, ScaleBLo, ScaleBHi, ScaleBloReg, TmpReg, ScratchBank, + PreMode); + emitGatherOdd(PreOS, ScaleALo, ScaleAHi, ScaleAhiReg, TmpReg, ScratchBank, + PreMode); + emitGatherOdd(PreOS, ScaleBLo, ScaleBHi, ScaleBhiReg, TmpReg, ScratchBank, + PreMode); + emitVgprCopy(PreOS, BCopyBase, BBase, BWidth, ScratchBank, PreMode); if (Plan->Scheme == AMaskScheme::Lane) { // pass-low keeps lanes 0-15 (low-16 subblocks); pass-high lanes 16-31. - emitLaneMaskCopy(PreOS, MaskS, 0x0000FFFFu, *SBase, ABase, AWidth); - emitLaneMaskCopy(HiOS, MaskS, 0xFFFF0000u, *SBase, ABase, AWidth); + emitLaneMaskCopy(PreOS, MaskS, 0x0000FFFFu, SBase, ABase, AWidth, + ScratchBank, PreMode); } else { // pass-low keeps the low-16 subblock VGPRs; pass-high the high-16 ones. - emitVgprSelectCopy(PreOS, /*KeepLow=*/true, *SBase, ABase, AWidth, - Plan->SubW); - emitVgprSelectCopy(HiOS, /*KeepLow=*/false, *SBase, ABase, AWidth, - Plan->SubW); + emitVgprSelectCopy(PreOS, /*KeepLow=*/true, SBase, ABase, AWidth, + Plan->SubW, ScratchBank, PreMode); } - SmallVector PreBytes = assembleInstructions(PreAsm, Ctx.LS); - SmallVector HiBytes = assembleInstructions(HiAsm, Ctx.LS); - if (PreBytes.empty() || HiBytes.empty()) - return failClosed(Ctx, DI, "preamble assembly failed"); + unsigned WmmaLoMode = *ActiveMode; + setVgprMsbBank(WmmaLoMode, VgprMsbOperand::Src0, ScratchBank); + setVgprMsbBank(WmmaLoMode, VgprMsbOperand::Src1, ScratchBank); + emitModeForOperands( + PreOS, PreMode, + {{VgprMsbOperand::Src0, ScratchBank}, + {VgprMsbOperand::Src1, ScratchBank}, + {VgprMsbOperand::Src2, getVgprMsbBank(WmmaLoMode, VgprMsbOperand::Src2)}, + {VgprMsbOperand::Dst, getVgprMsbBank(WmmaLoMode, VgprMsbOperand::Dst)}}); // pass-low WMMA: matrix A = masked copy, scales = even-byte gathers, src2 = // original C (preserved by the byte copy). - SmallVector WmmaLo = - rewriteScale16ToScale(Raw, DI.Size, VgprEncBase + *ScaleAloReg, - VgprEncBase + *ScaleBloReg, Ctx.LS); + SmallVector WmmaLo = rewriteScale16ToScale( + Raw, DI.Size, VgprEncBase + (ScaleAloReg % VgprBankSize), + VgprEncBase + (ScaleBloReg % VgprBankSize), Ctx.LS); if (WmmaLo.empty()) return failClosed(Ctx, DI, "pass-low WMMA rewrite failed"); - writeSrc0(WmmaLo.data(), VgprEncBase + *SBase); + writeSrc0(WmmaLo.data(), VgprEncBase + (SBase % VgprBankSize)); + writeSrc1(WmmaLo.data(), VgprEncBase + (BCopyBase % VgprBankSize)); // pass-high WMMA: odd-byte gathers, and src2 = D so it accumulates onto the // pass-low result. - SmallVector WmmaHi = - rewriteScale16ToScale(Raw, DI.Size, VgprEncBase + *ScaleAhiReg, - VgprEncBase + *ScaleBhiReg, Ctx.LS); + SmallVector WmmaHi = rewriteScale16ToScale( + Raw, DI.Size, VgprEncBase + (ScaleAhiReg % VgprBankSize), + VgprEncBase + (ScaleBhiReg % VgprBankSize), Ctx.LS); if (WmmaHi.empty()) return failClosed(Ctx, DI, "pass-high WMMA rewrite failed"); - writeSrc0(WmmaHi.data(), VgprEncBase + *SBase); + writeSrc0(WmmaHi.data(), VgprEncBase + (SBase % VgprBankSize)); + writeSrc1(WmmaHi.data(), VgprEncBase + (BCopyBase % VgprBankSize)); writeSrc2(WmmaHi.data(), VgprEncBase + extractVdst(Raw)); + unsigned HiMode = WmmaLoMode; + if (Plan->Scheme == AMaskScheme::Lane) { + emitLaneMaskCopy(HiOS, MaskS, 0xFFFF0000u, SBase, ABase, AWidth, + ScratchBank, HiMode); + } else { + emitVgprSelectCopy(HiOS, /*KeepLow=*/false, SBase, ABase, AWidth, + Plan->SubW, ScratchBank, HiMode); + } + unsigned WmmaHiMode = WmmaLoMode; + setVgprMsbBank(WmmaHiMode, VgprMsbOperand::Src2, OrigDstBank); + emitModeForOperands( + HiOS, HiMode, + {{VgprMsbOperand::Src0, ScratchBank}, + {VgprMsbOperand::Src1, ScratchBank}, + {VgprMsbOperand::Src2, OrigDstBank}, + {VgprMsbOperand::Dst, getVgprMsbBank(WmmaHiMode, VgprMsbOperand::Dst)}}); + + unsigned PostMode = WmmaHiMode; + unsigned ActiveSrc0 = getVgprMsbBank(*ActiveMode, VgprMsbOperand::Src0); + unsigned ActiveSrc1 = getVgprMsbBank(*ActiveMode, VgprMsbOperand::Src1); + unsigned ActiveSrc2 = getVgprMsbBank(*ActiveMode, VgprMsbOperand::Src2); + unsigned ActiveDst = getVgprMsbBank(*ActiveMode, VgprMsbOperand::Dst); + emitModeForOperands(PostOS, PostMode, + {{VgprMsbOperand::Src0, ActiveSrc0}, + {VgprMsbOperand::Src1, ActiveSrc1}, + {VgprMsbOperand::Src2, ActiveSrc2}, + {VgprMsbOperand::Dst, ActiveDst}}); + + SmallVector PreBytes = assembleInstructions(PreAsm, Ctx.LS); + SmallVector HiBytes = assembleInstructions(HiAsm, Ctx.LS); + SmallVector PostBytes; + if (!PostAsm.empty()) + PostBytes = assembleInstructions(PostAsm, Ctx.LS); + if (PreBytes.empty() || HiBytes.empty() || + (!PostAsm.empty() && PostBytes.empty())) + return failClosed(Ctx, DI, "mode-aware preamble assembly failed"); + // gfx1250 WMMA co-exec hazard: the pass-high copy (VALU) overwrites the // masked-A block the pass-low WMMA still reads, so it must not co-execute // with the in-flight WMMA. Insert the full required v_nop separation between @@ -454,6 +653,7 @@ static uint32_t patchWmmaScale16_16x16(PatchContext &Ctx, size_t Idx) { Replacement.append(VNop.begin(), VNop.end()); Replacement.append(HiBytes.begin(), HiBytes.end()); Replacement.append(WmmaHi.begin(), WmmaHi.end()); + Replacement.append(PostBytes.begin(), PostBytes.end()); unsigned Extra = Alloc.extraVgprsNeeded(); if (checkKernelVgprBump(Ctx, KernelName, Extra, PatchRequirement::Required) != @@ -481,8 +681,9 @@ static uint32_t patchWmmaScale16_16x16(PatchContext &Ctx, size_t Idx) { << utohexstr(DI.Offset) << " (" << (Plan->Scheme == AMaskScheme::Lane ? "lane-mask" : "vgpr-select") << ", A=v" << ABase << ":" << (ABase + AWidth - 1) << " -> masked v" - << *SBase << ", +" << Extra << " vgpr, " << A0Nops << " hazard v_nop, " - << Replacement.size() << " bytes)\n"; + << SBase << ", B copy=v" << BCopyBase << ":" << (BCopyBase + BWidth - 1) + << ", scratch bank " << ScratchBank << ", +" << Extra << " vgpr, " + << A0Nops << " hazard v_nop, " << Replacement.size() << " bytes)\n"; return 1; } diff --git a/amd/comgr/src/comgr-hotswap-patch-wmma-split.cpp b/amd/comgr/src/comgr-hotswap-patch-wmma-split.cpp index 23c2926966d39..2429d6d992e5b 100644 --- a/amd/comgr/src/comgr-hotswap-patch-wmma-split.cpp +++ b/amd/comgr/src/comgr-hotswap-patch-wmma-split.cpp @@ -966,13 +966,6 @@ findActiveVgprMsbMode(const PatchContext &Ctx, size_t Idx) { return static_cast(Ctx.VgprMsbModeBefore[Idx]); } -enum class VgprMsbOperand : unsigned { - Src0 = 0, - Src1 = 2, - Src2 = 4, - Dst = 6, -}; - unsigned getVgprMsbs(unsigned Mode, VgprMsbOperand Operand) { return (Mode >> static_cast(Operand)) & 0x3; } @@ -1131,6 +1124,53 @@ buildSplit32x16Asm(StringRef Replacement, const PrintedAsm &P, const WmmaOps &R, } // anonymous namespace +void ensureVgprMsbModes(PatchContext &Ctx) { + if (Ctx.VgprMsbModeBefore.empty()) + computeVgprMsbModes(Ctx); +} + +std::optional getActiveVgprMsbMode(PatchContext &Ctx, size_t Idx) { + ensureVgprMsbModes(Ctx); + return findActiveVgprMsbMode(Ctx, Idx); +} + +std::optional getLocallyEstablishedVgprMsbMode(PatchContext &Ctx, + size_t Idx) { + while (Idx > 0) { + const InternalDecodedInst &Prev = Ctx.Decoded[Idx - 1]; + const InternalDecodedInst &Current = Ctx.Decoded[Idx]; + if (Prev.Offset + Prev.Size != Current.Offset) + return std::nullopt; + + if (Ctx.DirectControlFlow.Targets.contains(Current.Offset)) + return std::nullopt; + for (uint64_t Entry : Ctx.DeclaredEntries) + if (Entry == Current.Offset) + return std::nullopt; + + if (std::optional Mode = getExactVgprMsbModeWritten(Prev, Ctx.LS)) + return Mode; + + if (Prev.Mnemonic == "" || + 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)))) + return std::nullopt; + --Idx; + } + return std::nullopt; +} + +unsigned getVgprMsbBank(unsigned Mode, VgprMsbOperand Operand) { + return getVgprMsbs(Mode, Operand); +} + +void setVgprMsbBank(unsigned &Mode, VgprMsbOperand Operand, unsigned Bank) { + setVgprMsbs(Mode, Operand, Bank); +} + // Return-value semantics (current shared dispatcher API in b0a0.cpp): // 0 = either "this patch did not match the instruction" OR "matched // but failed to apply" -- the dispatcher cannot distinguish the @@ -1221,8 +1261,7 @@ static uint32_t applyWmmaSplitPatchesImpl(PatchContext &Ctx, size_t Idx) { // the transition and restore it. K-splits always consult the mode (the // upper half reuses dst as src2). M-splits consult it only when a half // actually crosses v255. - if (Ctx.VgprMsbModeBefore.empty()) - computeVgprMsbModes(Ctx); + ensureVgprMsbModes(Ctx); bool UsesVgprMsbTransition = false; bool NeedsKnownVgprMsbMode = Match->Kind == SplitKind::Split128to64FP8BF8; @@ -1235,7 +1274,7 @@ static uint32_t applyWmmaSplitPatchesImpl(PatchContext &Ctx, size_t Idx) { unsigned ActiveVgprMsbMode = 0; if (NeedsKnownVgprMsbMode) { - std::optional Mode = findActiveVgprMsbMode(Ctx, Idx); + std::optional Mode = getActiveVgprMsbMode(Ctx, Idx); if (!Mode) { log() << "hotswap: error: WMMA split: cannot determine VGPR-MSB mode " "for " diff --git a/amd/comgr/test-lit/hotswap-control-flow-index-scale.s b/amd/comgr/test-lit/hotswap-control-flow-index-scale.s new file mode 100644 index 0000000000000..07417639276c7 --- /dev/null +++ b/amd/comgr/test-lit/hotswap-control-flow-index-scale.s @@ -0,0 +1,59 @@ +// COM: Exercise the public rewrite path with thousands of unrelated local +// COM: set-PC functions. Control-flow proof must remain fail-closed without +// COM: forming the Cartesian product of every set-PC and function range. + +// RUN: %clang -target amdgcn-amd-amdhsa -mcpu=gfx1250 -nostdlib %s -o %t.elf +// RUN: hotswap-rewrite %t.elf \ +// RUN: amdgcn-amd-amdhsa--gfx1250 amdgcn-amd-amdhsa--gfx1250 \ +// RUN: --output %t.out.elf \ +// RUN: | %FileCheck --check-prefix=API %s +// API: RESULT: SUCCESS +// RUN: %llvm-readelf --file-header --section-headers --program-headers \ +// RUN: %t.out.elf > /dev/null + +.amdgcn_target "amdgcn-amd-amdhsa--gfx1250" +.text + +.macro local_set_pc_function +.type local_set_pc_\@,@function +local_set_pc_\@: + s_set_pc_i64 s[0:1] +.Llocal_set_pc_end_\@: +.size local_set_pc_\@, .Llocal_set_pc_end_\@-local_set_pc_\@ +.endm + +.rept 4096 + local_set_pc_function +.endr + +.globl control_flow_index_kernel +.protected control_flow_index_kernel +.type control_flow_index_kernel,@function +control_flow_index_kernel: + s_endpgm +.Lcontrol_flow_index_kernel_end: +.size control_flow_index_kernel, .Lcontrol_flow_index_kernel_end-control_flow_index_kernel + +.rodata +.p2align 8 +.amdhsa_kernel control_flow_index_kernel + .amdhsa_next_free_vgpr 0 + .amdhsa_next_free_sgpr 0 +.end_amdhsa_kernel + +.amdgpu_metadata + amdhsa.version: + - 3 + - 0 + amdhsa.kernels: + - .name: control_flow_index_kernel + .symbol: control_flow_index_kernel.kd + .sgpr_count: 0 + .vgpr_count: 0 + .kernarg_segment_size: 0 + .group_segment_fixed_size: 0 + .private_segment_fixed_size: 0 + .kernarg_segment_align: 8 + .wavefront_size: 64 + .max_flat_workgroup_size: 256 +.end_amdgpu_metadata diff --git a/amd/comgr/test-lit/hotswap-data-only-empty-text.s b/amd/comgr/test-lit/hotswap-data-only-empty-text.s new file mode 100644 index 0000000000000..4929d793f5cd6 --- /dev/null +++ b/amd/comgr/test-lit/hotswap-data-only-empty-text.s @@ -0,0 +1,191 @@ +// COM: Data-only HIP device objects can carry global constants and variables +// COM: but no kernels or device functions. Such an object has a present, +// COM: zero-size .text section and an empty amdhsa.kernels array. There are no +// COM: instructions or kernel revision tags to transform, so return a +// COM: byte-identical successful output after validating that shape. + +// RUN: %clang -target amdgcn-amd-amdhsa -mcpu=gfx1250 -nostdlib %s -o %t.elf +// RUN: %llvm-readobj --sections --symbols --notes %t.elf \ +// RUN: | %FileCheck --check-prefix=SHAPE --implicit-check-not=.kd %s +// SHAPE: Name: .text +// SHAPE: Size: 0 +// SHAPE: Name: data_only_constant +// SHAPE: Type: Object +// SHAPE: AMDGPU Metadata: --- +// SHAPE-NEXT: amdhsa.kernels: [] + +// RUN: env AMD_COMGR_EMIT_VERBOSE_LOGS=1 hotswap-rewrite %t.elf \ +// RUN: amdgcn-amd-amdhsa--gfx1250 amdgcn-amd-amdhsa--gfx1250 \ +// RUN: --output %t.out.elf 2>&1 | %FileCheck --check-prefix=ACCEPT %s +// ACCEPT: hotswap: accepted data-only code object with empty .text; +// ACCEPT-SAME: returning a byte-identical copy. +// ACCEPT: RESULT: SUCCESS +// RUN: cmp %t.elf %t.out.elf +// RUN: hotswap-rewrite %t.out.elf \ +// RUN: amdgcn-amd-amdhsa--gfx1250 amdgcn-amd-amdhsa--gfx1250 \ +// RUN: --check-idempotent | %FileCheck --check-prefix=IDEM %s +// IDEM: IDEMPOTENT: YES + +// RUN: sed 's/^\.set claimed_function, 0$/.set claimed_function, 1/' \ +// RUN: %s > %t.function.s +// RUN: %clang -target amdgcn-amd-amdhsa -mcpu=gfx1250 -nostdlib \ +// RUN: %t.function.s -o %t.function.elf +// RUN: env AMD_COMGR_EMIT_VERBOSE_LOGS=1 hotswap-rewrite %t.function.elf \ +// RUN: amdgcn-amd-amdhsa--gfx1250 amdgcn-amd-amdhsa--gfx1250 \ +// RUN: --expect-status INVALID_ARGUMENT 2>&1 \ +// RUN: | %FileCheck --check-prefixes=FUNCTION,REJECT %s +// FUNCTION: hotswap: error: data-only object has a function/ifunc symbol +// FUNCTION-SAME: in empty .text. + +// RUN: sed 's/^\.set claimed_other_function, 0$/.set claimed_other_function, 1/' \ +// RUN: %s > %t.other-function.s +// RUN: %clang -target amdgcn-amd-amdhsa -mcpu=gfx1250 -nostdlib \ +// RUN: %t.other-function.s -o %t.other-function.elf +// RUN: env AMD_COMGR_EMIT_VERBOSE_LOGS=1 \ +// RUN: hotswap-rewrite %t.other-function.elf \ +// RUN: amdgcn-amd-amdhsa--gfx1250 amdgcn-amd-amdhsa--gfx1250 \ +// RUN: --expect-status INVALID_ARGUMENT 2>&1 \ +// RUN: | %FileCheck --check-prefixes=OTHER-FUNCTION,REJECT %s +// OTHER-FUNCTION: hotswap: error: data-only object has defined function/ifunc +// OTHER-FUNCTION-SAME: symbol 'claimed_other_function'. + +// RUN: sed 's/^\.set executable_section, 0$/.set executable_section, 1/' \ +// RUN: %s > %t.executable.s +// RUN: %clang -target amdgcn-amd-amdhsa -mcpu=gfx1250 -nostdlib \ +// RUN: %t.executable.s -o %t.executable.elf +// RUN: env AMD_COMGR_EMIT_VERBOSE_LOGS=1 hotswap-rewrite %t.executable.elf \ +// RUN: amdgcn-amd-amdhsa--gfx1250 amdgcn-amd-amdhsa--gfx1250 \ +// RUN: --expect-status INVALID_ARGUMENT 2>&1 \ +// RUN: | %FileCheck --check-prefixes=EXECUTABLE,REJECT %s +// EXECUTABLE: hotswap: error: data-only object has non-empty executable +// EXECUTABLE-SAME: section '.other_text'. + +// RUN: sed 's/^\.set claimed_descriptor, 0$/.set claimed_descriptor, 1/' \ +// RUN: %s > %t.descriptor.s +// RUN: %clang -target amdgcn-amd-amdhsa -mcpu=gfx1250 -nostdlib \ +// RUN: %t.descriptor.s -o %t.descriptor.elf +// RUN: env AMD_COMGR_EMIT_VERBOSE_LOGS=1 hotswap-rewrite %t.descriptor.elf \ +// RUN: amdgcn-amd-amdhsa--gfx1250 amdgcn-amd-amdhsa--gfx1250 \ +// RUN: --expect-status INVALID_ARGUMENT 2>&1 \ +// RUN: | %FileCheck --check-prefixes=DESCRIPTOR,REJECT %s +// DESCRIPTOR: hotswap: error: data-only object has kernel descriptor symbol +// DESCRIPTOR-SAME: 'claimed_kernel.kd'. + +// RUN: sed 's/^\.set claimed_kernel, 0$/.set claimed_kernel, 1/' \ +// RUN: %s > %t.kernel.s +// RUN: %clang -target amdgcn-amd-amdhsa -mcpu=gfx1250 -nostdlib \ +// RUN: %t.kernel.s -o %t.kernel.elf +// RUN: env AMD_COMGR_EMIT_VERBOSE_LOGS=1 hotswap-rewrite %t.kernel.elf \ +// RUN: amdgcn-amd-amdhsa--gfx1250 amdgcn-amd-amdhsa--gfx1250 \ +// RUN: --expect-status INVALID_ARGUMENT 2>&1 \ +// RUN: | %FileCheck --check-prefixes=KERNEL,REJECT %s +// KERNEL: hotswap: error: data-only AMDGPU metadata claims 1 kernel(s). + +// RUN: sed 's/^\.set malformed_metadata, 0$/.set malformed_metadata, 1/' \ +// RUN: %s > %t.malformed.s +// RUN: %clang -target amdgcn-amd-amdhsa -mcpu=gfx1250 -nostdlib \ +// RUN: %t.malformed.s -o %t.malformed.elf +// RUN: env AMD_COMGR_EMIT_VERBOSE_LOGS=1 hotswap-rewrite %t.malformed.elf \ +// RUN: amdgcn-amd-amdhsa--gfx1250 amdgcn-amd-amdhsa--gfx1250 \ +// RUN: --expect-status INVALID_ARGUMENT 2>&1 \ +// RUN: | %FileCheck --check-prefixes=MALFORMED,REJECT %s +// MALFORMED: hotswap: error: failed to parse data-only AMDGPU metadata note. + +// RUN: %llvm-objcopy --remove-section=.text %t.elf %t.missing-text.elf +// RUN: env AMD_COMGR_EMIT_VERBOSE_LOGS=1 hotswap-rewrite %t.missing-text.elf \ +// RUN: amdgcn-amd-amdhsa--gfx1250 amdgcn-amd-amdhsa--gfx1250 \ +// RUN: --expect-status INVALID_ARGUMENT 2>&1 \ +// RUN: | %FileCheck --check-prefix=MISSING %s +// MISSING: no .text section found +// MISSING: RESULT: INVALID_ARGUMENT + +// REJECT: hotswap: error: retargetCodeObject: +// REJECT-SAME: does not describe a valid data-only code object. +// REJECT: RESULT: INVALID_ARGUMENT + +.set claimed_function, 0 +.set claimed_other_function, 0 +.set executable_section, 0 +.set claimed_descriptor, 0 +.set claimed_kernel, 0 +.set malformed_metadata, 0 + +.amdgcn_target "amdgcn-amd-amdhsa--gfx1250" +.text +.if claimed_function +.globl claimed_function +.type claimed_function,@function +claimed_function: +.size claimed_function, .-claimed_function +.endif + +.rodata +.globl data_only_constant +.type data_only_constant,@object +.p2align 2 +data_only_constant: + .long 1 +.size data_only_constant, .-data_only_constant + +.if claimed_other_function +.globl claimed_other_function +.type claimed_other_function,@function +claimed_other_function: +.size claimed_other_function, .-claimed_other_function +.endif + +.if executable_section +.section .other_text,"ax",@progbits + v_nop +.endif + +.if claimed_descriptor +.globl claimed_kernel.kd +.type claimed_kernel.kd,@object +.p2align 6 +claimed_kernel.kd: + .zero 64 +.size claimed_kernel.kd, .-claimed_kernel.kd +.endif + +.if malformed_metadata +.section .note,"a",@note +.p2align 2 + .long 7 + .long 4 + .long 32 + .asciz "AMDGPU" +.p2align 2 + .byte 0xc1, 0xc1, 0xc1, 0xc1 +.p2align 2 +.else +.section .note,"a",@note +.p2align 2 + .long 7 + .long .Lmetadata_desc_end-.Lmetadata_desc_begin + .long 32 + .asciz "AMDGPU" +.p2align 2 +.Lmetadata_desc_begin: +.if claimed_kernel + // {"amdhsa.kernels": [{}]} + .byte 0x81, 0xae + .ascii "amdhsa.kernels" + .byte 0x91, 0x80 +.else + // Match the corpus note: + // {"amdhsa.kernels": [], "amdhsa.target": "...gfx1250", + // "amdhsa.version": [1, 2]} + .byte 0x83, 0xae + .ascii "amdhsa.kernels" + .byte 0x90, 0xad + .ascii "amdhsa.target" + .byte 0xba + .ascii "amdgcn-amd-amdhsa--gfx1250" + .byte 0xae + .ascii "amdhsa.version" + .byte 0x92, 0x01, 0x02 +.endif +.Lmetadata_desc_end: +.p2align 2 +.endif diff --git a/amd/comgr/test-lit/hotswap-reusable-pc-call-targets.s b/amd/comgr/test-lit/hotswap-reusable-pc-call-targets.s new file mode 100644 index 0000000000000..6d7c8d2b3224c --- /dev/null +++ b/amd/comgr/test-lit/hotswap-reusable-pc-call-targets.s @@ -0,0 +1,185 @@ +// COM: Production activation kernels select one of several local callees with +// COM: get-PC/carry materialization, merge the selected address, and reuse it +// COM: across many register calls. Resolve the finite reaching-target set so +// COM: an unrelated required far rewrite may safely use external gateway +// COM: padding. A selector bypass would leave the target unknown and must +// COM: continue to fail closed. + +// 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 amdgcn-amd-amdhsa--gfx1250 \ +// RUN: --output %t.out.elf 2>&1 | %FileCheck --check-prefix=LOG %s +// LOG: hotswap: resolved reusable PC-materialized call +// LOG-SAME: to 2 target(s) +// LOG-NOT: hotswap: unresolved call target +// LOG: hotswap: planned 1 shared far-dispatch gateway group(s) for 8 source site(s) +// LOG: RESULT: SUCCESS + +// RUN: sed 's/^\.set unsafe_selector, 0$/.set unsafe_selector, 1/' \ +// RUN: %s > %t.bypass.s +// RUN: %clang -target amdgcn-amd-amdhsa -mcpu=gfx1250 -nostdlib \ +// RUN: %t.bypass.s -o %t.bypass.elf +// RUN: env AMD_COMGR_EMIT_VERBOSE_LOGS=1 hotswap-rewrite %t.bypass.elf \ +// RUN: amdgcn-amd-amdhsa--gfx1250 amdgcn-amd-amdhsa--gfx1250 \ +// RUN: --expect-status ERROR 2>&1 \ +// RUN: | %FileCheck --check-prefixes=BYPASS,FAIL %s +// BYPASS: hotswap: unresolved call target +// FAIL: hotswap: unresolved control-flow target disables NOP-sled emission, +// FAIL-SAME: trampoline coalescing, source relocation, and .text gateways +// FAIL: hotswap: error: no safe short-branch gateway for far site +// FAIL: RESULT: ERROR + +// RUN: sed 's/^\.set outside_selector, 0$/.set outside_selector, 1/' \ +// RUN: %s > %t.outside.s +// RUN: %clang -target amdgcn-amd-amdhsa -mcpu=gfx1250 -nostdlib \ +// RUN: %t.outside.s -o %t.outside.elf +// RUN: env AMD_COMGR_EMIT_VERBOSE_LOGS=1 hotswap-rewrite %t.outside.elf \ +// RUN: amdgcn-amd-amdhsa--gfx1250 amdgcn-amd-amdhsa--gfx1250 \ +// RUN: --expect-status ERROR 2>&1 \ +// RUN: | %FileCheck --check-prefix=OUTSIDE %s +// OUTSIDE: hotswap: unresolved call target +// OUTSIDE-SAME: (reusable target outside .text) +// OUTSIDE: hotswap: unresolved control-flow target disables NOP-sled emission, +// OUTSIDE-SAME: trampoline coalescing, source relocation, and .text gateways +// OUTSIDE: hotswap: error: no safe short-branch gateway for far site +// OUTSIDE: RESULT: ERROR + +// RUN: %llvm-objdump -d %t.out.elf | %FileCheck --check-prefix=DISASM \ +// RUN: --implicit-check-not=s_add_pc_i64 %s +// DISASM-LABEL: : +// DISASM: s_swap_pc_i64 +// DISASM: s_get_pc_i64 +// DISASM: s_branch +// DISASM: s_get_pc_i64 +// DISASM-NEXT: s_branch +// DISASM-LABEL: : +// DISASM-NEXT: s_endpgm +// DISASM-NEXT: s_get_pc_i64 +// DISASM-NEXT: s_add_nc_u64 +// DISASM-NEXT: s_set_pc_i64 +// DISASM: ds_load_b32 v0, v2 offset:256 +// DISASM-NEXT: ds_load_b32 v1, v2 offset:768 + +// RUN: hotswap-rewrite %t.out.elf \ +// RUN: amdgcn-amd-amdhsa--gfx1250 amdgcn-amd-amdhsa--gfx1250 \ +// RUN: --check-idempotent | %FileCheck --check-prefix=IDEM %s +// IDEM: IDEMPOTENT: YES + +.set unsafe_selector, 0 +.set outside_selector, 0 + +.amdgcn_target "amdgcn-amd-amdhsa--gfx1250" +.text +.globl reusable_pc_targets +.p2align 8 +.type reusable_pc_targets,@function +reusable_pc_targets: +.if unsafe_selector + // This edge reaches the call without executing either get-PC sequence. + s_cmp_eq_u32 s0, 2 + s_cbranch_scc1 .Lselected +.endif +.if outside_selector + s_cmp_eq_u32 s0, 3 + s_cbranch_scc1 .Lselect_outside +.endif + s_cmp_eq_u32 s0, 0 + s_cbranch_scc1 .Lselect_second +.Lselect_first: + s_get_pc_i64 s[2:3] + s_add_co_i32 s4, callee_first-(.Lselect_first+4)-4, 4 + s_add_co_u32 s2, s2, s4 + s_add_co_ci_u32 s3, s3, 0 + s_branch .Lselected +.Lselect_second: + s_get_pc_i64 s[2:3] + s_add_co_i32 s4, callee_second-(.Lselect_second+4)-4, 4 + s_add_co_u32 s2, s2, s4 + s_add_co_ci_u32 s3, s3, 0 +.if outside_selector + s_branch .Lselected +.Lselect_outside: + s_get_pc_i64 s[2:3] + s_add_co_i32 s4, outside_text_end-(.Lselect_outside+4)-4, 4 + s_add_co_u32 s2, s2, s4 + s_add_co_ci_u32 s3, s3, 0 +.endif +.Lselected: + s_swap_pc_i64 s[6:7], s[2:3] + s_branch .Lpatch0 +.Lpatch0: + ds_load_2addr_stride64_b32 v[0:1], v2 offset0:1 offset1:3 + s_branch .Lpatch1 +.Lpatch1: + ds_load_2addr_stride64_b32 v[0:1], v2 offset0:1 offset1:3 + s_branch .Lpatch2 +.Lpatch2: + ds_load_2addr_stride64_b32 v[0:1], v2 offset0:1 offset1:3 + s_branch .Lpatch3 +.Lpatch3: + ds_load_2addr_stride64_b32 v[0:1], v2 offset0:1 offset1:3 + s_branch .Lpatch4 +.Lpatch4: + ds_load_2addr_stride64_b32 v[0:1], v2 offset0:1 offset1:3 + s_branch .Lpatch5 +.Lpatch5: + ds_load_2addr_stride64_b32 v[0:1], v2 offset0:1 offset1:3 + s_branch .Lpatch6 +.Lpatch6: + ds_load_2addr_stride64_b32 v[0:1], v2 offset0:1 offset1:3 + s_branch .Lpatch7 +.Lpatch7: + ds_load_2addr_stride64_b32 v[0:1], v2 offset0:1 offset1:3 +.Lpatch_done: + s_wait_dscnt 0x0 + s_endpgm +.size reusable_pc_targets, .-reusable_pc_targets + +.local callee_first +.type callee_first,@function +callee_first: + s_mov_b32 s8, 1 + s_set_pc_i64 s[6:7] +.size callee_first, .-callee_first + +.local callee_second +.type callee_second,@function +callee_second: + s_mov_b32 s8, 2 + s_set_pc_i64 s[6:7] +.size callee_second, .-callee_second + +.type gateway_barrier,@function +gateway_barrier: + s_endpgm +.size gateway_barrier, .-gateway_barrier +.fill 20, 1, 0 + +.rept 40000 + s_mov_b32 s10, s11 +.endr + +outside_text_end: +.rodata +.p2align 8 +.amdhsa_kernel reusable_pc_targets + .amdhsa_next_free_vgpr 3 + .amdhsa_next_free_sgpr 12 +.end_amdhsa_kernel + +.amdgpu_metadata + amdhsa.version: + - 3 + - 0 + amdhsa.kernels: + - .name: reusable_pc_targets + .symbol: reusable_pc_targets.kd + .sgpr_count: 12 + .vgpr_count: 3 + .kernarg_segment_size: 0 + .group_segment_fixed_size: 0 + .private_segment_fixed_size: 0 + .kernarg_segment_align: 8 + .wavefront_size: 64 + .max_flat_workgroup_size: 256 +.end_amdgpu_metadata diff --git a/amd/comgr/test-lit/hotswap-trampoline-ds-vgpr-msb.s b/amd/comgr/test-lit/hotswap-trampoline-ds-vgpr-msb.s new file mode 100644 index 0000000000000..6a925ecd93395 --- /dev/null +++ b/amd/comgr/test-lit/hotswap-trampoline-ds-vgpr-msb.s @@ -0,0 +1,62 @@ +// Verify that splitting a two-address DS load whose second destination crosses +// v255 rebases that half to v0 under a temporary destination VGPR-MSB mode. + +// 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 amdgcn-amd-amdhsa--gfx1250 \ +// RUN: --output %t.out.elf 2>&1 | %FileCheck --check-prefix=API %s +// API-NOT: error: +// API: RESULT: SUCCESS + +// RUN: %llvm-objdump -d %t.out.elf | %FileCheck --check-prefix=DISASM %s +// DISASM-LABEL: : +// DISASM: s_branch +// DISASM: ds_load_b64 v[254:255], v88 offset:680 +// DISASM-NEXT: s_set_vgpr_msb 64 +// DISASM-NEXT: ds_load_b64 v[0:1]{{.*v\[256:257\].*}}v88 offset:688 +// DISASM-NEXT: s_set_vgpr_msb 0x4000 +// DISASM-NEXT: s_wait_dscnt 0x0 +// DISASM: s_branch + +// RUN: hotswap-rewrite %t.out.elf \ +// RUN: amdgcn-amd-amdhsa--gfx1250 amdgcn-amd-amdhsa--gfx1250 \ +// RUN: --check-idempotent | %FileCheck --check-prefix=IDEM %s +// IDEM: IDEMPOTENT: YES + +.amdgcn_target "amdgcn-amd-amdhsa--gfx1250" +.text +.globl test_ds_vgpr_msb +.p2align 8 +.type test_ds_vgpr_msb,@function +test_ds_vgpr_msb: + s_set_vgpr_msb 0 + ds_load_2addr_b64 v[254:257], v88 offset0:85 offset1:86 + s_wait_dscnt 0x0 + s_endpgm +.Ltest_ds_vgpr_msb_end: +.size test_ds_vgpr_msb, .Ltest_ds_vgpr_msb_end-test_ds_vgpr_msb + +.rodata +.p2align 8 +.amdhsa_kernel test_ds_vgpr_msb + .amdhsa_next_free_vgpr 688 + .amdhsa_next_free_sgpr 2 + .amdhsa_wavefront_size32 1 +.end_amdhsa_kernel + +.amdgpu_metadata + amdhsa.version: + - 3 + - 0 + amdhsa.kernels: + - .name: test_ds_vgpr_msb + .symbol: test_ds_vgpr_msb.kd + .sgpr_count: 2 + .vgpr_count: 688 + .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-lit/hotswap-trampoline-long-branch.s b/amd/comgr/test-lit/hotswap-trampoline-long-branch.s index 6371b3f24f0fc..acc238e2e61c5 100644 --- a/amd/comgr/test-lit/hotswap-trampoline-long-branch.s +++ b/amd/comgr/test-lit/hotswap-trampoline-long-branch.s @@ -38,17 +38,84 @@ // RUN: | %FileCheck --check-prefix=IDEM %s // IDEM: IDEMPOTENT: YES -// COM: A kernel with no aligned two-SGPR block fails closed instead of -// COM: emitting the unsafe return or clobbering a live register. -// RUN: sed -e 's/s_mov_b64 vcc, -1/s_mov_b32 s105, 0/' \ -// RUN: -e 's/\.amdhsa_next_free_sgpr 12/.amdhsa_next_free_sgpr 106/' \ -// RUN: -e 's/\.sgpr_count: 14/.sgpr_count: 106/' %s > %t.full-sgpr.s +// COM: A kernel with no aligned numbered SGPR pair can reuse VCC when the +// COM: replacement does not consume its incoming value and the continuation +// COM: ends before reading it. +// RUN: sed -e 's/s_mov_b64 vcc, -1/s_mov_b32 s104, 0/' \ +// RUN: -e 's/\.amdhsa_next_free_sgpr 12/.amdhsa_next_free_sgpr 105/' \ +// RUN: -e 's/\.sgpr_count: 14/.sgpr_count: 105/' %s > %t.full-sgpr.s // RUN: %clang -target amdgcn-amd-amdhsa -mcpu=gfx1250 -nostdlib \ // RUN: %t.full-sgpr.s -o %t.full-sgpr.elf -// RUN: hotswap-rewrite %t.full-sgpr.elf \ +// RUN: env AMD_COMGR_EMIT_VERBOSE_LOGS=1 hotswap-rewrite %t.full-sgpr.elf \ // RUN: amdgcn-amd-amdhsa--gfx1250 amdgcn-amd-amdhsa--gfx1250 \ -// RUN: --expect-status ERROR | %FileCheck --check-prefix=FAIL %s -// FAIL: RESULT: ERROR +// RUN: --output %t.full-sgpr.out.elf 2>&1 \ +// RUN: | %FileCheck --check-prefix=FULL-LOG %s +// FULL-LOG: hotswap: safe far return: reusing dead VCC +// FULL-LOG: RESULT: SUCCESS +// RUN: %llvm-objdump -d %t.full-sgpr.out.elf \ +// RUN: | %FileCheck --check-prefix=FULL-DISASM %s +// FULL-DISASM: s_get_pc_i64 vcc +// FULL-DISASM-NEXT: s_add_nc_u64 vcc, vcc, +// FULL-DISASM-NEXT: s_set_pc_i64 vcc +// FULL-DISASM: s_get_pc_i64 vcc +// FULL-DISASM-NEXT: s_add_nc_u64 vcc, vcc, +// FULL-DISASM-NEXT: s_set_pc_i64 vcc + +// COM: A live VCC can instead use an already-allocated numbered pair after +// COM: CFG liveness proves that neither half is consumed by the replacement or +// COM: continuation before being redefined. +// RUN: sed '/^ tensor_load_to_lds/a\ s_cbranch_vccz 0' %t.full-sgpr.s \ +// RUN: > %t.local-pair.s +// RUN: %clang -target amdgcn-amd-amdhsa -mcpu=gfx1250 -nostdlib \ +// RUN: %t.local-pair.s -o %t.local-pair.elf +// RUN: env AMD_COMGR_EMIT_VERBOSE_LOGS=1 hotswap-rewrite %t.local-pair.elf \ +// RUN: amdgcn-amd-amdhsa--gfx1250 amdgcn-amd-amdhsa--gfx1250 \ +// RUN: --output %t.local-pair.out.elf 2>&1 \ +// RUN: | %FileCheck --check-prefix=LOCAL-PAIR-LOG %s +// LOCAL-PAIR-LOG: hotswap: safe far return: reusing locally dead s[104:105] +// LOCAL-PAIR-LOG: RESULT: SUCCESS + +// COM: Search every aligned pair, not just the highest eight. Every pair above +// COM: s[30:31] has a reachable incoming-value read; s[30:31] is overwritten +// COM: first and is therefore the highest locally dead pair. +// RUN: sed -e '/^ tensor_load_to_lds/a\ s_cbranch_vccz 0' \ +// RUN: -e 's|^// LOW-PAIR-ONLY:| |' %t.full-sgpr.s > %t.low-pair.s +// RUN: %clang -target amdgcn-amd-amdhsa -mcpu=gfx1250 -nostdlib \ +// RUN: %t.low-pair.s -o %t.low-pair.elf +// RUN: env AMD_COMGR_EMIT_VERBOSE_LOGS=1 hotswap-rewrite %t.low-pair.elf \ +// RUN: amdgcn-amd-amdhsa--gfx1250 amdgcn-amd-amdhsa--gfx1250 \ +// RUN: --output %t.low-pair.out.elf 2>&1 \ +// RUN: | %FileCheck --check-prefix=LOW-PAIR-LOG %s +// LOW-PAIR-LOG: hotswap: safe far return: reusing locally dead s[30:31] +// LOW-PAIR-LOG: RESULT: SUCCESS +// RUN: %llvm-objdump -d %t.low-pair.out.elf \ +// RUN: | %FileCheck --check-prefix=LOW-PAIR-DISASM %s +// LOW-PAIR-DISASM: s_get_pc_i64 s[30:31] +// LOW-PAIR-DISASM-NEXT: s_add_nc_u64 s[30:31], s[30:31], +// LOW-PAIR-DISASM-NEXT: s_set_pc_i64 s[30:31] + +// COM: When the continuation reads VCC before redefining it, a wave32 rewrite +// COM: preserves VCC_LO in the one remaining numbered SGPR. The source reaches +// COM: a save/set-PC gateway, and its tail becomes the restore landing pad. +// RUN: sed 's|^// LIVE-ONLY:| |' %t.full-sgpr.s \ +// RUN: > %t.live-vcc.s +// RUN: %clang -target amdgcn-amd-amdhsa -mcpu=gfx1250 -nostdlib \ +// RUN: %t.live-vcc.s -o %t.live-vcc.elf +// RUN: env AMD_COMGR_EMIT_VERBOSE_LOGS=1 hotswap-rewrite %t.live-vcc.elf \ +// RUN: amdgcn-amd-amdhsa--gfx1250 amdgcn-amd-amdhsa--gfx1250 \ +// RUN: --output %t.live-vcc.out.elf 2>&1 \ +// RUN: | %FileCheck --check-prefix=LIVE-LOG %s +// LIVE-LOG: hotswap: safe far return: preserving live wave32 VCC_LO in s105 +// LIVE-LOG: hotswap: assigned 1 SCC-neutral forward gateway(s) +// LIVE-LOG: RESULT: SUCCESS +// RUN: %llvm-objdump -d %t.live-vcc.out.elf \ +// RUN: | %FileCheck --check-prefix=LIVE-DISASM %s +// LIVE-DISASM-LABEL: : +// LIVE-DISASM: s_branch +// LIVE-DISASM-NEXT: s_mov_b32 vcc_lo, s105 +// LIVE-DISASM: s_cbranch_vccz +// LIVE-DISASM: s_mov_b32 s105, vcc_lo +// LIVE-DISASM-NEXT: s_get_pc_i64 vcc // COM: A metadata-less object also fails closed because scratch usage cannot // COM: be charged to its owning kernel. @@ -58,6 +125,7 @@ // RUN: hotswap-rewrite %t.nometa.elf \ // RUN: amdgcn-amd-amdhsa--gfx1250 amdgcn-amd-amdhsa--gfx1250 \ // RUN: --expect-status ERROR | %FileCheck --check-prefix=FAIL %s +// FAIL: RESULT: ERROR .amdgcn_target "amdgcn-amd-amdhsa--gfx1250" .text @@ -67,6 +135,14 @@ test_far: s_mov_b64 vcc, -1 tensor_load_to_lds s[0:3], s[4:11] +// LOW-PAIR-ONLY:s_mov_b64 s[30:31], 0 +// LOW-PAIR-ONLY:.irp live_reg, s32, s34, s36, s38, s40, s42, s44, s46, s48, s50, s52, s54, s56, s58, s60, s62, s64, s66, s68, s70, s72, s74, s76, s78, s80, s82, s84, s86, s88, s90, s92, s94, s96, s98, s100, s102, s104 +// LOW-PAIR-ONLY:s_mov_b32 s1, \live_reg +// LOW-PAIR-ONLY:.endr +// LIVE-ONLY:s_cbranch_vccz 0 +// LIVE-ONLY:.irp live_reg, s0, s2, s4, s6, s8, s10, s12, s14, s16, s18, s20, s22, s24, s26, s28, s30, s32, s34, s36, s38, s40, s42, s44, s46, s48, s50, s52, s54, s56, s58, s60, s62, s64, s66, s68, s70, s72, s74, s76, s78, s80, s82, s84, s86, s88, s90, s92, s94, s96, s98, s100, s102, s104 +// LIVE-ONLY:s_mov_b32 s1, \live_reg +// LIVE-ONLY:.endr s_endpgm .size test_far, .-test_far @@ -81,7 +157,19 @@ gateway_barrier: // ~160 KB of non-NOP filler so the appended trampoline pool is beyond // s_branch's +-128 KB reach from the tensor_load above (forces the // long-branch path). - .rept 40000 + .rept 20000 + s_mov_b32 s0, s1 + .endr + +// A safe midpoint sled gives a registerless far edge an s_branch island in +// each direction. +.type midpoint_gateway_barrier,@function +midpoint_gateway_barrier: + s_endpgm +.size midpoint_gateway_barrier, .-midpoint_gateway_barrier +.fill 32, 1, 0 + + .rept 20000 s_mov_b32 s0, s1 .endr .Ltest_far_end: @@ -91,6 +179,7 @@ gateway_barrier: .amdhsa_kernel test_far .amdhsa_next_free_vgpr 1 .amdhsa_next_free_sgpr 12 + .amdhsa_wavefront_size32 1 .end_amdhsa_kernel .amdgpu_metadata @@ -106,6 +195,6 @@ gateway_barrier: .group_segment_fixed_size: 0 .private_segment_fixed_size: 0 .kernarg_segment_align: 8 - .wavefront_size: 64 + .wavefront_size: 32 .max_flat_workgroup_size: 256 .end_amdgpu_metadata diff --git a/amd/comgr/test-lit/hotswap-trampoline-source-tail-islands.s b/amd/comgr/test-lit/hotswap-trampoline-source-tail-islands.s new file mode 100644 index 0000000000000..8b2221833a07f --- /dev/null +++ b/amd/comgr/test-lit/hotswap-trampoline-source-tail-islands.s @@ -0,0 +1,86 @@ +// COM: Far eight-byte patch sites have no padding inside their tiny owning +// COM: functions. Their unreachable second dwords form a registerless relay +// COM: chain to the appended trampoline pool. + +// 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 amdgcn-amd-amdhsa--gfx1250 \ +// RUN: --output %t.out.elf 2>&1 | %FileCheck --check-prefix=LOG %s +// LOG: hotswap: assigned 1 forward s_branch island chain(s) +// LOG: RESULT: SUCCESS + +// RUN: %llvm-objdump -d %t.out.elf | %FileCheck --check-prefix=DISASM %s +// DISASM-LABEL: : +// DISASM-NEXT: s_branch +// DISASM-NEXT: s_nop +// DISASM-LABEL: : +// DISASM-NEXT: s_branch +// DISASM-NEXT: s_branch + +// RUN: hotswap-rewrite %t.out.elf \ +// RUN: amdgcn-amd-amdhsa--gfx1250 amdgcn-amd-amdhsa--gfx1250 \ +// RUN: --check-idempotent | %FileCheck --check-prefix=IDEM %s +// IDEM: IDEMPOTENT: YES + +.amdgcn_target "amdgcn-amd-amdhsa--gfx1250" +.text +.globl source0 +.p2align 8 +.type source0,@function +source0: + ds_load_2addr_stride64_b32 v[0:1], v2 offset0:1 offset1:3 + s_endpgm +.size source0, .-source0 + +.rept 25000 + s_mov_b32 s0, s1 +.endr + +.globl source1 +.type source1,@function +source1: + ds_load_2addr_stride64_b32 v[0:1], v2 offset0:1 offset1:3 + s_endpgm +.size source1, .-source1 + +.rept 12500 + s_mov_b32 s0, s1 +.endr + +.rodata +.p2align 8 +.amdhsa_kernel source0 + .amdhsa_next_free_vgpr 3 + .amdhsa_next_free_sgpr 12 +.end_amdhsa_kernel +.amdhsa_kernel source1 + .amdhsa_next_free_vgpr 3 + .amdhsa_next_free_sgpr 12 +.end_amdhsa_kernel + +.amdgpu_metadata + amdhsa.version: + - 3 + - 0 + amdhsa.kernels: + - .name: source0 + .symbol: source0.kd + .sgpr_count: 14 + .vgpr_count: 3 + .kernarg_segment_size: 0 + .group_segment_fixed_size: 0 + .private_segment_fixed_size: 0 + .kernarg_segment_align: 8 + .wavefront_size: 64 + .max_flat_workgroup_size: 256 + - .name: source1 + .symbol: source1.kd + .sgpr_count: 14 + .vgpr_count: 3 + .kernarg_segment_size: 0 + .group_segment_fixed_size: 0 + .private_segment_fixed_size: 0 + .kernarg_segment_align: 8 + .wavefront_size: 64 + .max_flat_workgroup_size: 256 +.end_amdgpu_metadata diff --git a/amd/comgr/test-lit/hotswap-wmma-scale16-fp4.s b/amd/comgr/test-lit/hotswap-wmma-scale16-fp4.s index ad9608cd485d9..547117b26769f 100644 --- a/amd/comgr/test-lit/hotswap-wmma-scale16-fp4.s +++ b/amd/comgr/test-lit/hotswap-wmma-scale16-fp4.s @@ -17,14 +17,14 @@ // COM: VGPR select, never a lane mask, so no v_cndmask appears in either pass. // DISASM-NOT: v_cndmask_b32_e64 // DISASM: v_mov_b32{{(_e32)?}} v{{[0-9]+}}, 0 -// DISASM: v_wmma_scale_f32_16x16x128_f8f6f4 v[0:7], v[{{[0-9]+}}:{{[0-9]+}}], v[32:39], v[0:7],{{.*}}matrix_a_fmt:MATRIX_FMT_FP4 +// DISASM: v_wmma_scale_f32_16x16x128_f8f6f4 v[0:7], v[{{[0-9]+}}:{{[0-9]+}}], v[{{[0-9]+}}:{{[0-9]+}}], v[0:7],{{.*}}matrix_a_fmt:MATRIX_FMT_FP4 // COM: exactly one gfx1250 hazard v_nop before the pass-high masked-A VALU. // DISASM-COUNT-1: v_nop // DISASM-NEXT: v_mov_b32{{(_e32)?}} v{{[0-9]+}}, 0 // DISASM-NOT: v_cndmask_b32_e64 // COM: pass-high keeps the high-16 subblock VGPRs, odd-byte scale gather, // COM: accumulating onto pass-low through v[0:7]. -// DISASM: v_wmma_scale_f32_16x16x128_f8f6f4 v[0:7], v[{{[0-9]+}}:{{[0-9]+}}], v[32:39], v[0:7],{{.*}}matrix_a_fmt:MATRIX_FMT_FP4 +// DISASM: v_wmma_scale_f32_16x16x128_f8f6f4 v[0:7], v[{{[0-9]+}}:{{[0-9]+}}], v[{{[0-9]+}}:{{[0-9]+}}], v[0:7],{{.*}}matrix_a_fmt:MATRIX_FMT_FP4 .amdgcn_target "amdgcn-amd-amdhsa--gfx1250" .text diff --git a/amd/comgr/test-lit/hotswap-wmma-scale16-large-vgpr-count.s b/amd/comgr/test-lit/hotswap-wmma-scale16-large-vgpr-count.s new file mode 100644 index 0000000000000..d2f47c551058b --- /dev/null +++ b/amd/comgr/test-lit/hotswap-wmma-scale16-large-vgpr-count.s @@ -0,0 +1,75 @@ +// A kernel may allocate more than 256 physical VGPRs through gfx1250's +// VGPR-MSB mode. The scale16 lowering's generated assembly must still use +// encodable v0-v255 names and select its above-KD scratch bank explicitly. +// The bump is occupancy-neutral (the original and rewritten allocations both +// admit one wave/EU), despite the maximum-workgroup metadata asking for two. + +// 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 amdgcn-amd-amdhsa--gfx1250 \ +// RUN: --output %t.out.elf 2>&1 | %FileCheck --check-prefix=API %s +// API: wmma_scale16: exact K-split +// API-NOT: register index is out of range +// API-NOT: error: +// API: RESULT: SUCCESS + +// RUN: %llvm-objdump -d %t.out.elf | %FileCheck --check-prefix=DISASM %s +// DISASM-LABEL: : +// DISASM-NOT: v_wmma_scale16 +// DISASM: v_mov_b32_e32 v191 /*v703*/, v255 +// DISASM-NEXT: s_set_vgpr_msb 0xa0a1 +// DISASM-NEXT: v_mov_b32_e32 v192 /*v704*/, v0 /*v256*/ +// DISASM: s_set_vgpr_msb 0xa00a +// DISASM-NEXT: v_wmma_scale_f32_16x16x128_f8f6f4 v[38:45], v[182:189] /*v[694:701]*/, v[190:197] /*v[702:709]*/, 0, +// DISASM: s_set_vgpr_msb 0x880a +// DISASM-NEXT: v_wmma_scale_f32_16x16x128_f8f6f4 v[38:45], v[182:189] /*v[694:701]*/, v[190:197] /*v[702:709]*/, v[38:45], +// DISASM-NEXT: s_set_vgpr_msb 0xa00 +// DISASM: s_set_vgpr_msb 0xa00a +// DISASM-NEXT: v_wmma_scale_f32_16x16x128_f8f6f4 v[38:45], v[182:189] /*v[694:701]*/, v[190:197] /*v[702:709]*/, 0, +// DISASM: s_set_vgpr_msb 0x880a +// DISASM-NEXT: v_wmma_scale_f32_16x16x128_f8f6f4 v[38:45], v[182:189] /*v[694:701]*/, v[190:197] /*v[702:709]*/, v[38:45], +// DISASM-NEXT: s_set_vgpr_msb 0xa00 + +// RUN: hotswap-rewrite %t.out.elf \ +// RUN: amdgcn-amd-amdhsa--gfx1250 amdgcn-amd-amdhsa--gfx1250 \ +// RUN: --check-idempotent | %FileCheck --check-prefix=IDEM %s +// IDEM: IDEMPOTENT: YES + +.amdgcn_target "amdgcn-amd-amdhsa--gfx1250" +.text +.globl test_wmma_scale16_large_vgpr_count +.p2align 8 +.type test_wmma_scale16_large_vgpr_count,@function +test_wmma_scale16_large_vgpr_count: + s_set_vgpr_msb 0x100 + v_mov_b32 v10, v10 + v_wmma_scale16_f32_16x16x128_f8f6f4 v[38:45], v[174:181], v[254:261], 0, v[0:1], v[18:19] matrix_a_fmt:MATRIX_FMT_FP4 matrix_b_fmt:MATRIX_FMT_FP4 matrix_a_scale_fmt:MATRIX_SCALE_FMT_E4M3 matrix_b_scale_fmt:MATRIX_SCALE_FMT_E4M3 + v_wmma_scale16_f32_16x16x128_f8f6f4 v[38:45], v[174:181], v[240:247], 0, v[0:1], v[18:19] matrix_a_fmt:MATRIX_FMT_FP4 matrix_b_fmt:MATRIX_FMT_FP4 matrix_a_scale_fmt:MATRIX_SCALE_FMT_E4M3 matrix_b_scale_fmt:MATRIX_SCALE_FMT_E4M3 + s_endpgm +.Ltest_wmma_scale16_large_vgpr_count_end: +.size test_wmma_scale16_large_vgpr_count, .Ltest_wmma_scale16_large_vgpr_count_end-test_wmma_scale16_large_vgpr_count + +.rodata +.p2align 8 +.amdhsa_kernel test_wmma_scale16_large_vgpr_count + .amdhsa_next_free_vgpr 688 + .amdhsa_next_free_sgpr 2 + .amdhsa_wavefront_size32 1 +.end_amdhsa_kernel + +.amdgpu_metadata + amdhsa.version: + - 3 + - 0 + amdhsa.kernels: + - .name: test_wmma_scale16_large_vgpr_count + .symbol: test_wmma_scale16_large_vgpr_count.kd + .sgpr_count: 2 + .vgpr_count: 688 + .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-lit/hotswap-wmma-scale16.s b/amd/comgr/test-lit/hotswap-wmma-scale16.s index 6bcc9857bbcdf..e7dca7dbab363 100644 --- a/amd/comgr/test-lit/hotswap-wmma-scale16.s +++ b/amd/comgr/test-lit/hotswap-wmma-scale16.s @@ -1,8 +1,9 @@ // COM: The target gfx1250 has no block-16 scaled WMMA, so the rewrite lowers // COM: 16x16x128_f8f6f4 exactly into two block-32 WMMAs chained through the // COM: accumulator: each pass masks matrix A to one 16-K subblock (a lane mask -// COM: for FP8) and byte-gathers the matching block-16 scales. When the scratch -// COM: budget is unavailable it fails closed (see the 32x16 refuse test). +// COM: for FP8), copies matrix B into the same scratch bank, and byte-gathers +// COM: the matching block-16 scales. When the scratch budget is unavailable it +// COM: fails closed (see the 32x16 refuse test). // RUN: %clang -target amdgcn-amd-amdhsa -mcpu=gfx1250 -nostdlib %s -o %t.elf // RUN: hotswap-rewrite %t.elf \ @@ -18,13 +19,13 @@ // DISASM-NOT: v_wmma_scale16 // DISASM: s_mov_b32 s{{[0-9]+}}, 0xffff {{.*$}} // DISASM: v_cndmask_b32_e64 -// DISASM: v_wmma_scale_f32_16x16x128_f8f6f4 v[0:7], v[{{[0-9]+}}:{{[0-9]+}}], v[32:47], v[0:7], +// DISASM: v_wmma_scale_f32_16x16x128_f8f6f4 v[0:7], v[{{[0-9]+}}:{{[0-9]+}}], v[{{[0-9]+}}:{{[0-9]+}}], v[0:7], // COM: exactly one gfx1250 hazard v_nop before the pass-high lane-mask VALU that // COM: overwrites the masked-A scratch block. // DISASM-COUNT-1: v_nop // DISASM-NEXT: s_mov_b32 s{{[0-9]+}}, 0xffff0000 // DISASM: v_cndmask_b32_e64 -// DISASM: v_wmma_scale_f32_16x16x128_f8f6f4 v[0:7], v[{{[0-9]+}}:{{[0-9]+}}], v[32:47], v[0:7], +// DISASM: v_wmma_scale_f32_16x16x128_f8f6f4 v[0:7], v[{{[0-9]+}}:{{[0-9]+}}], v[{{[0-9]+}}:{{[0-9]+}}], v[0:7], .amdgcn_target "amdgcn-amd-amdhsa--gfx1250" .text diff --git a/amd/comgr/test-lit/lit.cfg.py b/amd/comgr/test-lit/lit.cfg.py index 9a04bbd6296db..c2cb1f386abdf 100644 --- a/amd/comgr/test-lit/lit.cfg.py +++ b/amd/comgr/test-lit/lit.cfg.py @@ -118,6 +118,9 @@ def _fwd(*parts): config.substitutions.append( ("%llvm-readelf", _fwd(config.llvm_tools_dir, "llvm-readelf")) ) +config.substitutions.append( + ("%llvm-readobj", _fwd(config.llvm_tools_dir, "llvm-readobj")) +) config.substitutions.append(("%ld.lld", _fwd(config.llvm_tools_dir, "ld.lld"))) config.substitutions.append(("%yaml2obj", _fwd(config.llvm_tools_dir, "yaml2obj"))) config.substitutions.append(("%FileCheck", _fwd(config.llvm_tools_dir, "FileCheck"))) diff --git a/amd/comgr/test-unit/HotswapMCTest.cpp b/amd/comgr/test-unit/HotswapMCTest.cpp index 5aa4120d9b9d5..29490a9933718 100644 --- a/amd/comgr/test-unit/HotswapMCTest.cpp +++ b/amd/comgr/test-unit/HotswapMCTest.cpp @@ -71,6 +71,106 @@ static TargetIdentifier makeGfx1250Ident() { return TI; } +static std::vector +decodeAsmSequence(const LLVMState &S, llvm::ArrayRef Lines) { + llvm::SmallVector Bytes; + for (llvm::StringRef Line : Lines) { + llvm::SmallVector Encoded = assembleSingleInst(Line, S); + EXPECT_FALSE(Encoded.empty()) << Line.str(); + Bytes.append(Encoded); + } + std::vector Decoded; + EXPECT_TRUE(decodeTextSection(Bytes.data(), Bytes.size(), S, Decoded)); + return Decoded; +} + +static bool scalarIncomingSgprIsUnsafe( + llvm::ArrayRef Decoded, const LLVMState &S, + uint64_t FunctionBegin, uint64_t FunctionEnd, uint64_t Continuation, + llvm::ArrayRef NumberedSgprs, unsigned Sgpr) { + auto FindInstruction = [&](uint64_t Offset) -> std::optional { + if (Offset < FunctionBegin || Offset >= FunctionEnd) + return std::nullopt; + auto It = llvm::lower_bound( + Decoded, Offset, [](const InternalDecodedInst &DI, uint64_t Target) { + return DI.Offset < Target; + }); + if (It == Decoded.end() || It->Offset != Offset) + return std::nullopt; + return It - Decoded.begin(); + }; + std::optional Start = FindInstruction(Continuation); + if (!Start) + return true; + + llvm::SmallVector Worklist(1, *Start); + llvm::DenseSet Visited; + while (!Worklist.empty()) { + size_t Index = Worklist.pop_back_val(); + if (!Visited.insert(Index).second) + continue; + const InternalDecodedInst &DI = Decoded[Index]; + if (!DI.DecodeSucceeded || !S.MIA || DI.Offset < FunctionBegin || + DI.Offset >= FunctionEnd) + return true; + + llvm::BitVector Uses(NumberedSgprs.size()); + llvm::BitVector Defs(NumberedSgprs.size()); + getNumberedSgprUsesAndDefs(DI, S, NumberedSgprs, Uses, Defs); + if (Uses.test(Sgpr)) + return true; + if (Defs.test(Sgpr) || DI.Inst.getOpcode() == S.SEndPgmOpcode || + DI.Inst.getOpcode() == S.SEndPgmSavedOpcode) + continue; + + auto AddSuccessor = [&](uint64_t Offset) { + std::optional Successor = FindInstruction(Offset); + if (!Successor) + return false; + Worklist.push_back(*Successor); + return true; + }; + if (S.MIA->isCall(DI.Inst) || S.MIA->isIndirectBranch(DI.Inst) || + S.MIA->isReturn(DI.Inst)) + return true; + if (S.MIA->isBranch(DI.Inst)) { + std::optional Target = evaluateDirectControlFlowTarget(DI, S); + if (!Target || !AddSuccessor(*Target)) + return true; + if (S.MIA->isUnconditionalBranch(DI.Inst)) + continue; + } else if (S.MIA->mayAffectControlFlow(DI.Inst, *S.MRI) && + !S.MIA->isBarrier(DI.Inst)) { + return true; + } + std::optional Fallthrough = + llvm::checkedAddUnsigned(DI.Offset, static_cast(DI.Size)); + if (!Fallthrough || !AddSuccessor(*Fallthrough)) + return true; + } + return false; +} + +static void +expectBatchSgprProofMatchesScalar(const LLVMState &S, + llvm::ArrayRef Lines) { + std::vector Decoded = decodeAsmSequence(S, Lines); + ASSERT_FALSE(Decoded.empty()); + std::optional> NumberedSgprs = + resolveNumberedSgprRegisters(*S.MRI, /*MaxSgprs=*/106); + ASSERT_TRUE(NumberedSgprs); + uint64_t FunctionEnd = Decoded.back().Offset + Decoded.back().Size; + std::optional Batch = unsafeIncomingNumberedSgprsInRange( + Decoded, S, /*FunctionBegin=*/0, FunctionEnd, /*Continuation=*/0, + *NumberedSgprs); + ASSERT_TRUE(Batch); + for (unsigned I = 0; I != NumberedSgprs->size(); ++I) + EXPECT_EQ(Batch->test(I), scalarIncomingSgprIsUnsafe( + Decoded, S, /*FunctionBegin=*/0, FunctionEnd, + /*Continuation=*/0, *NumberedSgprs, I)) + << "s" << I; +} + // Helper: decode the little-endian 32-bit dword at \p Bytes. static uint32_t readDword(const uint8_t *Bytes) { uint32_t V; @@ -452,6 +552,26 @@ TEST(EncodeSetPCLongBranch, ForwardLandsOnTarget) { EXPECT_EQ(From + MinInstSize + Delta, To); } +TEST(EncodeSetPCLongBranch, UsesVccWhenRequested) { + LLVMState S = initLLVM(makeGfx1250Ident()); + ASSERT_TRUE(S.Valid); + + std::optional> Out = encodeSetPCLongBranch( + S, /*FromOffset=*/0x1000, /*TargetOffset=*/0x81000, /*SgprBase=*/0, + /*UseVcc=*/true); + ASSERT_TRUE(Out); + + std::vector Decoded; + ASSERT_TRUE(decodeTextSection(Out->data(), Out->size(), S, Decoded)); + ASSERT_EQ(Decoded.size(), 3u); + for (const InternalDecodedInst &DI : Decoded) { + ASSERT_NE(DI.Inst.getNumOperands(), 0u); + ASSERT_TRUE(DI.Inst.getOperand(0).isReg()); + EXPECT_TRUE(S.MRI->regsOverlap( + llvm::MCRegister(DI.Inst.getOperand(0).getReg()), S.VCCRegister)); + } +} + TEST(EncodeSetPCLongBranch, InlineDisplacementUsesTwelveBytes) { LLVMState S = initLLVM(makeGfx1250Ident()); ASSERT_TRUE(S.Valid); @@ -490,6 +610,30 @@ TEST(FindNearestSetPcGateway, FitsActualSixteenByteEncoding) { EXPECT_EQ(Gateways[0].WritePos, 0x100u); } +TEST(FindNearestSetPcGateway, PrependsWave32VccSave) { + LLVMState S = initLLVM(makeGfx1250Ident()); + ASSERT_TRUE(S.Valid); + + std::vector Gateways = { + {/*Start=*/0x100, /*End=*/0x118, /*WritePos=*/0x100, + /*FunctionStart=*/0, /*FunctionEnd=*/0x1000}}; + llvm::Expected> GatewayOrErr = + findNearestSetPcGateway( + Gateways, S, /*FromOffset=*/0, /*TargetOffset=*/0x81000, + /*SgprBase=*/105, /*UseVcc=*/true, /*PreserveVcc=*/true); + ASSERT_TRUE((bool)GatewayOrErr) << llvm::toString(GatewayOrErr.takeError()); + std::optional &Gateway = *GatewayOrErr; + ASSERT_TRUE(Gateway); + + std::vector Decoded; + ASSERT_TRUE(decodeTextSection(Gateway->Bytes.data(), Gateway->Bytes.size(), S, + Decoded)); + ASSERT_EQ(Decoded.size(), 4u); + EXPECT_EQ(Decoded.front().Mnemonic, "s_mov_b32"); + EXPECT_EQ(Decoded[1].Mnemonic, "s_get_pc_i64"); + EXPECT_EQ(Decoded.back().Mnemonic, "s_set_pc_i64"); +} + TEST(FindNearestSetPcGateway, SkipsNearerUndersizedCandidate) { LLVMState S = initLLVM(makeGfx1250Ident()); ASSERT_TRUE(S.Valid); @@ -964,6 +1108,42 @@ TEST(CollectDirectBranchTargets, IgnoresSetPcWithoutTreatingItAsCall) { EXPECT_FALSE(Info->HasUnresolvedTargets); } +TEST(CollectDirectBranchTargets, + HandlesSparseSetPcAndFunctionRangesWithoutCartesianScan) { + LLVMState S = initLLVM(makeGfx1250Ident()); + ASSERT_TRUE(S.Valid); + llvm::SmallVector Bytes = + assembleSingleInst("s_set_pc_i64 s[8:9]", S); + ASSERT_FALSE(Bytes.empty()); + + std::vector Prototype; + ASSERT_TRUE(decodeTextSection(Bytes.data(), Bytes.size(), S, Prototype)); + ASSERT_EQ(Prototype.size(), 1u); + + constexpr size_t Count = 16384; + constexpr uint64_t TextSize = 1 << 20; + std::vector Decoded; + Decoded.reserve(Count); + llvm::SmallVector FunctionRanges; + FunctionRanges.reserve(Count); + for (size_t I = 0; I != Count; ++I) { + Decoded.push_back(Prototype[0]); + Decoded.back().Offset = I * MinInstSize; + // These validly ordered ranges do not cover any instruction. This pins + // the sparse return-to-range index: a full range scan for every set-PC + // would perform Count squared containment checks. + uint64_t Begin = TextSize + I * MinInstSize; + FunctionRanges.push_back({Begin, Begin + MinInstSize}); + } + + std::optional Info = + collectDirectBranchTargets(Decoded, S, /*TextAddr=*/0, TextSize, + /*DeclaredEntries=*/{}, FunctionRanges); + ASSERT_TRUE(Info); + EXPECT_TRUE(Info->Targets.empty()); + EXPECT_FALSE(Info->HasUnresolvedTargets); +} + TEST(CollectDirectBranchTargets, ResolvesProductionPcMaterializedCall) { LLVMState S = initLLVM(makeGfx1250Ident()); ASSERT_TRUE(S.Valid); @@ -1637,6 +1817,141 @@ TEST(AssembleDecode, SNopRoundTrip) { EXPECT_EQ(Decoded[0].Mnemonic, "s_nop"); } +TEST(RegisterLiveness, TiedAccumulatorDefCountsAsIncomingRead) { + LLVMState S = initLLVM(makeGfx1250Ident()); + ASSERT_TRUE(S.Valid); + + llvm::SmallVector Bytes = + assembleSingleInst("v_fmac_f32_e32 v5, v1, v2", S); + ASSERT_EQ(Bytes.size(), MinInstSize); + + std::vector Decoded; + ASSERT_TRUE(decodeTextSection(Bytes.data(), Bytes.size(), S, Decoded)); + ASSERT_EQ(Decoded.size(), 1u); + const InternalDecodedInst &DI = Decoded[0]; + const llvm::MCInstrDesc &Desc = S.MCII->get(DI.Inst.getOpcode()); + ASSERT_GE(Desc.getNumDefs(), 1u); + ASSERT_GE(DI.Inst.getNumOperands(), 1u); + ASSERT_TRUE(DI.Inst.getOperand(0).isReg()); + + bool HasTiedAccumulatorUse = false; + for (unsigned I = Desc.getNumDefs(); I != Desc.getNumOperands(); ++I) + HasTiedAccumulatorUse |= + Desc.getOperandConstraint(I, llvm::MCOI::TIED_TO) == 0; + ASSERT_TRUE(HasTiedAccumulatorUse); + + llvm::MCRegister Accumulator(DI.Inst.getOperand(0).getReg()); + EXPECT_TRUE(instructionReadsRegister(DI, S, Accumulator)); +} + +TEST(RegisterLiveness, BatchProofMatchesScalarAcrossControlFlow) { + LLVMState S = initLLVM(makeGfx1250Ident()); + ASSERT_TRUE(S.Valid); + + const llvm::StringRef BranchJoin[] = {"s_cbranch_scc0 1", "s_mov_b32 s30, 0", + "s_mov_b32 s0, s30", "s_endpgm"}; + expectBatchSgprProofMatchesScalar(S, BranchJoin); + + const llvm::StringRef Loop[] = {"s_mov_b32 s30, 0", "s_cbranch_scc0 -2", + "s_endpgm"}; + expectBatchSgprProofMatchesScalar(S, Loop); + + const llvm::StringRef DefBeforeOpaque[] = {"s_mov_b32 s30, 0", + "s_set_pc_i64 s[0:1]"}; + expectBatchSgprProofMatchesScalar(S, DefBeforeOpaque); + + const llvm::StringRef OpaqueBeforeDef[] = {"s_set_pc_i64 s[0:1]", + "s_mov_b32 s30, 0"}; + expectBatchSgprProofMatchesScalar(S, OpaqueBeforeDef); +} + +TEST(RegisterLiveness, NumberedSgprExtractionCoversAliasesAndTiedRmw) { + LLVMState S = initLLVM(makeGfx1250Ident()); + ASSERT_TRUE(S.Valid); + std::optional> NumberedSgprs = + resolveNumberedSgprRegisters(*S.MRI, /*MaxSgprs=*/106); + ASSERT_TRUE(NumberedSgprs); + + auto GetUseDef = [&](const InternalDecodedInst &DI) { + std::pair Result{ + llvm::BitVector(NumberedSgprs->size()), + llvm::BitVector(NumberedSgprs->size())}; + getNumberedSgprUsesAndDefs(DI, S, *NumberedSgprs, Result.first, + Result.second); + return Result; + }; + + std::vector Tuple = decodeAsmSequence( + S, llvm::ArrayRef({"s_mov_b64 s[0:1], s[30:31]"})); + ASSERT_EQ(Tuple.size(), 1u); + auto [TupleUses, TupleDefs] = GetUseDef(Tuple.front()); + EXPECT_TRUE(TupleUses.test(30)); + EXPECT_TRUE(TupleUses.test(31)); + EXPECT_TRUE(TupleDefs.test(0)); + EXPECT_TRUE(TupleDefs.test(1)); + + std::vector Rmw = decodeAsmSequence( + S, llvm::ArrayRef({"s_add_u32 s30, s30, 1"})); + ASSERT_EQ(Rmw.size(), 1u); + auto [RmwUses, RmwDefs] = GetUseDef(Rmw.front()); + EXPECT_TRUE(RmwUses.test(30)); + EXPECT_TRUE(RmwDefs.test(30)); + + llvm::MCRegister Low16; + for (unsigned I = 1; I != S.MRI->getNumRegs(); ++I) { + llvm::MCRegister Candidate(I); + if (llvm::StringRef(S.MRI->getName(Candidate)) == "SGPR30_LO16") { + Low16 = Candidate; + break; + } + } + ASSERT_TRUE(Low16.isValid()); + + std::vector Half = decodeAsmSequence( + S, llvm::ArrayRef({"s_mov_b32 s0, s1"})); + ASSERT_EQ(Half.size(), 1u); + ASSERT_GE(Half.front().Inst.getNumOperands(), 2u); + Half.front().Inst.getOperand(1).setReg(Low16); + auto [HalfUses, HalfDefs] = GetUseDef(Half.front()); + EXPECT_TRUE(HalfUses.test(30)); + EXPECT_FALSE(HalfUses.test(1)); + EXPECT_TRUE(HalfDefs.test(0)); + + Half.front().Inst.getOperand(0).setReg(Low16); + auto [HalfDefUses, HalfDefDefs] = GetUseDef(Half.front()); + EXPECT_TRUE(HalfDefUses.test(30)); + EXPECT_TRUE(HalfDefDefs.test(30)); +} + +TEST(RegisterLiveness, ReplacementTracksOnlyIncomingValues) { + LLVMState S = initLLVM(makeGfx1250Ident()); + ASSERT_TRUE(S.Valid); + std::optional> NumberedSgprs = + resolveNumberedSgprRegisters(*S.MRI, /*MaxSgprs=*/106); + ASSERT_TRUE(NumberedSgprs); + + llvm::SmallVector UseBeforeDef = + assembleInstructions("s_mov_b32 s0, s30\ns_mov_b32 s30, 0", S); + ASSERT_FALSE(UseBeforeDef.empty()); + llvm::BitVector UseBeforeDefUnsafe = + unsafeIncomingNumberedSgprsInReplacement(UseBeforeDef, S, *NumberedSgprs); + EXPECT_TRUE(UseBeforeDefUnsafe.test(30)); + + llvm::SmallVector DefBeforeUse = + assembleInstructions("s_mov_b32 s30, 0\ns_mov_b32 s0, s30", S); + ASSERT_FALSE(DefBeforeUse.empty()); + llvm::BitVector DefBeforeUseUnsafe = + unsafeIncomingNumberedSgprsInReplacement(DefBeforeUse, S, *NumberedSgprs); + EXPECT_FALSE(DefBeforeUseUnsafe.test(30)); + + llvm::SmallVector Opaque = + assembleSingleInst("s_set_pc_i64 s[0:1]", S); + ASSERT_FALSE(Opaque.empty()); + llvm::BitVector OpaqueUnsafe = + unsafeIncomingNumberedSgprsInReplacement(Opaque, S, *NumberedSgprs); + EXPECT_TRUE(OpaqueUnsafe.test(30)); +} + TEST(AssembleDecode, SingleInstructionRejectsSequence) { LLVMState S = initLLVM(makeGfx1250Ident()); ASSERT_TRUE(S.Valid);