From e29247ce80a13fd74fa8081119662f131aee7cbb Mon Sep 17 00:00:00 2001 From: Harsh Menon Date: Fri, 24 Jul 2026 20:40:50 -0700 Subject: [PATCH] comgr: prove zero tensor masks through local control flow --- amd/comgr/src/comgr-hotswap-b0a0.cpp | 33 +- amd/comgr/src/comgr-hotswap-internal.h | 5 + .../src/comgr-hotswap-patch-trampoline.cpp | 617 +++++++++++++++++- ...swap-trampoline-tensor-mask-already-zero.s | 110 ++++ ...ap-trampoline-tensor-mask-bounded-return.s | 99 +++ ...tswap-trampoline-tensor-mask-local-setpc.s | 156 +++++ ...swap-trampoline-tensor-mask-signed-setpc.s | 92 +++ 7 files changed, 1097 insertions(+), 15 deletions(-) create mode 100644 amd/comgr/test-lit/hotswap-trampoline-tensor-mask-already-zero.s create mode 100644 amd/comgr/test-lit/hotswap-trampoline-tensor-mask-bounded-return.s create mode 100644 amd/comgr/test-lit/hotswap-trampoline-tensor-mask-local-setpc.s create mode 100644 amd/comgr/test-lit/hotswap-trampoline-tensor-mask-signed-setpc.s diff --git a/amd/comgr/src/comgr-hotswap-b0a0.cpp b/amd/comgr/src/comgr-hotswap-b0a0.cpp index b71a44beaff13..732d87651afcd 100755 --- a/amd/comgr/src/comgr-hotswap-b0a0.cpp +++ b/amd/comgr/src/comgr-hotswap-b0a0.cpp @@ -4036,8 +4036,37 @@ std::optional collectDirectBranchTargets( Info.Targets.insert(*Continuation); } } - for (const BoundedSetPcReturn &Return : BoundedReturns) - Info.BoundedIndirectTransfers.insert(Decoded[Return.InstIndex].Offset); + DenseSet DecodedOffsets; + for (const InternalDecodedInst &DI : Decoded) + if (DI.DecodeSucceeded) + DecodedOffsets.insert(DI.Offset); + for (const BoundedSetPcReturn &Return : BoundedReturns) { + uint64_t ReturnOffset = Decoded[Return.InstIndex].Offset; + SmallVector LocalTargets; + for (uint64_t Target : Return.Targets) { + if (Target >= TextSize) + continue; + if (!DecodedOffsets.contains(Target)) { + log() << "hotswap: audited bounded return at 0x" + << utohexstr(ReturnOffset) << " has non-boundary local target 0x" + << utohexstr(Target) << "\n"; + return std::nullopt; + } + LocalTargets.push_back(Target); + } + llvm::sort(LocalTargets); + LocalTargets.erase(std::unique(LocalTargets.begin(), LocalTargets.end()), + LocalTargets.end()); + auto Inserted = + Info.BoundedIndirectTargets.try_emplace(ReturnOffset, LocalTargets); + if (!Inserted.second && Inserted.first->second != LocalTargets) { + log() << "hotswap: conflicting audited target sets for bounded return at " + "0x" + << utohexstr(ReturnOffset) << "\n"; + return std::nullopt; + } + Info.BoundedIndirectTransfers.insert(ReturnOffset); + } if (!IndirectControlFlowClosed && HasUnboundedIndirectEntries) Info.HasUnboundedIndirectEntries = true; return Info; diff --git a/amd/comgr/src/comgr-hotswap-internal.h b/amd/comgr/src/comgr-hotswap-internal.h index 7bf556ad8f38e..29e46bc8fae26 100644 --- a/amd/comgr/src/comgr-hotswap-internal.h +++ b/amd/comgr/src/comgr-hotswap-internal.h @@ -1286,6 +1286,11 @@ struct DirectControlFlowInfo { // These do not make every instruction in their containing function a // potential indirect destination. llvm::DenseSet BoundedIndirectTransfers; + // Exact .text-relative decoded-boundary targets for each proven finite + // register transfer, keyed by that transfer's decoded offset. An empty + // target vector means the transfer is proven to leave local .text. + llvm::DenseMap> + BoundedIndirectTargets; // A reachable indirect transfer can enter bytes that are not represented // by an original instruction or symbol, including synthetic source tails // created while planning gateways. Keep this distinct from unresolved call diff --git a/amd/comgr/src/comgr-hotswap-patch-trampoline.cpp b/amd/comgr/src/comgr-hotswap-patch-trampoline.cpp index ef3c2cf05802f..9edaf24ca4c1c 100755 --- a/amd/comgr/src/comgr-hotswap-patch-trampoline.cpp +++ b/amd/comgr/src/comgr-hotswap-patch-trampoline.cpp @@ -34,6 +34,7 @@ #include "llvm/ADT/StringExtras.h" #include "llvm/ADT/StringRef.h" #include "llvm/ADT/StringSwitch.h" +#include "llvm/MC/MCExpr.h" #include "llvm/MC/MCRegisterInfo.h" #include "llvm/Support/raw_ostream.h" @@ -1028,6 +1029,290 @@ struct TensorFunctionCfg { } }; +enum class TensorLocalSetPcShape { Linear, SignedTwoArm }; + +struct TensorLocalSetPcResolution { + uint64_t Target = 0; + size_t SequenceBeginIndex = 0; + size_t SequenceEndIndex = 0; + TensorLocalSetPcShape Shape = TensorLocalSetPcShape::Linear; +}; + +bool isExactTensorRegisterOperand(const MCInst &Inst, unsigned OperandIndex, + MCRegister Reg) { + return OperandIndex < Inst.getNumOperands() && + Inst.getOperand(OperandIndex).isReg() && + Inst.getOperand(OperandIndex).getReg() == Reg; +} + +std::optional evaluateTensorUint32Operand(const MCOperand &Operand) { + if (Operand.isImm()) + return static_cast(Operand.getImm()); + if (!Operand.isExpr()) + return std::nullopt; + int64_t Value = 0; + if (!Operand.getExpr()->evaluateAsAbsolute(Value)) + return std::nullopt; + return static_cast(Value); +} + +// Resolve the compiler-emitted reusable-PC tail jump used by large tensor +// kernels: +// +// s_get_pc_i64 Pair +// s_add_co_i32 Delta, Imm0, Imm1 +// s_add_co_u32 Pair.lo, Pair.lo, Delta +// s_add_co_ci_u32 Pair.hi, Pair.hi, 0 +// s_set_pc_i64 Pair +// +// Keep this recognizer deliberately narrow. All five instructions must be +// adjacent and use the exact register relationships above; any variation +// remains unresolved and makes the tensor definition proof fail closed. +std::optional +resolveTensorLinearSetPcTarget(const PatchContext &Ctx, size_t SetPcIndex, + size_t BeginIndex, + const ElfView::FunctionTextRange &Range) { + if (SetPcIndex < BeginIndex || SetPcIndex - BeginIndex < 4) + return std::nullopt; + + const InternalDecodedInst &GetPc = Ctx.Decoded[SetPcIndex - 4]; + const InternalDecodedInst &MakeDelta = Ctx.Decoded[SetPcIndex - 3]; + const InternalDecodedInst &AddLow = Ctx.Decoded[SetPcIndex - 2]; + const InternalDecodedInst &AddHigh = Ctx.Decoded[SetPcIndex - 1]; + const InternalDecodedInst &SetPc = Ctx.Decoded[SetPcIndex]; + if (!GetPc.DecodeSucceeded || !MakeDelta.DecodeSucceeded || + !AddLow.DecodeSucceeded || !AddHigh.DecodeSucceeded || + !SetPc.DecodeSucceeded || + GetPc.Inst.getOpcode() != Ctx.LS.SGetPcI64Opcode || + MakeDelta.Mnemonic != "s_add_co_i32" || + AddLow.Mnemonic != "s_add_co_u32" || + AddHigh.Mnemonic != "s_add_co_ci_u32" || + SetPc.Inst.getOpcode() != Ctx.LS.SSetPcI64Opcode || + GetPc.Inst.getNumOperands() != 1 || + MakeDelta.Inst.getNumOperands() != 3 || + AddLow.Inst.getNumOperands() != 3 || AddHigh.Inst.getNumOperands() != 3 || + SetPc.Inst.getNumOperands() != 1) + return std::nullopt; + + auto IsImmediatelyBefore = [](const InternalDecodedInst &Before, + const InternalDecodedInst &After) { + return Before.Offset <= + std::numeric_limits::max() - Before.Size && + Before.Offset + Before.Size == After.Offset; + }; + if (!IsImmediatelyBefore(GetPc, MakeDelta) || + !IsImmediatelyBefore(MakeDelta, AddLow) || + !IsImmediatelyBefore(AddLow, AddHigh) || + !IsImmediatelyBefore(AddHigh, SetPc)) + return std::nullopt; + if (GetPc.Offset < Range.Begin || SetPc.Offset >= Range.End || + SetPc.Size > Range.End - SetPc.Offset) + return std::nullopt; + + std::optional FirstAddend = + evaluateTensorUint32Operand(MakeDelta.Inst.getOperand(1)); + std::optional SecondAddend = + evaluateTensorUint32Operand(MakeDelta.Inst.getOperand(2)); + + const MCOperand &GetPcPair = GetPc.Inst.getOperand(0); + const MCOperand &SetPcPair = SetPc.Inst.getOperand(0); + const MCOperand &DeltaDst = MakeDelta.Inst.getOperand(0); + if (!GetPcPair.isReg() || !GetPcPair.getReg() || !SetPcPair.isReg() || + SetPcPair.getReg() != GetPcPair.getReg() || !DeltaDst.isReg() || + !DeltaDst.getReg() || !FirstAddend || !SecondAddend) + return std::nullopt; + + MCRegister Pair(GetPcPair.getReg()); + MCRegister DeltaReg(DeltaDst.getReg()); + if (Ctx.LS.MRI->regsOverlap(Pair, DeltaReg)) + return std::nullopt; + + auto SameRegOperand = [](const MCInst &Inst, unsigned OperandIndex, + MCRegister Reg) { + return Inst.getOperand(OperandIndex).isReg() && + Inst.getOperand(OperandIndex).getReg() == Reg; + }; + + if (!AddLow.Inst.getOperand(0).isReg() || + !AddLow.Inst.getOperand(0).getReg() || + !AddHigh.Inst.getOperand(0).isReg() || + !AddHigh.Inst.getOperand(0).getReg()) + return std::nullopt; + MCRegister Low(AddLow.Inst.getOperand(0).getReg()); + MCRegister High(AddHigh.Inst.getOperand(0).getReg()); + if (!SameRegOperand(AddLow.Inst, 1, Low) || + !SameRegOperand(AddLow.Inst, 2, DeltaReg) || + !SameRegOperand(AddHigh.Inst, 1, High) || + !AddHigh.Inst.getOperand(2).isImm() || + AddHigh.Inst.getOperand(2).getImm() != 0) + return std::nullopt; + + std::optional LowIndex = getSgprIndex(Low, *Ctx.LS.MRI); + std::optional HighIndex = getSgprIndex(High, *Ctx.LS.MRI); + if (!LowIndex || !HighIndex || *HighIndex != *LowIndex + 1 || + !Ctx.LS.MRI->regsOverlap(Low, Pair) || + !Ctx.LS.MRI->regsOverlap(High, Pair)) + return std::nullopt; + + uint32_t Delta = *FirstAddend + *SecondAddend; + std::optional PcValue = checkedAddUint64( + GetPc.Offset, GetPc.Size, "tensor CFG reusable-PC instruction"); + if (!PcValue) + return std::nullopt; + std::optional Target = + checkedAddUint64(*PcValue, Delta, "tensor CFG reusable-PC target"); + if (!Target) + return std::nullopt; + return TensorLocalSetPcResolution{*Target, SetPcIndex - 4, SetPcIndex, + TensorLocalSetPcShape::Linear}; +} + +// Resolve Tensile's signed-direction reusable-PC transfer. Both arms compute +// the same modulo-2^64 target: +// +// s_get_pc_i64 Pair +// s_add_co_i32 Delta, Imm0, Imm1 +// s_cmp_ge_i32 Delta, 0 +// s_cbranch_scc1 Positive +// s_abs_i32 Delta, Delta +// s_sub_co_u32 Pair.lo, Pair.lo, Delta +// s_sub_co_ci_u32 Pair.hi, Pair.hi, 0 +// s_set_pc_i64 Pair +// Positive: +// s_add_co_u32 Pair.lo, Pair.lo, Delta +// s_add_co_ci_u32 Pair.hi, Pair.hi, 0 +// s_set_pc_i64 Pair +// +// The caller separately verifies that no control-flow edge enters either arm +// or the materialization interior. +std::optional +resolveTensorSignedSetPcTarget(const PatchContext &Ctx, size_t SetPcIndex, + size_t BeginIndex, size_t EndIndex, + const ElfView::FunctionTextRange &Range) { + for (size_t SetPcPosition : {size_t{7}, size_t{10}}) { + if (SetPcIndex < BeginIndex || SetPcIndex - BeginIndex < SetPcPosition) + continue; + size_t FirstIndex = SetPcIndex - SetPcPosition; + if (FirstIndex > EndIndex || EndIndex - FirstIndex <= 10) + continue; + + const InternalDecodedInst &GetPc = Ctx.Decoded[FirstIndex]; + const InternalDecodedInst &MakeDelta = Ctx.Decoded[FirstIndex + 1]; + const InternalDecodedInst &Compare = Ctx.Decoded[FirstIndex + 2]; + const InternalDecodedInst &Branch = Ctx.Decoded[FirstIndex + 3]; + const InternalDecodedInst &Abs = Ctx.Decoded[FirstIndex + 4]; + const InternalDecodedInst &SubLow = Ctx.Decoded[FirstIndex + 5]; + const InternalDecodedInst &SubHigh = Ctx.Decoded[FirstIndex + 6]; + const InternalDecodedInst &NegativeSetPc = Ctx.Decoded[FirstIndex + 7]; + const InternalDecodedInst &AddLow = Ctx.Decoded[FirstIndex + 8]; + const InternalDecodedInst &AddHigh = Ctx.Decoded[FirstIndex + 9]; + const InternalDecodedInst &PositiveSetPc = Ctx.Decoded[FirstIndex + 10]; + + bool AllDecodedAndAdjacent = true; + for (size_t I = FirstIndex; I <= FirstIndex + 10; ++I) { + if (!Ctx.Decoded[I].DecodeSucceeded) { + AllDecodedAndAdjacent = false; + break; + } + if (I == FirstIndex + 10) + continue; + const InternalDecodedInst &Current = Ctx.Decoded[I]; + if (Current.Offset > + std::numeric_limits::max() - Current.Size || + Current.Offset + Current.Size != Ctx.Decoded[I + 1].Offset) { + AllDecodedAndAdjacent = false; + break; + } + } + if (!AllDecodedAndAdjacent || GetPc.Offset < Range.Begin || + PositiveSetPc.Offset >= Range.End || + PositiveSetPc.Size > Range.End - PositiveSetPc.Offset || + GetPc.Inst.getOpcode() != Ctx.LS.SGetPcI64Opcode || + GetPc.Inst.getNumOperands() != 1 || !GetPc.Inst.getOperand(0).isReg() || + !GetPc.Inst.getOperand(0).getReg() || + MakeDelta.Mnemonic != "s_add_co_i32" || + MakeDelta.Inst.getNumOperands() != 3 || + !MakeDelta.Inst.getOperand(0).isReg() || + !MakeDelta.Inst.getOperand(0).getReg() || + !MakeDelta.Inst.getOperand(2).isImm()) + continue; + + MCRegister Pair(GetPc.Inst.getOperand(0).getReg()); + MCRegister DeltaReg(MakeDelta.Inst.getOperand(0).getReg()); + if (Ctx.LS.MRI->regsOverlap(DeltaReg, Pair) || + Compare.Mnemonic != "s_cmp_ge_i32" || + Compare.Inst.getNumOperands() != 2 || + !isExactTensorRegisterOperand(Compare.Inst, 0, DeltaReg) || + !Compare.Inst.getOperand(1).isImm() || + Compare.Inst.getOperand(1).getImm() != 0 || + Branch.Mnemonic != "s_cbranch_scc1" || Abs.Mnemonic != "s_abs_i32" || + Abs.Inst.getNumOperands() != 2 || + !isExactTensorRegisterOperand(Abs.Inst, 0, DeltaReg) || + !isExactTensorRegisterOperand(Abs.Inst, 1, DeltaReg) || + NegativeSetPc.Inst.getOpcode() != Ctx.LS.SSetPcI64Opcode || + NegativeSetPc.Inst.getNumOperands() != 1 || + !isExactTensorRegisterOperand(NegativeSetPc.Inst, 0, Pair) || + PositiveSetPc.Inst.getOpcode() != Ctx.LS.SSetPcI64Opcode || + PositiveSetPc.Inst.getNumOperands() != 1 || + !isExactTensorRegisterOperand(PositiveSetPc.Inst, 0, Pair)) + continue; + + uint64_t PositiveTarget = 0; + if (!Ctx.LS.MIA->evaluateBranch(Branch.Inst, Branch.Offset, Branch.Size, + PositiveTarget) || + PositiveTarget != AddLow.Offset) + continue; + + auto MatchesPairArithmetic = [&](const InternalDecodedInst &Low, + const InternalDecodedInst &High, + StringRef LowMnemonic, + StringRef HighMnemonic) { + if (Low.Mnemonic != LowMnemonic || High.Mnemonic != HighMnemonic || + Low.Inst.getNumOperands() != 3 || !Low.Inst.getOperand(0).isReg() || + !Low.Inst.getOperand(1).isReg() || !Low.Inst.getOperand(0).getReg() || + Low.Inst.getOperand(0).getReg() != Low.Inst.getOperand(1).getReg() || + !isExactTensorRegisterOperand(Low.Inst, 2, DeltaReg) || + High.Inst.getNumOperands() != 3 || !High.Inst.getOperand(0).isReg() || + !High.Inst.getOperand(1).isReg() || + !High.Inst.getOperand(0).getReg() || + High.Inst.getOperand(0).getReg() != + High.Inst.getOperand(1).getReg() || + !High.Inst.getOperand(2).isImm() || + High.Inst.getOperand(2).getImm() != 0) + return false; + MCRegister LowReg(Low.Inst.getOperand(0).getReg()); + MCRegister HighReg(High.Inst.getOperand(0).getReg()); + std::optional LowIndex = getSgprIndex(LowReg, *Ctx.LS.MRI); + std::optional HighIndex = getSgprIndex(HighReg, *Ctx.LS.MRI); + return LowIndex && HighIndex && *HighIndex == *LowIndex + 1 && + Ctx.LS.MRI->regsOverlap(LowReg, Pair) && + Ctx.LS.MRI->regsOverlap(HighReg, Pair); + }; + if (!MatchesPairArithmetic(SubLow, SubHigh, "s_sub_co_u32", + "s_sub_co_ci_u32") || + !MatchesPairArithmetic(AddLow, AddHigh, "s_add_co_u32", + "s_add_co_ci_u32")) + continue; + + std::optional FirstAddend = + evaluateTensorUint32Operand(MakeDelta.Inst.getOperand(1)); + if (!FirstAddend) + continue; + uint32_t DeltaBits = + *FirstAddend + + static_cast(MakeDelta.Inst.getOperand(2).getImm()); + int64_t SignedDelta = static_cast(DeltaBits); + std::optional PcValue = checkedAddUint64( + GetPc.Offset, GetPc.Size, "tensor CFG signed reusable-PC value"); + if (!PcValue) + continue; + uint64_t Target = *PcValue + static_cast(SignedDelta); + return TensorLocalSetPcResolution{Target, FirstIndex, FirstIndex + 10, + TensorLocalSetPcShape::SignedTwoArm}; + } + return std::nullopt; +} + std::optional buildTensorFunctionCfg(const PatchContext &Ctx, const ElfView::FunctionTextRange &Range) { @@ -1042,8 +1327,21 @@ buildTensorFunctionCfg(const PatchContext &Ctx, while (EndIndex < Ctx.Decoded.size() && Ctx.Decoded[EndIndex].Offset < Range.End) ++EndIndex; - if (BeginIndex == EndIndex || Ctx.Decoded[BeginIndex].Offset != Range.Begin) + if (BeginIndex == EndIndex || Ctx.Decoded[BeginIndex].Offset != Range.Begin) { + log() << "hotswap: tensor CFG rejected range [0x" << utohexstr(Range.Begin) + << ", 0x" << utohexstr(Range.End) + << "): range does not begin at a decoded instruction\n"; return std::nullopt; + } + for (uint64_t Entry : Ctx.DeclaredEntries) { + if (Entry <= Range.Begin || Entry >= Range.End) + continue; + log() << "hotswap: tensor CFG rejected range [0x" << utohexstr(Range.Begin) + << ", 0x" << utohexstr(Range.End) + << "): declared entry at interior offset 0x" << utohexstr(Entry) + << "\n"; + return std::nullopt; + } TensorFunctionCfg Graph; Graph.BeginIndex = BeginIndex; @@ -1055,10 +1353,50 @@ buildTensorFunctionCfg(const PatchContext &Ctx, for (size_t I = BeginIndex; I < EndIndex; ++I) IndexAtOffset[Ctx.Decoded[I].Offset] = I - BeginIndex; + struct PendingLocalSetPc { + size_t FirstLocalIndex; + size_t LastLocalIndex; + size_t SetPcLocalIndex; + std::optional TargetLocalIndex; + TensorLocalSetPcShape Shape; + bool Audited; + }; + SmallVector PendingLocalSetPcs; + + auto AddAuditedBoundedEdges = + [&](size_t LocalIndex, + const InternalDecodedInst &DI) -> std::optional { + auto Bounded = Ctx.DirectControlFlow.BoundedIndirectTargets.find(DI.Offset); + if (Bounded == Ctx.DirectControlFlow.BoundedIndirectTargets.end()) + return std::nullopt; + for (uint64_t Target : Bounded->second) { + if (Target < Range.Begin || Target >= Range.End) { + log() << "hotswap: tensor CFG rejected bounded transfer at 0x" + << utohexstr(DI.Offset) << " with out-of-range target 0x" + << utohexstr(Target) << "\n"; + return false; + } + DenseMap::const_iterator TargetIt = + IndexAtOffset.find(Target); + if (TargetIt == IndexAtOffset.end()) { + log() << "hotswap: tensor CFG rejected bounded transfer at 0x" + << utohexstr(DI.Offset) << " with non-boundary target 0x" + << utohexstr(Target) << "\n"; + return false; + } + Graph.addEdge(LocalIndex, TargetIt->second); + } + return true; + }; + for (size_t I = BeginIndex; I < EndIndex; ++I) { const InternalDecodedInst &DI = Ctx.Decoded[I]; - if (!DI.DecodeSucceeded) + if (!DI.DecodeSucceeded) { + log() << "hotswap: tensor CFG rejected range [0x" + << utohexstr(Range.Begin) << ", 0x" << utohexstr(Range.End) + << "): undecoded instruction at 0x" << utohexstr(DI.Offset) << "\n"; return std::nullopt; + } size_t LocalIndex = I - BeginIndex; bool HasFallthrough = I + 1 < EndIndex; @@ -1066,42 +1404,253 @@ buildTensorFunctionCfg(const PatchContext &Ctx, DI.Inst.getOpcode() == Ctx.LS.SEndPgmSavedOpcode) continue; - if (Ctx.LS.MIA->isCall(DI.Inst) || Ctx.LS.MIA->isIndirectBranch(DI.Inst) || - Ctx.LS.MIA->isReturn(DI.Inst)) + // Calls return to the next instruction, so retain their local fallthrough + // edge. The reaching-definition checks below reject a call only while a + // changed descriptor or SCC value is live; a call elsewhere in a broad + // symbol-less function range must not make an otherwise local proof fail. + if (Ctx.LS.MIA->isCall(DI.Inst)) { + if (!HasFallthrough) { + log() << "hotswap: tensor CFG rejected call without fallthrough at 0x" + << utohexstr(DI.Offset) << "\n"; + return std::nullopt; + } + Graph.addEdge(LocalIndex, LocalIndex + 1); + continue; + } + + if (Ctx.LS.MIA->isIndirectBranch(DI.Inst) || + Ctx.LS.MIA->isReturn(DI.Inst)) { + std::optional Added = AddAuditedBoundedEdges(LocalIndex, DI); + if (Added) { + if (!*Added) + return std::nullopt; + continue; + } + log() << "hotswap: tensor CFG rejected unresolved control flow at 0x" + << utohexstr(DI.Offset) << " (" << DI.Mnemonic << ")\n"; return std::nullopt; + } if (Ctx.LS.MIA->isBranch(DI.Inst)) { uint64_t Target = 0; - if (!Ctx.LS.MIA->evaluateBranch(DI.Inst, DI.Offset, DI.Size, Target)) - return std::nullopt; + bool ResolvedSetPc = false; + if (!Ctx.LS.MIA->evaluateBranch(DI.Inst, DI.Offset, DI.Size, Target)) { + std::optional SetPc = + resolveTensorLinearSetPcTarget(Ctx, I, BeginIndex, Range); + if (!SetPc) + SetPc = resolveTensorSignedSetPcTarget(Ctx, I, BeginIndex, EndIndex, + Range); + bool AuditedSetPc = false; + if (SetPc) { + auto Audited = + Ctx.DirectControlFlow.BoundedIndirectTargets.find(DI.Offset); + AuditedSetPc = + Audited != Ctx.DirectControlFlow.BoundedIndirectTargets.end() && + Audited->second.size() == 1 && + Audited->second.front() == SetPc->Target; + } + if (SetPc) { + Target = SetPc->Target; + ResolvedSetPc = true; + PendingLocalSetPcs.push_back({SetPc->SequenceBeginIndex - BeginIndex, + SetPc->SequenceEndIndex - BeginIndex, + LocalIndex, std::nullopt, SetPc->Shape, + AuditedSetPc}); + } else { + std::optional Added = AddAuditedBoundedEdges(LocalIndex, DI); + if (Added) { + if (!*Added) + return std::nullopt; + continue; + } + log() << "hotswap: tensor CFG rejected unresolved branch at 0x" + << utohexstr(DI.Offset) << " (" << DI.Mnemonic << ")\n"; + return std::nullopt; + } + } + + std::optional TargetLocalIndex; if (Target != Range.End) { DenseMap::const_iterator TargetIt = IndexAtOffset.find(Target); - if (TargetIt == IndexAtOffset.end()) + if (TargetIt == IndexAtOffset.end()) { + log() << "hotswap: tensor CFG rejected non-local branch target 0x" + << utohexstr(Target) << " from 0x" << utohexstr(DI.Offset) + << "\n"; return std::nullopt; - Graph.addEdge(LocalIndex, TargetIt->second); + } + TargetLocalIndex = TargetIt->second; + } + + if (ResolvedSetPc) { + PendingLocalSetPcs.back().TargetLocalIndex = TargetLocalIndex; + continue; } + if (TargetLocalIndex) + Graph.addEdge(LocalIndex, *TargetLocalIndex); if (Ctx.LS.MIA->isConditionalBranch(DI.Inst)) { - if (!HasFallthrough) + if (!HasFallthrough) { + log() << "hotswap: tensor CFG rejected conditional branch without " + "fallthrough at 0x" + << utohexstr(DI.Offset) << "\n"; return std::nullopt; + } Graph.addEdge(LocalIndex, LocalIndex + 1); } else if (!Ctx.LS.MIA->isUnconditionalBranch(DI.Inst)) { + log() << "hotswap: tensor CFG rejected unclassified branch at 0x" + << utohexstr(DI.Offset) << "\n"; return std::nullopt; } continue; } const MCInstrDesc &Desc = Ctx.LS.MCII->get(DI.Inst.getOpcode()); - if (Desc.mayAffectControlFlow(DI.Inst, *Ctx.LS.MRI)) + if (Desc.mayAffectControlFlow(DI.Inst, *Ctx.LS.MRI)) { + log() << "hotswap: tensor CFG rejected control-flow instruction at 0x" + << utohexstr(DI.Offset) << " (" << DI.Mnemonic << ")\n"; return std::nullopt; + } if (HasFallthrough) Graph.addEdge(LocalIndex, LocalIndex + 1); } + // Add every resolved transfer before validating sequence interiors so a + // second reusable-PC jump targeting another sequence's interior is visible + // as an alternate predecessor too. + for (const PendingLocalSetPc &Pending : PendingLocalSetPcs) + if (Pending.TargetLocalIndex) + Graph.addEdge(Pending.SetPcLocalIndex, *Pending.TargetLocalIndex); + + // Entering after s_get_pc_i64 can reuse stale pair/delta state, invalidating + // the computed target. Require every sequence-interior instruction to have + // only its exact arm predecessor as an entry. The signed form's positive + // add arm is entered by its conditional branch; every other instruction is + // entered by its immediate predecessor. This rejects direct targets, call + // continuations, function entry, and targets from another recognized + // reusable-PC transfer. + for (const PendingLocalSetPc &Pending : PendingLocalSetPcs) { + for (size_t LocalIndex = Pending.FirstLocalIndex + 1; + LocalIndex <= Pending.LastLocalIndex; ++LocalIndex) { + size_t ExpectedPredecessor = LocalIndex - 1; + if (Pending.Shape == TensorLocalSetPcShape::SignedTwoArm && + LocalIndex == Pending.FirstLocalIndex + 8) + ExpectedPredecessor = Pending.FirstLocalIndex + 3; + ArrayRef Predecessors = Graph.Predecessors[LocalIndex]; + if (Predecessors.size() == 1 && + Predecessors.front() == ExpectedPredecessor) + continue; + const InternalDecodedInst &Interior = + Ctx.Decoded[BeginIndex + LocalIndex]; + log() << "hotswap: tensor CFG rejected alternate entry into reusable-PC " + "sequence at 0x" + << utohexstr(Interior.Offset) << "\n"; + return std::nullopt; + } + if (!Pending.Audited) { + const InternalDecodedInst &SetPc = + Ctx.Decoded[BeginIndex + Pending.SetPcLocalIndex]; + log() << "hotswap: tensor CFG rejected unaudited reusable-PC transfer at " + "0x" + << utohexstr(SetPc.Offset) << "\n"; + return std::nullopt; + } + } + return Graph; } +bool writesBaseWithKnownZeroLow16(const InternalDecodedInst &DI, + MCRegister BaseMCReg, const LLVMState &LS) { + const MCInst &Inst = DI.Inst; + if (Inst.getNumOperands() < 2 || !Inst.getOperand(0).isReg() || + !Inst.getOperand(0).getReg() || + !LS.MRI->regsOverlap(MCRegister(Inst.getOperand(0).getReg()), BaseMCReg)) + return false; + + if (DI.Mnemonic == "s_mov_b32") + return Inst.getOperand(1).isImm() && + (static_cast(Inst.getOperand(1).getImm()) & 0xffffu) == 0; + + // An immediate AND with zeros in bits [15:0] forces the result's low half + // to zero regardless of the other input. + return Inst.getOpcode() == LS.SAndB32Opcode && Inst.getNumOperands() >= 3 && + Inst.getOperand(2).isImm() && + (static_cast(Inst.getOperand(2).getImm()) & 0xffffu) == 0; +} + +// Prove that the descriptor base already has a zero workgroup_mask at the +// tensor. Walk every reachable predecessor path backward. A path is complete +// only at an immediate zero definition (or zero-forcing AND); otherwise it may +// cross exact self-writes that cannot introduce a low bit. Calls, unresolved +// control flow, non-preserving definitions, and entry without a proven +// definition all reject. Reads are harmless because this proof changes no +// value. +bool isTensorMaskAlreadyZero(const PatchContext &Ctx, size_t TensorIdx, + MCRegister BaseMCReg) { + const InternalDecodedInst &Tensor = Ctx.Decoded[TensorIdx]; + std::optional Range = + Ctx.Elf.findFunctionTextRangeAtOffset(Tensor.Offset); + if (!Range) + return false; + + std::optional Graph = buildTensorFunctionCfg(Ctx, *Range); + if (!Graph || TensorIdx < Graph->BeginIndex || TensorIdx >= Graph->EndIndex) + return false; + + size_t TensorLocal = TensorIdx - Graph->BeginIndex; + SmallVector Reachable(Graph->Successors.size(), 0); + SmallVector Worklist; + Worklist.push_back(0); + while (!Worklist.empty()) { + size_t LocalIndex = Worklist.pop_back_val(); + if (Reachable[LocalIndex] != 0) + continue; + Reachable[LocalIndex] = 1; + for (size_t Successor : Graph->Successors[LocalIndex]) + Worklist.push_back(Successor); + } + if (Reachable[TensorLocal] == 0) + return false; + + SmallVector Visited(Graph->Successors.size(), 0); + for (size_t Predecessor : Graph->Predecessors[TensorLocal]) + if (Reachable[Predecessor] != 0) + Worklist.push_back(Predecessor); + if (Worklist.empty()) + return false; + + while (!Worklist.empty()) { + size_t LocalIndex = Worklist.pop_back_val(); + if (Visited[LocalIndex] != 0) + continue; + Visited[LocalIndex] = 1; + + const InternalDecodedInst &DI = Ctx.Decoded[Graph->BeginIndex + LocalIndex]; + if (Ctx.LS.MIA->isCall(DI.Inst)) + return false; + + if (instructionDefinesBase(DI, BaseMCReg, Ctx.LS)) { + if (writesBaseWithKnownZeroLow16(DI, BaseMCReg, Ctx.LS)) + continue; + if (!writesBasePreservingZeroLow16(DI, BaseMCReg, Ctx.LS)) + return false; + } + + bool HasReachablePredecessor = false; + for (size_t Predecessor : Graph->Predecessors[LocalIndex]) { + if (Reachable[Predecessor] == 0) + continue; + HasReachablePredecessor = true; + Worklist.push_back(Predecessor); + } + if (!HasReachablePredecessor) + return false; + } + + return true; +} + bool isMaskDefinitionSafe(const PatchContext &Ctx, const TensorFunctionCfg &Graph, size_t MaskIndex, MCRegister BaseMCReg) { @@ -1130,14 +1679,27 @@ bool isMaskDefinitionSafe(const PatchContext &Ctx, const InternalDecodedInst &DI = Ctx.Decoded[Graph.BeginIndex + LocalIndex]; const MCInstrDesc &Desc = Ctx.LS.MCII->get(DI.Inst.getOpcode()); + if (Ctx.LS.MIA->isCall(DI.Inst)) { + log() << "hotswap: tensor mask definition 0x" + << utohexstr(Ctx.Decoded[MaskIndex].Offset) + << " rejected: call while descriptor state is live at 0x" + << utohexstr(DI.Offset) << "\n"; + return false; + } + bool PreservesBaseValue = (State & BaseValueLive) != 0 && writesBasePreservingZeroLow16(DI, BaseMCReg, Ctx.LS); if ((State & BaseValueLive) != 0) { if (!PreservesBaseValue) { if (instructionReadsRegister(DI, BaseMCReg, Ctx.LS) && - !isTensorDescriptorUseOnly(DI, BaseMCReg, Ctx.LS)) + !isTensorDescriptorUseOnly(DI, BaseMCReg, Ctx.LS)) { + log() << "hotswap: tensor mask definition 0x" + << utohexstr(Ctx.Decoded[MaskIndex].Offset) + << " rejected: descriptor base read at 0x" + << utohexstr(DI.Offset) << " by " << DI.Mnemonic << "\n"; return false; + } if (instructionDefinesBase(DI, BaseMCReg, Ctx.LS)) State &= ~BaseValueLive; } @@ -1146,8 +1708,13 @@ bool isMaskDefinitionSafe(const PatchContext &Ctx, // A preserving writer propagates the changed base and may derive a changed // SCC value from it, so its SCC definition does not end the safety check. if ((State & SccValueLive) != 0) { - if (instReadsScc(DI.Inst, Desc, *Ctx.LS.MRI)) + if (instReadsScc(DI.Inst, Desc, *Ctx.LS.MRI)) { + log() << "hotswap: tensor mask definition 0x" + << utohexstr(Ctx.Decoded[MaskIndex].Offset) + << " rejected: changed SCC read at 0x" << utohexstr(DI.Offset) + << " by " << DI.Mnemonic << "\n"; return false; + } if (instWritesScc(DI.Inst, Desc, *Ctx.LS.MRI) && !PreservesBaseValue) State &= ~SccValueLive; } else if (PreservesBaseValue && @@ -1217,6 +1784,13 @@ TensorMaskDef findTensorMaskSetDefinitions(const PatchContext &Ctx, size_t InstIndex = Graph->BeginIndex + LocalIndex; const InternalDecodedInst &DI = Ctx.Decoded[InstIndex]; + if (Ctx.LS.MIA->isCall(DI.Inst)) { + log() << "hotswap: tensor mask definition search for tensor 0x" + << utohexstr(Tensor.Offset) << " rejected: call at 0x" + << utohexstr(DI.Offset) << "\n"; + return TensorMaskDef::NotApplicable; + } + bool PreservesBaseValue = writesBasePreservingZeroLow16(DI, BaseMCReg, Ctx.LS); if (instructionDefinesBase(DI, BaseMCReg, Ctx.LS)) { @@ -1227,14 +1801,24 @@ TensorMaskDef findTensorMaskSetDefinitions(const PatchContext &Ctx, } if (isClearedMaskAndOnBase(DI, BaseMCReg, Ctx.LS)) continue; + if (!PreservesBaseValue) + log() << "hotswap: tensor mask definition search for tensor 0x" + << utohexstr(Tensor.Offset) + << " rejected: non-preserving base definition at 0x" + << utohexstr(DI.Offset) << " by " << DI.Mnemonic << "\n"; if (!PreservesBaseValue) return TensorMaskDef::NotApplicable; } if (!PreservesBaseValue && instructionReadsRegister(DI, BaseMCReg, Ctx.LS) && - !isTensorDescriptorUseOnly(DI, BaseMCReg, Ctx.LS)) + !isTensorDescriptorUseOnly(DI, BaseMCReg, Ctx.LS)) { + log() << "hotswap: tensor mask definition search for tensor 0x" + << utohexstr(Tensor.Offset) + << " rejected: descriptor base read at 0x" << utohexstr(DI.Offset) + << " by " << DI.Mnemonic << "\n"; return TensorMaskDef::NotApplicable; + } bool HasReachablePredecessor = false; for (size_t Predecessor : Graph->Predecessors[LocalIndex]) { @@ -1384,6 +1968,13 @@ bool patchTensorLoadToLdsA0(PatchContext &Ctx, size_t Idx) { if (isAlreadyTensorMaskPatched(Ctx, Idx, BaseMCReg)) return false; + if (isTensorMaskAlreadyZero(Ctx, Idx, BaseMCReg)) { + log() << "hotswap: tensor_load_to_lds: descriptor workgroup_mask is " + "already zero on every path\n"; + DI.Mnemonic = ""; + return false; + } + SmallVector MaskSets; TensorMaskDef Result = findTensorMaskSetDefinitions(Ctx, Idx, BaseMCReg, MaskSets); diff --git a/amd/comgr/test-lit/hotswap-trampoline-tensor-mask-already-zero.s b/amd/comgr/test-lit/hotswap-trampoline-tensor-mask-already-zero.s new file mode 100644 index 0000000000000..c031bfcb70148 --- /dev/null +++ b/amd/comgr/test-lit/hotswap-trampoline-tensor-mask-already-zero.s @@ -0,0 +1,110 @@ +// COM: A tensor descriptor whose base has a provably zero low half needs no +// COM: rewrite. Both CFG paths seed s4 with zero and then use only self-writes +// COM: that cannot introduce a low bit. Foreign reads after the tensor are +// COM: harmless because this proof changes neither the definition nor the +// COM: tensor. A variant with one nonzero seed must reject the no-op proof and +// COM: use the save/clear/tensor/restore fallback. + +// 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 \ +// RUN: | %FileCheck --check-prefix=ZERO %s +// ZERO: descriptor workgroup_mask is already zero on every path +// ZERO: hotswap: applied 0 instruction patches +// ZERO: RESULT: SUCCESS +// RUN: cmp %t.elf %t.out.elf + +// COM: An unresolved set-PC is a possible alternate entry to the tensor. The +// COM: modeled zero path cannot prove anything about s4 on that hidden edge. +// RUN: sed 's/^\.set opaque_tensor_entry, 0$/.set opaque_tensor_entry, 1/' \ +// RUN: %s > %t.opaque.s +// RUN: %clang -target amdgcn-amd-amdhsa -mcpu=gfx1250 -nostdlib \ +// RUN: %t.opaque.s -o %t.opaque.elf +// RUN: env AMD_COMGR_EMIT_VERBOSE_LOGS=1 hotswap-rewrite %t.opaque.elf \ +// RUN: amdgcn-amd-amdhsa--gfx1250 amdgcn-amd-amdhsa--gfx1250 \ +// RUN: --output %t.opaque.out.elf 2>&1 \ +// RUN: | %FileCheck --check-prefix=NONZERO %s + +// RUN: sed 's/^\.set nonzero_path, 0$/.set nonzero_path, 1/' \ +// RUN: %s > %t.nonzero.s +// RUN: %clang -target amdgcn-amd-amdhsa -mcpu=gfx1250 -nostdlib \ +// RUN: %t.nonzero.s -o %t.nonzero.elf +// RUN: env AMD_COMGR_EMIT_VERBOSE_LOGS=1 hotswap-rewrite %t.nonzero.elf \ +// RUN: amdgcn-amd-amdhsa--gfx1250 amdgcn-amd-amdhsa--gfx1250 \ +// RUN: --output %t.nonzero.out.elf 2>&1 \ +// RUN: | %FileCheck --check-prefix=NONZERO %s +// NONZERO-NOT: descriptor workgroup_mask is already zero on every path +// NONZERO: tensor_load_to_lds: s4 live, save/restore via s12 +// NONZERO: RESULT: SUCCESS +// RUN: %llvm-objdump -d %t.nonzero.out.elf \ +// RUN: | %FileCheck --check-prefix=NONZERO-DIS %s +// NONZERO-DIS-LABEL: : +// NONZERO-DIS: s_and_b32 s4, s4, 0xfff7ffff +// NONZERO-DIS: s_mov_b32 s12, s4 +// NONZERO-DIS-NEXT: s_pack_hh_b32_b16 s4, 0, s4 +// NONZERO-DIS-NEXT: tensor_load_to_lds s[0:3], s[4:11] +// NONZERO-DIS-NEXT: s_mov_b32 s4, s12 + +.set nonzero_path, 0 +.set opaque_tensor_entry, 0 + +.amdgcn_target "amdgcn-amd-amdhsa--gfx1250" +.text +.globl test_tensor_mask_already_zero +.p2align 8 +.type test_tensor_mask_already_zero,@function +test_tensor_mask_already_zero: +.if opaque_tensor_entry + s_cbranch_scc0 .Lopaque_entry +.endif + s_cbranch_execz .Lsecond + s_mov_b32 s4, 0 + s_branch .Lbuild +.Lsecond: +.if nonzero_path + s_mov_b32 s4, 1 +.else + s_mov_b32 s4, 0 +.endif +.Lbuild: + s_and_b32 s4, s4, 0xfffcffff + s_or_b32 s4, s4, 0x10000 + s_and_b32 s4, s4, 0xfff7ffff + tensor_load_to_lds s[0:3], s[4:11] + s_mov_b32 s0, s4 + s_endpgm + .rept 8 + s_nop 0 + .endr +.if opaque_tensor_entry +.Lopaque_entry: + s_mov_b32 s4, 1 + s_set_pc_i64 s[10:11] +.endif +.Ltest_tensor_mask_already_zero_end: +.size test_tensor_mask_already_zero, .Ltest_tensor_mask_already_zero_end-test_tensor_mask_already_zero + +.rodata +.p2align 8 +.amdhsa_kernel test_tensor_mask_already_zero + .amdhsa_next_free_vgpr 1 + .amdhsa_next_free_sgpr 12 +.end_amdhsa_kernel + +.amdgpu_metadata + amdhsa.version: + - 3 + - 0 + amdhsa.kernels: + - .name: test_tensor_mask_already_zero + .symbol: test_tensor_mask_already_zero.kd + .sgpr_count: 12 + .vgpr_count: 1 + .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-tensor-mask-bounded-return.s b/amd/comgr/test-lit/hotswap-trampoline-tensor-mask-bounded-return.s new file mode 100644 index 0000000000000..e63eea7b086a2 --- /dev/null +++ b/amd/comgr/test-lit/hotswap-trampoline-tensor-mask-bounded-return.s @@ -0,0 +1,99 @@ +// COM: Tensor CFGs may consume exact target sets from the object-wide bounded +// COM: return audit. The default helper returns to an in-function continuation, +// COM: so its otherwise-indirect set-PC is a closed edge and the zero-mask +// COM: proof succeeds. The variant's proven return target is outside the +// COM: tensor function, which this local CFG rejects conservatively. + +// 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=LOCAL %s +// LOCAL: descriptor workgroup_mask is already zero on every path +// LOCAL: hotswap: applied 0 instruction patches +// LOCAL: RESULT: SUCCESS +// RUN: cmp %t.elf %t.out.elf + +// RUN: sed 's/^\.set outside_return, 0$/.set outside_return, 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: --output %t.outside.out.elf 2>&1 \ +// RUN: | %FileCheck --check-prefix=OUTSIDE %s +// OUTSIDE: tensor CFG rejected bounded transfer at 0x{{[0-9A-F]+}} with out-of-range target +// OUTSIDE-NOT: descriptor workgroup_mask is already zero on every path +// OUTSIDE: tensor_load_to_lds: s4 dead, no save/restore needed +// OUTSIDE: RESULT: SUCCESS + +.set outside_return, 0 + +.amdgcn_target "amdgcn-amd-amdhsa--gfx1250" +.text + +.if outside_return +.type tensor_bounded_return_callee,@function +tensor_bounded_return_callee: + s_mov_b32 s4, 0 + tensor_load_to_lds s[0:3], s[4:11] + s_mov_b32 s4, 0 + s_set_pc_i64 s[0:1] +.Ltensor_bounded_return_callee_end: +.size tensor_bounded_return_callee, .Ltensor_bounded_return_callee_end-tensor_bounded_return_callee + +.globl test_tensor_mask_bounded_return +.p2align 8 +.type test_tensor_mask_bounded_return,@function +test_tensor_mask_bounded_return: +.Loutside_getpc: + s_get_pc_i64 s[2:3] + // The aligned caller begins at .text+0x100; captured PC is .text+0x104. + s_add_nc_u64 s[2:3], s[2:3], -260 + s_swap_pc_i64 s[0:1], s[2:3] + s_endpgm +.Ltest_tensor_mask_bounded_return_end: +.size test_tensor_mask_bounded_return, .Ltest_tensor_mask_bounded_return_end-test_tensor_mask_bounded_return +.else +.globl test_tensor_mask_bounded_return +.p2align 8 +.type test_tensor_mask_bounded_return,@function +test_tensor_mask_bounded_return: + s_mov_b32 s4, 0 + tensor_load_to_lds s[0:3], s[4:11] + s_mov_b32 s4, 0 +.Llocal_getpc: + s_get_pc_i64 s[2:3] + // Captured PC is .text+0x18 and the helper begins at .text+0x24. + s_add_nc_u64 s[2:3], s[2:3], 12 + s_swap_pc_i64 s[0:1], s[2:3] +.Lcontinuation: + s_endpgm +.Lhelper: + s_set_pc_i64 s[0:1] +.Ltest_tensor_mask_bounded_return_end: +.size test_tensor_mask_bounded_return, .Ltest_tensor_mask_bounded_return_end-test_tensor_mask_bounded_return +.endif + +.rodata +.p2align 8 +.amdhsa_kernel test_tensor_mask_bounded_return + .amdhsa_next_free_vgpr 1 + .amdhsa_next_free_sgpr 12 +.end_amdhsa_kernel + +.amdgpu_metadata + amdhsa.version: + - 3 + - 0 + amdhsa.kernels: + - .name: test_tensor_mask_bounded_return + .symbol: test_tensor_mask_bounded_return.kd + .sgpr_count: 12 + .vgpr_count: 1 + .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-tensor-mask-local-setpc.s b/amd/comgr/test-lit/hotswap-trampoline-tensor-mask-local-setpc.s new file mode 100644 index 0000000000000..94e635b19b259 --- /dev/null +++ b/amd/comgr/test-lit/hotswap-trampoline-tensor-mask-local-setpc.s @@ -0,0 +1,156 @@ +// COM: A broad tensor function may contain a compiler-emitted reusable-PC +// COM: tail jump. Resolve only the exact adjacent get-PC/carry-add/set-PC +// COM: sequence to its in-range instruction boundary. Both the tensor path and +// COM: the jump target kill s4 after the tensor mask definition, so the +// COM: definition-time low16 clear is safe. + +// 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 \ +// RUN: | %FileCheck --check-prefix=API %s +// API: cleared workgroup_mask at descriptor definition 0x{{[0-9A-F]+}} (s4) +// API-NOT: tensor_load_to_lds: s4 +// API: RESULT: SUCCESS +// RUN: %llvm-objdump -d %t.out.elf | %FileCheck --check-prefix=DIS %s +// DIS-LABEL: : +// DIS: s_and_b32 s4, s4, 0xfff70000 +// DIS-NOT: s_pack_hh_b32_b16 +// DIS: tensor_load_to_lds s[0:3], s[4:11] +// 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 wrong_pair, 0 +.set intervening_def, 0 +.set outside_target, 0 +.set interior_entry, 0 +.set declared_entry, 0 + +// COM: A wrong transfer pair, an intervening target-pair definition, or an +// COM: otherwise exact sequence whose target is outside the function must +// COM: remain opaque. Each variant therefore rejects the definition rewrite +// COM: and uses the at-site fallback. +// RUN: sed 's/^\.set wrong_pair, 0$/.set wrong_pair, 1/' \ +// RUN: %s > %t.wrong-pair.s +// RUN: %clang -target amdgcn-amd-amdhsa -mcpu=gfx1250 -nostdlib \ +// RUN: %t.wrong-pair.s -o %t.wrong-pair.elf +// RUN: env AMD_COMGR_EMIT_VERBOSE_LOGS=1 hotswap-rewrite %t.wrong-pair.elf \ +// RUN: amdgcn-amd-amdhsa--gfx1250 amdgcn-amd-amdhsa--gfx1250 \ +// RUN: --output %t.wrong-pair.out.elf 2>&1 \ +// RUN: | %FileCheck --check-prefix=OPAQUE %s +// RUN: sed 's/^\.set intervening_def, 0$/.set intervening_def, 1/' \ +// RUN: %s > %t.intervening.s +// RUN: %clang -target amdgcn-amd-amdhsa -mcpu=gfx1250 -nostdlib \ +// RUN: %t.intervening.s -o %t.intervening.elf +// RUN: env AMD_COMGR_EMIT_VERBOSE_LOGS=1 hotswap-rewrite %t.intervening.elf \ +// RUN: amdgcn-amd-amdhsa--gfx1250 amdgcn-amd-amdhsa--gfx1250 \ +// RUN: --output %t.intervening.out.elf 2>&1 \ +// RUN: | %FileCheck --check-prefix=OPAQUE %s +// RUN: sed 's/^\.set outside_target, 0$/.set outside_target, 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: --output %t.outside.out.elf 2>&1 \ +// RUN: | %FileCheck --check-prefix=OPAQUE %s +// RUN: sed 's/^\.set interior_entry, 0$/.set interior_entry, 1/' \ +// RUN: %s > %t.interior.s +// RUN: %clang -target amdgcn-amd-amdhsa -mcpu=gfx1250 -nostdlib \ +// RUN: %t.interior.s -o %t.interior.elf +// RUN: env AMD_COMGR_EMIT_VERBOSE_LOGS=1 hotswap-rewrite %t.interior.elf \ +// RUN: amdgcn-amd-amdhsa--gfx1250 amdgcn-amd-amdhsa--gfx1250 \ +// RUN: --output %t.interior.out.elf 2>&1 \ +// RUN: | %FileCheck --check-prefix=OPAQUE %s +// RUN: sed 's/^\.set declared_entry, 0$/.set declared_entry, 1/' \ +// RUN: %s > %t.declared.s +// RUN: %clang -target amdgcn-amd-amdhsa -mcpu=gfx1250 -nostdlib \ +// RUN: %t.declared.s -o %t.declared.elf +// RUN: env AMD_COMGR_EMIT_VERBOSE_LOGS=1 hotswap-rewrite %t.declared.elf \ +// RUN: amdgcn-amd-amdhsa--gfx1250 amdgcn-amd-amdhsa--gfx1250 \ +// RUN: --output %t.declared.out.elf 2>&1 \ +// RUN: | %FileCheck --check-prefix=OPAQUE %s +// OPAQUE-NOT: cleared workgroup_mask at descriptor definition +// OPAQUE: tensor_load_to_lds: s4 dead, no save/restore needed +// OPAQUE: RESULT: SUCCESS + +.amdgcn_target "amdgcn-amd-amdhsa--gfx1250" +.text +.globl test_tensor_mask_local_setpc +.p2align 8 +.type test_tensor_mask_local_setpc,@function +test_tensor_mask_local_setpc: + s_and_b32 s4, s4, 0xfff7ffff + s_cmp_eq_u32 s0, s0 +.if interior_entry + s_cbranch_scc0 .Ladd_low +.endif + s_cbranch_execz .Ltensor +.Lgetpc: + s_get_pc_i64 s[70:71] +.if outside_target + s_add_co_i32 s72, .Loutside_target-(.Lgetpc+4)-4, 4 +.else + s_add_co_i32 s72, .Ltarget-(.Lgetpc+4)-4, 4 +.endif +.Ladd_low: +.if declared_entry +.globl test_tensor_mask_local_setpc_interior +.type test_tensor_mask_local_setpc_interior,@function +test_tensor_mask_local_setpc_interior: +.endif + s_add_co_u32 s70, s70, s72 +.if declared_entry +.size test_tensor_mask_local_setpc_interior, .-test_tensor_mask_local_setpc_interior +.endif + s_add_co_ci_u32 s71, s71, 0 +.if intervening_def + s_mov_b32 s70, 0 +.endif +.if wrong_pair + s_set_pc_i64 s[74:75] +.else + s_set_pc_i64 s[70:71] +.endif +.Ltensor: + tensor_load_to_lds s[0:3], s[4:11] + s_mov_b32 s4, 0 + s_cmp_eq_u32 s0, s0 + s_endpgm +.Ltarget: + s_mov_b32 s4, 0 + s_cmp_eq_u32 s0, s0 + s_endpgm +.Ltest_tensor_mask_local_setpc_end: +.size test_tensor_mask_local_setpc, .Ltest_tensor_mask_local_setpc_end-test_tensor_mask_local_setpc + +.Loutside: + s_nop 0 +.Loutside_target: + s_endpgm + +.rodata +.p2align 8 +.amdhsa_kernel test_tensor_mask_local_setpc + .amdhsa_next_free_vgpr 1 + .amdhsa_next_free_sgpr 76 +.end_amdhsa_kernel + +.amdgpu_metadata + amdhsa.version: + - 3 + - 0 + amdhsa.kernels: + - .name: test_tensor_mask_local_setpc + .symbol: test_tensor_mask_local_setpc.kd + .sgpr_count: 76 + .vgpr_count: 1 + .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-tensor-mask-signed-setpc.s b/amd/comgr/test-lit/hotswap-trampoline-tensor-mask-signed-setpc.s new file mode 100644 index 0000000000000..db6fb8cc611c1 --- /dev/null +++ b/amd/comgr/test-lit/hotswap-trampoline-tensor-mask-signed-setpc.s @@ -0,0 +1,92 @@ +// COM: Recognize only Tensile's exact contiguous signed-direction reusable-PC +// COM: transfer. Both arithmetic arms compute the same local target. A direct +// COM: entry into the positive arm invalidates that target proof and forces +// COM: the tensor's at-site fallback. + +// 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=LOCAL %s +// LOCAL: cleared workgroup_mask at descriptor definition 0x{{[0-9A-F]+}} (s4) +// LOCAL-NOT: tensor_load_to_lds: s4 +// LOCAL: RESULT: SUCCESS +// RUN: %llvm-objdump -d %t.out.elf | %FileCheck --check-prefix=DIS %s +// DIS: s_and_b32 s4, s4, 0xfff70000 +// DIS-NOT: s_pack_hh_b32_b16 +// DIS: tensor_load_to_lds s[0:3], s[4:11] + +// RUN: sed 's/^\.set alternate_arm_entry, 0$/.set alternate_arm_entry, 1/' \ +// RUN: %s > %t.alternate.s +// RUN: %clang -target amdgcn-amd-amdhsa -mcpu=gfx1250 -nostdlib \ +// RUN: %t.alternate.s -o %t.alternate.elf +// RUN: env AMD_COMGR_EMIT_VERBOSE_LOGS=1 hotswap-rewrite %t.alternate.elf \ +// RUN: amdgcn-amd-amdhsa--gfx1250 amdgcn-amd-amdhsa--gfx1250 \ +// RUN: --output %t.alternate.out.elf 2>&1 \ +// RUN: | %FileCheck --check-prefix=ALTERNATE %s +// ALTERNATE: tensor CFG rejected alternate entry into reusable-PC sequence +// ALTERNATE-NOT: cleared workgroup_mask at descriptor definition +// ALTERNATE: tensor_load_to_lds: s4 dead, no save/restore needed +// ALTERNATE: RESULT: SUCCESS + +.set alternate_arm_entry, 0 + +.amdgcn_target "amdgcn-amd-amdhsa--gfx1250" +.text +.globl test_tensor_mask_signed_setpc +.p2align 8 +.type test_tensor_mask_signed_setpc,@function +test_tensor_mask_signed_setpc: + s_and_b32 s4, s4, 0xfff7ffff + s_cmp_eq_u32 s0, s0 +.if alternate_arm_entry + s_cbranch_scc0 .Lpositive +.endif + s_cbranch_execz .Ltensor +.Lgetpc: + s_get_pc_i64 s[70:71] + s_add_co_i32 s72, .Ltarget-(.Lgetpc+4), 0 + s_cmp_ge_i32 s72, 0 + s_cbranch_scc1 .Lpositive + s_abs_i32 s72, s72 + s_sub_co_u32 s70, s70, s72 + s_sub_co_ci_u32 s71, s71, 0 + s_set_pc_i64 s[70:71] +.Lpositive: + s_add_co_u32 s70, s70, s72 + s_add_co_ci_u32 s71, s71, 0 + s_set_pc_i64 s[70:71] +.Ltensor: + tensor_load_to_lds s[0:3], s[4:11] + s_mov_b32 s4, 0 + s_cmp_eq_u32 s0, s0 + s_endpgm +.Ltarget: + s_mov_b32 s4, 0 + s_cmp_eq_u32 s0, s0 + s_endpgm +.Ltest_tensor_mask_signed_setpc_end: +.size test_tensor_mask_signed_setpc, .Ltest_tensor_mask_signed_setpc_end-test_tensor_mask_signed_setpc + +.rodata +.p2align 8 +.amdhsa_kernel test_tensor_mask_signed_setpc + .amdhsa_next_free_vgpr 1 + .amdhsa_next_free_sgpr 73 +.end_amdhsa_kernel + +.amdgpu_metadata + amdhsa.version: + - 3 + - 0 + amdhsa.kernels: + - .name: test_tensor_mask_signed_setpc + .symbol: test_tensor_mask_signed_setpc.kd + .sgpr_count: 73 + .vgpr_count: 1 + .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