[comgr][hotswap] Close finite materialized-call components

amd/comgr/src/comgr-hotswap-b0a0.cpp

@@ -4223,17 +4223,22 @@ static FiniteControlFlowAudit auditFiniteIndirectControlFlow(
     if (Reachable.test(SetPc) && !BoundedSetPc.test(SetPc))
       markUnboundedIndirectEntry();
 
-  for (size_t InstIndex : Index.UnboundedIndirectIndices)
-    if (Reachable.test(InstIndex))
-      markUnboundedIndirectEntry();
-
   // Every call is also an indirect entry source until either a finite local
   // target or a finite external target has been recorded for it.
   BitVector FiniteCalls(Decoded.size());
   for (const KnownCallSite &Call : Index.Calls)
     FiniteCalls.set(Call.InstIndex);
   for (const ExternalCallContinuation &Call : Index.ExternalCallContinuations)
     FiniteCalls.set(Call.InstIndex);
+
+  // MC may also classify a register call as an indirect branch. Do not let
+  // that generic classification create a false unbounded self-edge after the
+  // exact call target and continuation have been admitted to this same joint
+  // audit. Unknown calls remain unbounded in the call-specific loop below.
+  for (size_t InstIndex : Index.UnboundedIndirectIndices)
+    if (Reachable.test(InstIndex) && !FiniteCalls.test(InstIndex))
+      markUnboundedIndirectEntry();
+
   for (size_t InstIndex : Index.BranchOrCallIndices) {
     if (!Reachable.test(InstIndex) || !LS.MIA->isCall(Decoded[InstIndex].Inst))
       continue;
@@ -4258,6 +4263,7 @@ hasKnownControlFlowEntry(ArrayRef<uint64_t> DeclaredEntries,
                          ArrayRef<BoundedSetPcReturn> BoundedReturns,
                          const DenseMap<size_t, size_t> &BoundedReturnPositions,
                          const ControlFlowScanIndex &Index,
+                         bool HasUnboundedIndirectEntries,
                          uint64_t SequenceStart, uint64_t SequenceEnd) {
   for (uint64_t Entry : DeclaredEntries)
     if (Entry > SequenceStart && Entry <= SequenceEnd)
@@ -4277,9 +4283,19 @@ hasKnownControlFlowEntry(ArrayRef<uint64_t> DeclaredEntries,
   // 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)
+  if (HasUnboundedIndirectEntries)
     return true;
 
+  auto EntersSequence = [&](uint64_t Target) {
+    return Target > SequenceStart && Target <= SequenceEnd;
+  };
+  for (const KnownCallSite &Call : Index.Calls)
+    if (EntersSequence(Call.Target) || EntersSequence(Call.Continuation))
+      return true;
+  for (const ExternalCallContinuation &Call : Index.ExternalCallContinuations)
+    if (EntersSequence(Call.Continuation))
+      return true;
+
   SmallVector<DirectTargetSource, 16>::const_iterator First =
       llvm::upper_bound(Index.DirectTargetsByTarget, SequenceStart,
                         [](uint64_t Target, const DirectTargetSource &Source) {
@@ -5493,6 +5509,7 @@ static std::optional<WellFormedAbiEntrySet> validateWellFormedAbiEntrySet(
   // Phase 2 may now authorize only the exact source indices proven above
   // while validating ordinary s30-return functions.
   DenseSet<size_t> CanonicalAbiReturns;
+  DenseSet<uint64_t> CanonicalReturningFunctionBegins;
   for (const auto &Entry : ReturnsByFunction) {
     const ElfView::FunctionTextRange *Range =
         ReturnFunctionRanges.lookup(Entry.first);
@@ -5508,6 +5525,11 @@ static std::optional<WellFormedAbiEntrySet> validateWellFormedAbiEntrySet(
             << " does not use a provable canonical s30 call frame\n";
       return std::nullopt;
     }
+    for (size_t Return : Entry.second)
+      if (CanonicalAbiReturns.contains(Return)) {
+        CanonicalReturningFunctionBegins.insert(Entry.first.first - TextAddr);
+        break;
+      }
   }
 
   bool SawAbiCall = false;
@@ -5606,12 +5628,57 @@ static std::optional<WellFormedAbiEntrySet> validateWellFormedAbiEntrySet(
       return std::nullopt;
     }
   }
-  if (!SawAbiCall || !SawAbiReturn) {
+  bool SawExactCanonicalCall = false;
+  if (!SawAbiCall && SawAbiReturn)
+    for (const KnownCallSite &Call : Index.Calls)
+      if (Index.MaterializedCalls.contains(Call.InstIndex) &&
+          Call.ReturnRegister == AbiLinkPair &&
+          CanonicalReturningFunctionBegins.contains(Call.Target)) {
+        SawExactCanonicalCall = true;
+        break;
+      }
+  if (!SawAbiReturn || (!SawAbiCall && !SawExactCanonicalCall)) {
     log() << "hotswap: linked-code-object ABI entry-set fallback rejected: "
              "no opaque s30 call/return pair\n";
     return std::nullopt;
   }
 
+  if (!SawAbiCall) {
+    auto HasMaterializationEntry =
+        [&](const PcMaterializedCallInfo &Materialized) {
+          auto IsInterior = [&](uint64_t Offset) {
+            return Offset > Materialized.SequenceStart &&
+                   Offset <= Materialized.SequenceEnd;
+          };
+          for (uint64_t Entry : Result.Targets)
+            if (IsInterior(Entry))
+              return true;
+          for (const DirectTargetSource &Source : Index.DirectTargetsByTarget)
+            if (IsInterior(Source.Target))
+              return true;
+          for (const KnownCallSite &Call : Index.Calls)
+            if (IsInterior(Call.Target) || IsInterior(Call.Continuation))
+              return true;
+          for (const auto &Call : CallContinuations)
+            if (IsInterior(Call.second))
+              return true;
+          for (const FiniteSetPcTransfer *Transfer : SelectedSetPcs)
+            if (Transfer->LocalTargetIndex &&
+                IsInterior(Decoded[*Transfer->LocalTargetIndex].Offset))
+              return true;
+          return false;
+        };
+    for (const auto &Entry : Index.MaterializedCalls)
+      if (Result.Calls.contains(Entry.first) &&
+          HasMaterializationEntry(Entry.second)) {
+        log() << "hotswap: exact materialized-call/canonical-return closure "
+                 "rejected: alternate entry inside call materialization "
+                 "ending at 0x"
+              << utohexstr(Decoded[Entry.first].Offset) << "\n";
+        return std::nullopt;
+      }
+  }
+
   auto hasInteriorEntry = [&](const FiniteSetPcTransfer &Candidate) {
     uint64_t Begin = Decoded[Candidate.SequenceBeginIndex].Offset;
     std::optional<uint64_t> End =
@@ -5654,11 +5721,12 @@ static std::optional<WellFormedAbiEntrySet> validateWellFormedAbiEntrySet(
 
   for (const auto &Call : CallContinuations)
     Result.Targets.insert(Call.second);
-  log() << "hotswap: accepted well-formed linked-code-object ABI entry set "
-           "for "
-        << Result.Calls.size() << " register call(s), " << Result.SetPcs.size()
-        << " set-PC transfer(s), and " << Result.Targets.size()
-        << " finite local entry point(s)\n";
+  log() << "hotswap: accepted "
+        << (SawAbiCall ? "well-formed linked-code-object ABI entry set"
+                       : "exact materialized-call/canonical-return closure")
+        << " for " << Result.Calls.size() << " register call(s), "
+        << Result.SetPcs.size() << " set-PC transfer(s), and "
+        << Result.Targets.size() << " finite local entry point(s)\n";
   return Result;
 }
 
@@ -5892,7 +5960,8 @@ std::optional<DirectControlFlowInfo> collectDirectBranchTargets(
       continue;
     if (hasKnownControlFlowEntry(
             DeclaredEntries, BoundedReturns, BoundedReturnPositions, *Index,
-            Entry.second.SequenceStart, Entry.second.SequenceEnd)) {
+            HasUnboundedIndirectEntries, Entry.second.SequenceStart,
+            Entry.second.SequenceEnd)) {
       ReusableCalls[I].clear();
       continue;
     }
@@ -5953,10 +6022,10 @@ std::optional<DirectControlFlowInfo> collectDirectBranchTargets(
         DenseMap<size_t, PcMaterializedCallInfo>::const_iterator Materialized =
             Index->MaterializedCalls.find(InstIndex);
         if (Materialized != Index->MaterializedCalls.end() &&
-            !hasKnownControlFlowEntry(DeclaredEntries, BoundedReturns,
-                                      BoundedReturnPositions, *Index,
-                                      Materialized->second.SequenceStart,
-                                      Materialized->second.SequenceEnd))
+            !hasKnownControlFlowEntry(
+                DeclaredEntries, BoundedReturns, BoundedReturnPositions, *Index,
+                HasUnboundedIndirectEntries, Materialized->second.SequenceStart,
+                Materialized->second.SequenceEnd))
           Target = Materialized->second.Target;
       }
       if (!ReusableCalls[InstIndex].empty()) {

amd/comgr/test-lit/hotswap-finite-materialized-call-closure.s

@@ -0,0 +1,65 @@
+// COM: An exact PC-materialized singleton call enters a defined local helper.
+// COM: The helper uses s[30:31] as scratch only after saving the incoming link
+// COM: in callee-saved VGPR lanes, then restores it before returning. Prove the
+// COM: call, canonical return, and call continuation together without falling
+// COM: back to an object-wide unknown indirect-entry assumption.
+
+// RUN: %clang -target amdgcn-amd-amdhsa -mcpu=gfx1250 -nostdlib %s -o %t.elf
+// RUN: env AMD_COMGR_EMIT_VERBOSE_LOGS=1 hotswap-rewrite %t.elf \
+// RUN:   amdgcn-amd-amdhsa--gfx1250:gfx1250-b0-specific+ \
+// RUN:   amdgcn-amd-amdhsa--gfx1250:gfx1250-b0-specific- \
+// RUN:   --output %t.out.elf 2>&1 | %FileCheck --check-prefix=LOG %s
+// LOG: hotswap: accepted exact materialized-call/canonical-return closure for 1 register call(s)
+// LOG-NOT: hotswap: unresolved call target
+// LOG-NOT: hotswap: unresolved control-flow target disables
+// LOG: RESULT: SUCCESS
+
+.amdgcn_target "amdgcn-amd-amdhsa--gfx1250"
+.text
+
+.globl exact_closure_kernel
+.type exact_closure_kernel,@function
+exact_closure_kernel:
+  s_get_pc_i64 s[0:1]
+  // PC after get-PC is 4; exact_closure_helper begins at 16.
+  s_add_nc_u64 s[0:1], s[0:1], 12
+  s_swap_pc_i64 s[30:31], s[0:1]
+  s_endpgm
+.size exact_closure_kernel, .-exact_closure_kernel
+
+.local exact_closure_helper
+.type exact_closure_helper,@function
+exact_closure_helper:
+  v_writelane_b32 v40, s30, 0
+  v_writelane_b32 v41, s31, 1
+  s_mov_b32 s30, 0
+  s_mov_b32 s31, 0
+  v_readlane_b32 s30, v40, 0
+  v_readlane_b32 s31, v41, 1
+  s_set_pc_i64 s[30:31]
+  s_endpgm
+.size exact_closure_helper, .-exact_closure_helper
+
+.rodata
+.p2align 8
+.amdhsa_kernel exact_closure_kernel
+  .amdhsa_next_free_vgpr 42
+  .amdhsa_next_free_sgpr 42
+.end_amdhsa_kernel
+
+.amdgpu_metadata
+  amdhsa.version:
+    - 3
+    - 0
+  amdhsa.kernels:
+    - .name: exact_closure_kernel
+      .symbol: exact_closure_kernel.kd
+      .sgpr_count: 42
+      .vgpr_count: 42
+      .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

amd/comgr/test-lit/hotswap-materialized-call-joint-gateway.s

@@ -0,0 +1,101 @@
+// COM: The production corpus contains compiler PC-materialized s30 calls into
+// COM: a helper with exact materialized set-PC jumps and an s30 return. Those
+// COM: three edge families form one finite component. MC also classifies the
+// COM: swap-call as an indirect branch; that generic classification must not
+// COM: create an unbounded self-edge after the same call has been proven
+// COM: finite. Closing the component keeps external zero padding available as
+// COM: a gateway for an otherwise-stranded far eight-byte DS2 patch.
+
+// 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-prefixes=LOG,API %s
+// LOG: hotswap: accepted exact materialized-call/canonical-return closure
+// LOG-NOT: hotswap: unresolved call target
+// LOG-NOT: hotswap: unresolved control-flow target disables
+// API: RESULT: SUCCESS
+
+// RUN: %llvm-objdump -d %t.out.elf | %FileCheck --check-prefix=DISASM %s
+// DISASM-NOT: ds_load_2addr
+// DISASM: ds_load_b64
+// DISASM-NEXT: ds_load_b64
+
+// 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
+
+.local joint_helper
+.type joint_helper,@function
+joint_helper:
+  // Match the production helper: preserve its incoming link in callee-saved
+  // VGPR lanes, use s30:s31 as scratch, then restore it before returning.
+  v_writelane_b32 v40, s30, 0
+  v_writelane_b32 v41, s31, 1
+  s_mov_b32 s30, 0
+  s_mov_b32 s31, 0
+
+  // This exact local set-PC jump is needed to reach the restore/return path.
+  s_get_pc_i64 s[4:5]
+.Lhelper_pc:
+  s_add_nc_u64 s[4:5], s[4:5], .Lhelper_restore-.Lhelper_pc
+  s_set_pc_i64 s[4:5]
+  s_endpgm
+.Lhelper_restore:
+  v_readlane_b32 s30, v40, 0
+  v_readlane_b32 s31, v41, 1
+  s_set_pc_i64 s[30:31]
+  s_endpgm
+.size joint_helper, .-joint_helper
+
+.globl test_materialized_call_joint_gateway
+.p2align 8
+.type test_materialized_call_joint_gateway,@function
+test_materialized_call_joint_gateway:
+  ds_load_2addr_stride64_b64 v[0:3], v4 offset0:1 offset1:2
+
+  // The call target is the exact local helper entry. Its continuation is the
+  // following s_endpgm instruction.
+  s_get_pc_i64 s[0:1]
+.Lcaller_pc:
+  s_add_nc_u64 s[0:1], s[0:1], joint_helper-.Lcaller_pc
+  s_swap_pc_i64 s[30:31], s[0:1]
+  s_endpgm
+.size test_materialized_call_joint_gateway, .-test_materialized_call_joint_gateway
+
+// Safe external gateway space. A falsely-unbounded call clears this map and
+// makes the far eight-byte patch fail closed.
+.fill 64, 1, 0
+
+// Keep the appended trampoline pool beyond one signed s_branch span.
+.rept 40000
+  s_mov_b32 s0, s1
+.endr
+
+.rodata
+.p2align 8
+.amdhsa_kernel test_materialized_call_joint_gateway
+  .amdhsa_next_free_vgpr 42
+  .amdhsa_next_free_sgpr 42
+.end_amdhsa_kernel
+
+.amdgpu_metadata
+  amdhsa.version:
+    - 3
+    - 0
+  amdhsa.kernels:
+    - .name: test_materialized_call_joint_gateway
+      .symbol: test_materialized_call_joint_gateway.kd
+      .sgpr_count: 42
+      .vgpr_count: 42
+      .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