minseobshin11 · Oct 22, 2026
@@ -84,6 +84,7 @@ experimental/range_replay.h           public payload struct, status enum, begin/
         |     range_state.cpp         eligibility bookkeeping, the recorder, decline channels
         |     replay_callbacks.cpp    CONFIG / PASS / CLOSE delivery, pass-count decisions
         |     executor.cpp            snapshot, kernarg staging, the pass loop, restore
+        |     untracked_memory.cpp    kernel arguments that point outside the snapshot
         |     retained_kernarg.cpp    the kernarg block kept for each agent's next range
         |     digest.cpp              region hashing and divergence counting
         |
@@ -107,7 +108,7 @@ Range replay v1 is scoped to what the existing snapshot can restore correctly.
 | Not supported in v1 | Consequence |
 |---|---|
 | More than one queue or agent per range | declined (`MULTI_QUEUE` / `MULTI_AGENT`) |
-| Unified / managed memory, host memory, fine-grained memory, VM-mapped allocations | not restored between passes; not detected — see the divergence check |
+| Unified / managed memory, host memory, fine-grained memory, VM-mapped allocations | not restored between passes; declined (`SNAPSHOT_FAILED`) when an explicit kernel argument points at them, otherwise not detected — see the divergence check |
 | Host state | never rewound; a range whose kernels consume values the host recomputed inside it is not a replay candidate |
 | HIP graph launches inside a range | declined (`GRAPH_LAUNCH`) |
 | Code objects loaded or unloaded inside a range (lazy module loading, JIT, `hipModuleUnload`) | declined (`CODE_OBJECT_CHANGED_IN_RANGE`); run a range once before measuring it |

@@ -25,7 +25,8 @@ are skipped, and the reason arrives as the `CLOSE` callback's `status`.
 | Kernarg bytes available at replay time | HIP recycles a kernarg block as soon as its kernel completes | bytes copied at record time |
 | Recorded code stays loaded, and its module variables are in the snapshot | An unload frees code a recorded packet points at; a code object loaded after the entry snapshot has `__device__` / `__constant__` globals the snapshot never saw | declined |
 | Dispatch packets re-submittable | The recorded packets must reach the GPU through the same path the application used | requires the interception path; queue interposition declines |
-| Everything read is in the snapshot | Unified memory, host memory, and host-side state are not captured | **not checkable** — see the divergence check |
+| Everything written is in the snapshot | Unified memory, host memory, fine-grained and VM-mapped allocations are not captured, so a pass would apply its writes to them again | declined when a dispatch's explicit kernel arguments point at such memory; pointers stored inside other memory are **not checkable** |
+| Everything read is in the snapshot | Host-side state is not captured | **not checkable** — see the divergence check |
 
 ## The decline decision table
 
@@ -45,6 +46,7 @@ GPU dependencies so the table below can be unit tested directly.
 | Any code object is unloaded while the range is bound | `CODE_OBJECT_CHANGED_IN_RANGE` | cross-thread, folded in |
 | Device allocation created or freed inside the range | `ALLOCATION_CHANGED_IN_RANGE` | close time |
 | Snapshot capture or a later snapshot failed | `SNAPSHOT_FAILED` | snapshot time |
+| A dispatch's kernel arguments point at memory the snapshot does not cover | `SNAPSHOT_FAILED` | submission time |
 | Kernarg staging allocation failed | `STAGING_FAILED` | close time |
 | Queue interposition is enabled | `UNSUPPORTED_QUEUE_PATH` | `begin` |
 | Nothing was recorded | `NO_DISPATCH` | close time |
@@ -161,9 +163,11 @@ timings.
 
 ## The divergence check
 
-The checks above cover interference the SDK can see. They cannot cover a range that reads state the
-snapshot does not capture — unified or managed memory, fine-grained or host memory,
-virtual-memory-mapped allocations, or a value the host recomputed inside the range. Such a range
+The checks above cover interference the SDK can see. A range whose explicit kernel arguments point
+at memory the snapshot does not capture — unified or managed memory, fine-grained or host memory,
+virtual-memory-mapped allocations — is declined (`SNAPSHOT_FAILED`) when the dispatch is recorded.
+What remains unseen is a pointer to such memory stored inside other memory or packed into a by-value
+word that is not 8-byte aligned, and a value the host recomputed inside the range. Such a range
 passes every check and is replayed, and its passes quietly describe different inputs.
 
 Setting `ROCPROF_RANGE_REPLAY_VERIFY` turns on a check for exactly that residue. The executor hashes
