minseobshin11 · Aug 28, 2026 · Oct 2, 2026 · Oct 2, 2026 · Oct 7, 2026 · Oct 27, 2026
@@ -91,8 +91,8 @@ experimental/range_replay.h           public payload struct, status enum, begin/
         +-- hsa/queue.cpp             recording hook, foreign-dispatch detection
         +-- hsa/async_copy.cpp        device-write detection
         |
-        +-- kernel_replay/            reused unchanged: memory_tracker, memory_snapshot,
-                                      local_context
+        +-- kernel_replay/            shared: memory_tracker, GPU/pinned snapshot planner,
+                                      batched blit restore, local_context
 ```
 
 `range_state.cpp` and `digest.cpp` are deliberately free of GPU dependencies. The eligibility

@@ -150,9 +150,10 @@ restored after the last pass so the application resumes with its own results —
 final replayed pass left behind. Skipping it would silently hand the application the replay's
 output.
 
-A failed `restore()` is fatal rather than a decline. A partial host-to-device copy leaves device
-memory in a state that is neither the application's nor the snapshot's, and continuing would hand
-that corruption to the application. This matches kernel replay's convention for the same situation.
+A failed `restore()` is fatal rather than a decline. A partial host-to-device copy or failed device
+blit leaves device memory in a state that is neither the application's nor the snapshot's, and
+continuing would hand that corruption to the application. This matches kernel replay's convention
+for the same situation.
 
 Each dispatch in a pass has its AQL barrier bit forced on, so the passes are serialized
 dispatch-by-dispatch. This is stricter than the application, whose packets may have been free to
@@ -167,8 +168,14 @@ virtual-memory-mapped allocations, or a value the host recomputed inside the ran
 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
-each snapshot region (FNV-1a) after the application's execution and again after the final replayed
-pass, and reports how many regions differ as the `CLOSE` callback's `divergence_count`.
+each snapshot region after the application's execution and again after the final replayed pass, and
+reports how many regions differ as the `CLOSE` callback's `divergence_count`. Hashing uses the
+SSE4.2 CRC32 instruction when the host provides it and falls back to FNV-1a otherwise; the digests
+are compared only against digests produced by the same binary in the same process and are never
+persisted, so the hash function may change between releases. The check costs one extra
+device-to-host copy of the snapshot per digest plus the hash itself; with the CRC32 path the copy
+dominates, and on hosts without SSE4.2 the fallback hash is slower than the copy and dominates
+instead.
 
 Regions are ordered by device address before hashing, because the allocation inventory itself is
 unordered and the two snapshots must compare positionally. A region-count mismatch is reported as

@@ -40,6 +40,7 @@
 #include "lib/rocprofiler-sdk/hsa/signal_pool.hpp"
 #include "lib/rocprofiler-sdk/kernel_dispatch/profiling_time.hpp"
 #include "lib/rocprofiler-sdk/kernel_dispatch/tracing.hpp"
+#include "lib/rocprofiler-sdk/kernel_replay/blit-copy.hpp"
 #include "lib/rocprofiler-sdk/kernel_replay/local_context.hpp"
 #include "lib/rocprofiler-sdk/kernel_replay/memory_snapshot.hpp"
 #include "lib/rocprofiler-sdk/kernel_replay/replay_callbacks.hpp"
@@ -1003,17 +1004,40 @@ WriteInterceptor(const void* packets,
             // amd-vkale).
             replay_drain_agent_or_fatal(replay_agent);
 
+            const auto gpu_backend_available =
+                kernel_replay::blit::prepare(queue) == HSA_STATUS_SUCCESS;
+            if(!gpu_backend_available)
+                ROCP_WARNING << "kernel replay: internal blit kernel unavailable; trying "
+                                "pinned-host snapshot backing";
+
+            // snap() holds the inventory read lock while this runs, so the blit must complete
+            // before it returns.
+            const auto capture_copy = [&](const auto& regions) {
+                auto pending = std::vector<kernel_replay::blit::packet_info>{};
+                auto status  = kernel_replay::blit::copy(queue, writer, regions, pending);
+                if(status != HSA_STATUS_SUCCESS || pending.empty()) return status;
+                return pending.back().wait();
+            };
+
             // Save this agent's tracked device allocations so every pass runs against identical
-            // inputs. snap() returns ok=false if it could not capture the complete set (host memory
-            // pressure, a failed copy, or module-scope variables it could not enumerate). It logs
-            // which of those it hit.
-            const auto snapshot = kernel_replay::memory_snapshot::snap(replay_agent);
+            // inputs. Plan whole-region placement against GPU then pinned-host budgets and
+            // materialize at most one arena per tier before capture starts.
+            const auto snapshot = kernel_replay::memory_snapshot::snap(
+                kernel_replay::memory_snapshot::capture_context_t{
+                    replay_agent,
+                    queue.get_agent().near_cpu(),
+                    queue.get_agent().gpu_pool(),
+                    queue.get_agent().cpu_pool(),
+                    kernel_replay::memory_snapshot::capture_mode::automatic,
+                    gpu_backend_available,
+                },
+                capture_copy);
 
             // Snapshot incomplete: restoring a partial snapshot between passes would corrupt
             // application data, so decline replay. Close the CONFIG sequence, free our drain
             // signal, and run this dispatch once still under the writer lock with its original
             // completion signal, then return.
-            if(!snapshot.ok)
+            if(!snapshot.complete())
             {
                 LOG_FIRST_N(WARNING, 1) << "kernel replay: snapshot capture incomplete; running "
                                            "this dispatch once without replay";
@@ -1040,17 +1064,9 @@ WriteInterceptor(const void* packets,
             auto local_ctx_tls_guard =
                 kernel_replay::scoped_local_context_control{context::get_active_contexts()};
 
-            // Per-pass loop: PASS enter -> submit -> drain the async handler -> PASS exit -> ask
-            // the tool whether to continue -> restore device memory before the next pass.
-            for(uint64_t pass = 0;; ++pass)
-            {
-                const bool is_final =
-                    !replay_plan.indefinite && (pass == replay_plan.total_passes - 1);
-
-                auto pass_state = kernel_replay::pass_context_state_t{};
-                kernel_replay::execute_pass_phase_enter(
-                    replay_plan, pass, thr_id, internal_corr_id, ancestor_corr_id, pass_state);
+            auto restore_packets = std::vector<kernel_replay::blit::packet_info>{};
 
+            auto submit_and_drain_pass = [&]() {
                 process_packet_batch(packets_arr,
                                      1,
                                      forward_to_writer,
@@ -1059,27 +1075,63 @@ WriteInterceptor(const void* packets,
 
                 // Drain this pass's async handler (separate HSA thread: reads counters, emits
                 // records, releases signals/corr-id refs) before PASS EXIT / continue-decision /
-                // restore() / next submit, else we race its record delivery and reuse buffers and
-                // signals it still holds. This also implies GPU drain. Exactly one handler is in
-                // flight per pass (we drain before each submit, under the agent writer lock).
-                ROCP_FATAL_IF(queue.active_async_packets() > 1) << fmt::format(
-                    "kernel replay: more than one async handler in flight during a replay pass");
+                // next submit, else we race its record delivery and reuse buffers and signals it
+                // still holds. This also implies GPU drain. Exactly one handler is in flight per
+                // pass (we drain before each submit, under the agent writer lock).
+                ROCP_FATAL_IF(queue.active_async_packets() > 1)
+                    << fmt::format("kernel replay: more than one async handler in flight during "
+                                   "a replay pass");
                 replay_drain_or_fatal(queue);
 
+                // The target kernel was queued behind the restore dependency barrier. Its
+                // completion proves the blit also completed, so the packet can be retired without
+                // another wait.
+                for(auto& restore : restore_packets)
+                    ROCP_FATAL_IF(restore.retire() != HSA_STATUS_SUCCESS)
+                        << "kernel replay: asynchronous restore blit failed";
+                restore_packets.clear();
+            };
+
+            // Per-pass loop: PASS enter -> optional restore -> submit -> drain the async handler ->
+            // PASS exit -> ask the tool whether to continue.
+            for(uint64_t pass = 0;; ++pass)
+            {
+                const bool is_final =
+                    !replay_plan.indefinite && (pass == replay_plan.total_passes - 1);
+
+                auto pass_state = kernel_replay::pass_context_state_t{};
+                kernel_replay::execute_pass_phase_enter(
+                    replay_plan, pass, thr_id, internal_corr_id, ancestor_corr_id, pass_state);
+
+                if(pass == 0)
+                {
+                    submit_and_drain_pass();
+                }
+                else
+                {
+                    // memory_snapshot::restore invokes this callback while holding the inventory
+                    // read lock. Submit the target directly behind the blit and drain it before
+                    // returning, so the lock remains held until both packets complete.
+                    const auto batch_copy = [&](const auto& regions) {
+                        auto status =
+                            kernel_replay::blit::copy(queue, writer, regions, restore_packets);
+                        if(status != HSA_STATUS_SUCCESS) return status;
+
+                        submit_and_drain_pass();
+                        return HSA_STATUS_SUCCESS;
+                    };
+                    ROCP_FATAL_IF(!kernel_replay::memory_snapshot::restore(snapshot, batch_copy))
+                        << fmt::format("kernel replay: restore failed between passes (partial "
+                                       "host->device copy). Aborting rather than continuing with "
+                                       "corrupted device memory");
+                }
+
                 kernel_replay::execute_pass_phase_exit(replay_plan, pass, pass_state);
 
                 // Stop once the tool (or the fixed pass count) says we're done; the last executed
                 // pass leaves device memory as the app expects, so no restore follows the break.
                 if(!kernel_replay::should_continue_replay(replay_plan, pass_state, pass, is_final))
                     break;
-
-                // Restore device memory between passes so the next pass sees identical inputs.
-                // A failed host->device copy leaves the snapshot only partially applied; continuing
-                // would submit the next pass over corrupted memory and (because the final pass
-                // skips restore) would also leave that corruption visible to the application.
-                ROCP_FATAL_IF(!kernel_replay::memory_snapshot::restore(snapshot)) << fmt::format(
-                    "kernel replay: restore failed between passes (partial host->device copy); "
-                    "aborting rather than continuing with corrupted device memory");
             }
 
             kernel_replay::execute_config_phase_exit(

@@ -140,36 +140,6 @@ replay_ring_submit(const Queue& queue, const rocprofiler_packet* packets, size_t
     return true;
 }
 
-bool
-replay_barrier_and_wait(const Queue& queue, hsa_signal_t completion)
-{
-    using namespace std::chrono_literals;
-
-    auto barrier   = hsa_barrier_and_packet_t{};
-    barrier.header = HSA_PACKET_TYPE_BARRIER_AND << HSA_PACKET_HEADER_TYPE;
-    barrier.header |= 1 << HSA_PACKET_HEADER_BARRIER;
-    barrier.completion_signal = completion;
-
-    auto packet = rocprofiler_packet{barrier};
-
-    const auto& core = queue.core_api();
-    core.hsa_signal_store_screlease_fn(completion, 1);
-
-    if(!replay_ring_submit(queue, &packet, 1)) return false;
-
-    replay_wait_or_fatal(
-        [&]() {
-            return core.hsa_signal_wait_scacquire_fn(completion,
-                                                     HSA_SIGNAL_CONDITION_EQ,
-                                                     0,
-                                                     std::chrono::nanoseconds{5s}.count(),
-                                                     HSA_WAIT_STATE_BLOCKED) == 0;
-        },
-        "this queue's prior GPU work");
-
-    return true;
-}
-
 void
 set_interceptor_passthrough(bool enabled)
 {

@@ -125,11 +125,6 @@ replay_drain_agent_or_fatal(hsa_agent_t agent);
 bool
 replay_ring_submit(const Queue& queue, const rocprofiler_packet* packets, size_t count);
 
-// Submit a barrier packet carrying `completion` and wait for it, fencing the CPU against all GPU
-// work previously submitted to `queue`. Returns false if the submission itself failed.
-bool
-replay_barrier_and_wait(const Queue& queue, hsa_signal_t completion);
-
 // While set on this thread, WriteInterceptor forwards packets untouched. Range replay's re-submit
 // path transforms the recorded packets once (via Queue::invoke_write_interceptor) and then writes
 // the result to the ring, which re-enters interception on this same thread; without this flag the

@@ -1,15 +1,63 @@
 #
 set(ROCPROFILER_LIB_KERNEL_REPLAY_HEADERS
-    kernel_replay.hpp local_context.hpp memory_snapshot.hpp memory_tracker.hpp
-    replay_callbacks.hpp utils.hpp)
+    blit-copy-kernel.hpp
+    blit-copy.hpp
+    kernel_replay.hpp
+    local_context.hpp
+    memory_snapshot.hpp
+    memory_tracker.hpp
+    replay_callbacks.hpp
+    snapshot-plan.hpp
+    snapshot-storage.hpp
+    utils.hpp)
 set(ROCPROFILER_LIB_KERNEL_REPLAY_SOURCES
-    kernel_replay.cpp local_context.cpp memory_snapshot.cpp memory_tracker.cpp
-    replay_callbacks.cpp utils.cpp)
+    blit-copy.cpp kernel_replay.cpp local_context.cpp memory_snapshot.cpp
+    memory_tracker.cpp replay_callbacks.cpp utils.cpp)
 
 target_sources(
     rocprofiler-sdk-object-library PRIVATE ${ROCPROFILER_LIB_KERNEL_REPLAY_SOURCES}
                                            ${ROCPROFILER_LIB_KERNEL_REPLAY_HEADERS})
 
+find_program(
+    kernel_replay_hipcc
+    NAMES hipcc
+    HINTS ${ROCM_PATH} ENV ROCM_PATH /opt/rocm
+    PATHS ${ROCM_PATH} ENV ROCM_PATH /opt/rocm
+    PATH_SUFFIXES bin llvm/bin REQUIRED)
+
+set(kernel_replay_blit_targets ${GPU_TARGETS})
+list(REMOVE_DUPLICATES kernel_replay_blit_targets)
+set(kernel_replay_blit_code_objects)
+foreach(target_id IN LISTS kernel_replay_blit_targets)
+    set(output_file ${CMAKE_CURRENT_BINARY_DIR}/${target_id}_kernel_replay_blit.hsaco)
+    add_custom_command(
+        OUTPUT ${output_file}
+        COMMAND
+            ${kernel_replay_hipcc} --genco --no-gpu-bundle-output
+            --offload-arch=${target_id} -mcode-object-version=5 -O3 -Xoffload-linker
+            --build-id=sha1 -o ${output_file} ${CMAKE_CURRENT_SOURCE_DIR}/blit-copy.hip
+        DEPENDS ${CMAKE_CURRENT_SOURCE_DIR}/blit-copy.hip
+                ${CMAKE_CURRENT_SOURCE_DIR}/blit-copy-kernel.hpp
+        COMMENT "Building ${target_id} kernel replay blit code object")
+    list(APPEND kernel_replay_blit_code_objects ${output_file})
+endforeach()
+
+add_custom_target(kernel-replay-blit-code-objects
+                  DEPENDS ${kernel_replay_blit_code_objects})
+add_dependencies(rocprofiler-sdk-object-library kernel-replay-blit-code-objects)
+
+install(
+    FILES ${kernel_replay_blit_code_objects}
+    DESTINATION ${CMAKE_INSTALL_DATADIR}/rocprofiler-sdk/kernel-replay
+    COMPONENT core)
+
+target_compile_definitions(
+    rocprofiler-sdk-object-library
+    PRIVATE
+        ROCPROFILER_KERNEL_REPLAY_BLIT_BUILD_DIR="${CMAKE_CURRENT_BINARY_DIR}"
+        ROCPROFILER_KERNEL_REPLAY_BLIT_INSTALL_DIR="${CMAKE_INSTALL_FULL_DATADIR}/rocprofiler-sdk/kernel-replay"
+    )
+
 if(ROCPROFILER_BUILD_TESTS)
     add_subdirectory(tests)
 endif()
@@ -0,0 +1,70 @@
+// 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.
+
+#pragma once
+
+#include <cstdint>
+
+namespace rocprofiler
+{
+namespace kernel_replay
+{
+namespace blit
+{
+namespace kernel_abi
+{
+inline constexpr auto bytes_per_item = std::uint64_t{16};
+inline constexpr auto workgroup_size = std::uint16_t{1024};
+
+constexpr std::uint64_t
+items_for_size(std::uint64_t size)
+{
+    return size / bytes_per_item + static_cast<std::uint64_t>(size % bytes_per_item != 0);
+}
+
+constexpr std::uint64_t
+first_local_item(std::uint64_t thread, std::uint64_t stride, std::uint64_t logical_base)
+{
+    const auto base_mod = logical_base % stride;
+    return (thread >= base_mod) ? (thread - base_mod) : (stride - (base_mod - thread));
+}
+
+struct copy_descriptor_t
+{
+    std::uint64_t size           = 0;
+    std::uint64_t source_address = 0;
+    std::uint64_t dest_address   = 0;
+};
+
+struct kernel_args_t
+{
+    std::uint64_t descriptors_address = 0;
+    std::uint64_t descriptor_count    = 0;
+    std::uint64_t launched_items      = 0;
+};
+
+static_assert(sizeof(copy_descriptor_t) == 24);
+static_assert(sizeof(kernel_args_t) == 24);
+}  // namespace kernel_abi
+}  // namespace blit
+}  // namespace kernel_replay
+}  // namespace rocprofiler