@@ -234,6 +238,7 @@ All paths are relative to `projects/rocprofiler-sdk/`.
 | Open-range fast gate | `source/lib/rocprofiler-sdk/range_replay/range_state.cpp` | `any_range_open()` |
 | Replay window and pass loop | `source/lib/rocprofiler-sdk/range_replay/executor.cpp` | `execute_range()` |
 | Entry snapshot | `source/lib/rocprofiler-sdk/range_replay/executor.cpp` | `ensure_entry_snapshot()` |
+| Kernel arguments outside the snapshot | `source/lib/rocprofiler-sdk/range_replay/untracked_memory.cpp` | `decline_untracked()` |
 | Kernarg staging | `source/lib/rocprofiler-sdk/range_replay/executor.cpp` | `kernarg_staging` |
 | Kernarg block retention | `source/lib/rocprofiler-sdk/range_replay/retained_kernarg.cpp` | `kernarg_block_fits()`, `take_retained_kernarg_block()`, `retain_kernarg_block()` |
 | Divergence hashing | `source/lib/rocprofiler-sdk/range_replay/digest.cpp` | `hash_bytes()`, `count_divergent()` |

@@ -59,6 +59,8 @@ typedef enum rocprofiler_range_replay_status_t  // NOLINT(performance-enum-size)
                                                           ///< agent inside the range
     ROCPROFILER_RANGE_REPLAY_STATUS_PROGRAM_TOO_LARGE,    ///< Range exceeded the record budget
     ROCPROFILER_RANGE_REPLAY_STATUS_SNAPSHOT_FAILED,      ///< Device-memory snapshot was incomplete
+                                                          ///< or does not cover memory a kernel
+                                                          ///< argument points at
     ROCPROFILER_RANGE_REPLAY_STATUS_STAGING_FAILED,  ///< Kernarg staging or signal setup failed
     ROCPROFILER_RANGE_REPLAY_STATUS_ALLOCATION_CHANGED_IN_RANGE,  ///< A device allocation was made
                                                                   ///< or freed inside the range
@@ -101,7 +103,10 @@ typedef enum rocprofiler_range_replay_status_t  // NOLINT(performance-enum-size)
  * while it is open. Repeatability rests on the same
  * device-memory snapshot kernel replay uses (coarse-grained device allocations owned by the agent
  * plus module-scope @c __device__ / @c __constant__ variables), so unified/managed memory, host and
- * fine-grained memory, and virtual-memory-mapped allocations are not restored between passes.
+ * fine-grained memory, and virtual-memory-mapped allocations are not restored between passes. A
+ * range whose explicit kernel arguments point at such memory is declined
+ * (@c ROCPROFILER_RANGE_REPLAY_STATUS_SNAPSHOT_FAILED); a pointer stored inside other memory is
+ * not seen.
  * Replayed passes are serialized dispatch-by-dispatch, so concurrency-sensitive measurements differ
  * from pass 0. Host state is not rewound at all: a range whose kernels consume values the host
  * recomputed inside the range is not a candidate for replay.

@@ -1,10 +1,10 @@
 #
 set(ROCPROFILER_LIB_RANGE_REPLAY_HEADERS
     digest.hpp executor.hpp queue_hooks.hpp range_replay.hpp range_state.hpp
-    replay_callbacks.hpp retained_kernarg.hpp)
+    replay_callbacks.hpp retained_kernarg.hpp untracked_memory.hpp)
 set(ROCPROFILER_LIB_RANGE_REPLAY_SOURCES
     api.cpp digest.cpp executor.cpp queue_hooks.cpp range_replay.cpp range_state.cpp
-    replay_callbacks.cpp retained_kernarg.cpp)
+    replay_callbacks.cpp retained_kernarg.cpp untracked_memory.cpp)
 
 target_sources(
     rocprofiler-sdk-object-library PRIVATE ${ROCPROFILER_LIB_RANGE_REPLAY_SOURCES}

@@ -44,6 +44,7 @@ submission_hook(const hsa::Queue*                     queue,
     {
         note_submission(_queue, packets, pkt_count, graph_launch_active);
         ensure_entry_snapshot(*open_range, _queue, writer);
+        decline_untracked(*open_range);
     }
     else if(const auto* rocp_agent = _queue.get_agent().get_rocp_agent(); rocp_agent != nullptr)
     {

@@ -39,6 +39,7 @@
 #include "lib/rocprofiler-sdk/hsa/rocprofiler_packet.hpp"
 #include "lib/rocprofiler-sdk/kernel_replay/memory_snapshot.hpp"
 #include "lib/rocprofiler-sdk/range_replay/replay_callbacks.hpp"
+#include "lib/rocprofiler-sdk/range_replay/untracked_memory.hpp"
 
 #include <rocprofiler-sdk/agent.h>
 #include <rocprofiler-sdk/experimental/range_replay.h>
@@ -50,6 +51,7 @@
 #include <cstddef>
 #include <cstdint>
 #include <memory>
+#include <unordered_set>
 #include <vector>
 
 namespace rocprofiler
@@ -166,6 +168,12 @@ struct range_context_t
     // restore, so it is declined rather than replayed against partially-restored inputs.
     std::vector<void*> snapshot_ptrs = {};
 
+    // Untracked-memory check (untracked_memory.hpp): the snapshot's spans, kernarg values already
+    // looked up, and how many recorded dispatches have been checked.
+    std::vector<untracked_memory::address_span_t> covered           = {};
+    std::unordered_set<uint64_t>                  seen              = {};
+    size_t                                        untracked_scanned = 0;
+
     std::shared_ptr<external_decline_t> external = {};
 
     // What the tool asked for at CONFIG, and the contexts this range's PASS/CLOSE callbacks go to.

@@ -9,7 +9,7 @@ include(GoogleTest)
 # Range eligibility bookkeeping, divergence checking, the queue hook's early-return paths
 # and the kernarg block retention policy are pure host logic with no GPU or HSA runtime
 # dependency, so these are plain CXX unit tests that run unconditionally.
-foreach(_TEST range-state digest queue-hooks retained-kernarg)
+foreach(_TEST range-state digest queue-hooks retained-kernarg untracked-memory)
     string(REPLACE "-" "_" _SOURCE_STEM "${_TEST}")
     set(_SOURCE "${_SOURCE_STEM}.cpp")
 

@@ -0,0 +1,131 @@
+// MIT License
+//
+// Copyright (c) 2026 Advanced Micro Devices, Inc. All rights reserved.
+//
+// Permission is hereby granted, free of charge, to any person obtaining a copy
+// of this software and associated documentation files (the "Software"), to deal
+// in the Software without restriction, including without limitation the rights
+// to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
+// copies of the Software, and to permit persons to whom the Software is
+// furnished to do so, subject to the following conditions:
+//
+// The above copyright notice and this permission notice shall be included in
+// all copies or substantial portions of the Software.
+//
+// THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+// IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+// FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT.  IN NO EVENT SHALL THE
+// AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+// LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+// OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
+// THE SOFTWARE.
+
+#include "lib/rocprofiler-sdk/range_replay/untracked_memory.hpp"
+#include "lib/rocprofiler-sdk/range_replay/range_state.hpp"
+
+#include <gtest/gtest.h>
+
+#include <cstdint>
+#include <cstring>
+#include <vector>
+
+namespace rr = ::rocprofiler::range_replay;
+namespace um = ::rocprofiler::range_replay::untracked_memory;
+
+namespace
+{
+using type_t = um::pointer_facts_t::type_t;
+
+constexpr uint64_t pinned_addr  = 0x7f0000001000;
+constexpr uint64_t tracked_addr = 0x7f0000100000;
+
+um::pointer_facts_t
+facts_of(type_t type, bool owner_is_cpu = false, bool in_host_mapping = false)
+{
+    auto f                     = um::pointer_facts_t{};
+    f.type                     = type;
+    f.owner_is_cpu             = owner_is_cpu;
+    f.in_writable_host_mapping = in_host_mapping;
+    return f;
+}
+
+// Host-pinned at pinned_addr; everything else unknown to HSA.
+um::pointer_facts_t
+pinned_only(uint64_t address)
+{
+    return address == pinned_addr ? facts_of(type_t::hsa, true) : facts_of(type_t::unknown);
+}
+
+std::vector<uint8_t>
+kernarg_with(const std::vector<uint64_t>& words, size_t trailing = 0)
+{
+    auto out = std::vector<uint8_t>(words.size() * sizeof(uint64_t) + trailing, 0);
+    std::memcpy(out.data(), words.data(), words.size() * sizeof(uint64_t));
+    return out;
+}
+
+rr::range_context_t
+recording_context()
+{
+    auto ctx           = rr::range_context_t{};
+    ctx.record         = rr::range_record_t{7};
+    ctx.snapshot_taken = true;
+    return ctx;
+}
+
+void
+record_dispatch(rr::range_context_t& ctx, const std::vector<uint64_t>& words)
+{
+    auto d    = rr::recorded_dispatch_t{};
+    d.kernarg = kernarg_with(words);
+    ASSERT_TRUE(ctx.record.add_dispatch(std::move(d)));
+}
+}  // namespace
+
+TEST(range_replay_untracked_memory, classifies_memory_the_snapshot_does_not_restore)
+{
+    EXPECT_EQ(um::classify_pointer(facts_of(type_t::hsa, true)), um::kind_t::host_pinned);
+    EXPECT_EQ(um::classify_pointer(facts_of(type_t::hsa, false)), um::kind_t::device_other);
+    EXPECT_EQ(um::classify_pointer(facts_of(type_t::reserved)), um::kind_t::managed);
+    EXPECT_EQ(um::classify_pointer(facts_of(type_t::vmem)), um::kind_t::virtual_mem);
+    EXPECT_EQ(um::classify_pointer(facts_of(type_t::locked)), um::kind_t::locked_host);
+    EXPECT_EQ(um::classify_pointer(facts_of(type_t::unknown, false, true)),
+              um::kind_t::host_pageable);
+    EXPECT_EQ(um::classify_pointer(facts_of(type_t::unknown)), std::nullopt)
+        << "not memory a kernel can reach";
+}
+
+TEST(range_replay_untracked_memory, ignores_scalars_covered_memory_and_hidden_arguments)
+{
+    const auto covered = std::vector<um::address_span_t>{{tracked_addr, tracked_addr + 4096}};
+    const auto facts   = [](uint64_t) { return facts_of(type_t::hsa, true); };
+
+    // A small scalar, then a pointer into the snapshot, then the hidden block with a stale pointer.
+    auto kernarg = kernarg_with({42, tracked_addr + 8}, um::implicit_arg_block_bytes);
+    std::memcpy(kernarg.data() + 2 * sizeof(uint64_t), &pinned_addr, sizeof(pinned_addr));
+
+    auto seen = std::unordered_set<uint64_t>{};
+    EXPECT_FALSE(um::find_untracked(kernarg.data(), kernarg.size(), covered, facts, seen));
+}
+
+TEST(range_replay_untracked_memory, a_dispatch_handed_untracked_memory_declines_the_range)
+{
+    auto ctx = recording_context();
+    record_dispatch(ctx, {42});
+    rr::decline_untracked(ctx, pinned_only);
+    EXPECT_TRUE(ctx.record.eligible());
+
+    record_dispatch(ctx, {7, pinned_addr});
+    rr::decline_untracked(ctx, pinned_only);
+    EXPECT_FALSE(ctx.record.eligible());
+    EXPECT_EQ(ctx.record.status(), ROCPROFILER_RANGE_REPLAY_STATUS_SNAPSHOT_FAILED);
+}
+
+TEST(range_replay_untracked_memory, nothing_is_checked_before_the_entry_snapshot)
+{
+    auto ctx           = recording_context();
+    ctx.snapshot_taken = false;
+    record_dispatch(ctx, {pinned_addr});
+    rr::decline_untracked(ctx, pinned_only);
+    EXPECT_TRUE(ctx.record.eligible());
+}