Skip to content

[Core] Add whole-hook BatchQueue/SliceQueue handoff, pageable staging pool, and capture leases - #173

Merged
zhaoyboo merged 1 commit into
mainfrom
feat/streaming-batch-handoff
Nov 5, 2026
Merged

zhaoyboo merged 1 commit into
mainfrom
feat/streaming-batch-handoff

Conversation

@zhaoyboo

@zhaoyboo zhaoyboo commented Nov 4, 2026

Copy link
Copy Markdown
Collaborator

Depends on #172.

Summary

In-process consumers without the GIL on the drain path: a SliceQueue
(per-request slices) or BatchQueue (one whole tensor per HookPoint per step
plus a segments table), pooled pageable buffers, capture leases that track
when a step's data has been consumed, and a delivery barrier.

Implementation details

  • capture_lease.h, drop_notice.h, slice_queue.h, batch_queue.h
    (rank-ordered eviction that evicts nothing unless evicting every less
    urgent item admits the newcomer), buffer_pool.h (reference-counted pool
    state, so delivered tensors may outlive the engine; bounded size classes),
    batch submit in the P2P thread with bounds-checked segment rows, capture
    timestamps, request_flush, drain_to_sink_and_wait (returns once the
    consumer thread has exited), consumer-failure accounting
    (submit_exceptions, last_sink_error, opt-in abort).
  • MonitoringEngine(ring_consumer=...) (a BatchQueue selects the batch
    handoff); RingWithoutConsumerWarning.
  • docs/ring-consumers.md.

Compatibility and scope

Existing consumers are unchanged.

Validation

  • test_ring_consumer_queues.py (11 passed, including eviction that cannot
    admit the newcomer), test_ring_consumers_gpu.py (9 passed,
    including tensors outliving the engine).
  • Native: test_ring_engine 317 assertions, including a
    delivery barrier that returns after the consumer thread exits (it hangs
    without the fix).
  • CPU tests: 2323 passed, 458 skipped.

@Samfisheryu

Copy link
Copy Markdown
Collaborator

Independent audit of PR #173 at 7f984280231ef21c51be4d58c8562b353a649e98. Four findings below cover correctness, bounded retention, lifecycle ownership, and the delivery barrier. GPU validation and performance benchmarks have not run; they remain deferred until the repository owner explicitly authorizes them. The evidence below consists of executed CPU reproducers and a separately labeled source-level timeout finding.

1. [P1] BatchQueue silently drops supported HF batched token and attention captures

Trigger: select BatchQueue / batch_submit=True and produce HF batched token data [batch, tokens, hidden] or HF batched attention [batch, heads, q, k].

Expected: the documented per-slice fallback delivers those captures, or construction/configuration rejects an unsupported shape before accepting captures. Loss must not be invisible.

Actual: the binding installs only batch_sink, leaving submit_fn empty. The whole-batch branch requires flattened && !is_attn; HF inputs fail that condition and then return at the empty submit_fn check. Neither sink receives data, and submit_exceptions remains zero. This contradicts the documented HF per-slice fallback and makes capture completeness unreliable.

Executed CPU evidence: the actual P2P postprocessing code with the same callback arrangement as the binding delivered each HF case once through SubmitFn, but delivered neither case through BatchSink:

legacy_submit_fn hf_batched_token delivered_slices=1 delivered_batches=0 submit_exceptions=0
legacy_submit_fn hf_batched_attention delivered_slices=1 delivered_batches=0 submit_exceptions=0
batch_sink hf_batched_token delivered_slices=0 delivered_batches=0 submit_exceptions=0
batch_sink hf_batched_attention delivered_slices=0 delivered_batches=0 submit_exceptions=0

Packed token is a positive control. Packed attention is explicitly excluded from the regression claim: both the legacy and new modes already skip it. The complete control output and reproducible harness are embedded below finding 2.

Fix direction: provide an actual handoff for HF fallback slices, for example a supported one-request batch representation, or reject that consumer/layout combination before capture starts. Add token and attention tests that assert delivery, metadata, and loss/error accounting.

2. [P1] Queue retention accounting ignores the pooled backing allocation

Trigger: retain many small pooled tensors, such as sparse/head captures, in a BatchQueue with a byte cap.

Expected: the advertised retained-byte bound accounts for the backing allocations kept alive by queued items, with a documented policy for metadata overhead and shared ownership.

Actual: make_batch_item stores tensor.nbytes(), but BufferPool::acquire allocates a rounded class whose minimum is 64 KiB and exposes only the requested logical byte length through from_blob. Live queued blocks do not count toward the pool's separate idle-byte cap. Small payloads therefore retain at least 2,048 times their counted payload bytes as backing capacity in this example, before metadata overhead.

Executed CPU evidence: 32 independent 32-byte pooled items, whose logical sum exactly equals max_bytes=1024:

pool_cap=1024 retained_bytes=1024 qsize=32 actual_allocator_bytes=2098432 pool_idle=0 drops=0

actual_allocator_bytes is the sum of glibc malloc_usable_size results: allocated backing capacity, not a measured resident-RSS value. Exact allocator overhead can vary across systems; the 32 separately retained 64-KiB size classes are fixed by the implementation. This is not the documented single-oversized-item exception: every item is individually small, and there are 32 of them.

Fix direction: preserve the allocation's actual capacity alongside the tensor and charge that retained capacity when queueing; bound/charge material metadata and segments as appropriate. Replacing tensor.nbytes() with tensor.storage().nbytes() alone does not fix a from_blob storage created with the logical length. Add a small-payload/many-item cap test in addition to oversized-item tests.

Reproduce findings 1–2 on CPU: complete C++ harness, portable build command, and observed control output

Run in a checkout of 7f984280231ef21c51be4d58c8562b353a649e98. This directly compiles the real p2p_thread.cpp and uses the real queue/pool headers. It stubs only the unused drain construction and never starts a worker or calls a CUDA API; an accidental cudaFreeHost call aborts. CUDA headers are required for type declarations. The installed ATen library may transitively depend on CUDA runtime libraries, but syscall traces of the executed original harnesses showed no /dev/nvidia, /dev/dri, or libcuda.so access.

The two blocks below were executed as separate CPU harnesses against this PR. They are combined here without changing their test logic, to make the report self-contained. The combined wrapper itself has not been rerun. The legacy-vs-batch control was executed against #173; the pool case and original HF-token loss case were also independently executed against #173 and the cumulative #174 head.

Save the following as /tmp/pr173_cpu_audit.cpp:

// CPU-only: actual P2P postprocessing, legacy SubmitFn vs new BatchSink.
// No CUDA API call. Drain construction has no running worker or GPU allocation.
#include <ATen/ATen.h>
#include <iostream>
#include <malloc.h>
#include "ring/drain_thread.h"
#include "ring/record_consumer.h"
#include "ring/tensor_meta.h"
#include "ring/buffer_pool.h"
#include "ring/batch_queue.h"
#define private public
#include "ring/p2p_thread.h"
#undef private
extern "C" cudaError_t cudaFreeHost(void*) { std::abort(); }
namespace ring {
DrainThread::DrainThread(RingState& rs, PinnedStaging& st, const RingConfig& cfg)
    : ring_(rs), staging_(st), cfg_(cfg) {}
DrainThread::~DrainThread() noexcept = default;
}
int main() {
  for (bool batch_mode : {false, true}) {
    ring::RingState state{};
    ring::PinnedStaging staging;
    ring::RingConfig cfg;
    cfg.batch_submit = batch_mode;
    ring::DrainThread drain(state, staging, cfg);
    ring_py::TensorMetaFifo fifo;
    size_t slices = 0, batches = 0;
    ring::SubmitFn slice;
    ring::BatchSink batch;
    if (batch_mode) {
      batch = [&](const std::string&, int32_t, const std::string&, int32_t,
                  uint64_t, bool, at::Tensor, std::vector<std::string>,
                  at::Tensor, ring_py::CaptureLease, int64_t) { ++batches; };
    } else {
      slice = [&](const std::string&, int32_t, const std::string&, const std::string&,
                  int32_t, int32_t, int32_t, at::Tensor, int64_t,
                  ring_py::CaptureLease) { ++slices; };
    }
    ring::P2PThread p2p(drain, fifo, cfg, slice, batch);
    auto check = [&](const char* label, bool flattened, int hook, std::vector<int64_t> shape) {
      auto* ctx = new ring_py::StepContext;
      ctx->model_id = "m";
      ctx->flattened = flattened;
      ctx->requests = {{"r", 0, 2, 0, 0, -1}};
      ring_py::TensorMeta meta;
      meta.hook_type = hook;
      meta.shape = shape;
      meta.dtype = int(at::kFloat);
      meta.last_in_step = true;
      std::vector<ring_py::TensorMeta> metas{meta};
      fifo.push_step(ctx, metas);
      auto tensor = at::zeros(shape, at::TensorOptions().dtype(at::kFloat));
      ring::DrainTask task;
      task.tensor_total_bytes = tensor.nbytes();
      const auto previous_slices = slices, previous_batches = batches;
      p2p.do_post_processing(tensor, task);
      std::cout << (batch_mode ? "batch_sink" : "legacy_submit_fn") << " " << label
                << " delivered_slices=" << slices - previous_slices
                << " delivered_batches=" << batches - previous_batches
                << " submit_exceptions=" << p2p.submit_exceptions() << '\n';
    };
    check("packed_token", true, ring_py::HOOK_TYPE_RESID_PRE, {2, 4});
    check("hf_batched_token", false, ring_py::HOOK_TYPE_RESID_PRE, {1, 2, 4});
    check("hf_batched_attention", false, ring_py::HOOK_TYPE_PATTERN, {1, 1, 2, 2});
    check("packed_attention_existing_unsupported", true, ring_py::HOOK_TYPE_PATTERN, {1, 2, 2});
  }

    ring::BufferPool pool(256 << 20);
    ring_py::BatchQueue queue(1024);
    size_t allocator_bytes = 0;
    for (size_t i = 0; i < 32; ++i) {
        ring_py::HookBatchItem item;
        item.tensor = pool.acquire(32);
        item.nbytes = item.tensor.nbytes();
        allocator_bytes += malloc_usable_size(item.tensor.data_ptr());
        item.act_name = "hook_resid_pre_head";
        item.req_ids = {"r"};
        queue.push(std::move(item));
    }
    std::cout << "pool_cap=1024 retained_bytes=" << queue.retained_bytes()
              << " qsize=" << queue.size() << " actual_allocator_bytes=" << allocator_bytes
              << " pool_idle=" << pool.idle_bytes() << " drops=" << queue.dropped() << '\n';
}

Use the Python environment with the installed Torch version used to build DMI. Set CUDA_HOME if your CUDA headers are not in /usr/local/cuda. This obtains Torch paths and ABI from that environment instead of relying on paths from the audit machine:

export CUDA_VISIBLE_DEVICES=""
export OMP_NUM_THREADS=1
export MKL_NUM_THREADS=1
AUDIT_CUDA_HOME="${CUDA_HOME:-/usr/local/cuda}"
mapfile -t AUDIT_TORCH_INFO < <(python - <<'PYTORCH'
from pathlib import Path
import torch
root = Path(torch.__file__).resolve().parent
print(root / "include")
print(root / "include" / "torch" / "csrc" / "api" / "include")
print(root / "lib")
print(int(torch.compiled_with_cxx11_abi()))
PYTORCH
)
nice -n 15 g++ -std=c++17 -O0 -ffunction-sections -fdata-sections \
  "-D_GLIBCXX_USE_CXX11_ABI=${AUDIT_TORCH_INFO[3]}" \
  -I native/csrc -I "$AUDIT_CUDA_HOME/include" \
  -I "${AUDIT_TORCH_INFO[0]}" -I "${AUDIT_TORCH_INFO[1]}" \
  /tmp/pr173_cpu_audit.cpp native/csrc/ring/p2p_thread.cpp \
  -Wl,--gc-sections -L "${AUDIT_TORCH_INFO[2]}" \
  "-Wl,-rpath,${AUDIT_TORCH_INFO[2]}" \
  -ltorch_cpu -lc10 -pthread -o /tmp/pr173_cpu_audit
/tmp/pr173_cpu_audit

The audit's actual successful builds used -std=c++17 and Torch's reported ABI value 1. That Torch installation emitted warnings about C++20 bit-field initializer extensions. -std=c++20 is recommended for those headers, but C++20 was not the audit-run compiler mode.

Observed output of the extended delivery-control harness:

legacy_submit_fn packed_token delivered_slices=1 delivered_batches=0 submit_exceptions=0
legacy_submit_fn hf_batched_token delivered_slices=1 delivered_batches=0 submit_exceptions=0
legacy_submit_fn hf_batched_attention delivered_slices=1 delivered_batches=0 submit_exceptions=0
legacy_submit_fn packed_attention_existing_unsupported delivered_slices=0 delivered_batches=0 submit_exceptions=0
batch_sink packed_token delivered_slices=0 delivered_batches=1 submit_exceptions=0
batch_sink hf_batched_token delivered_slices=0 delivered_batches=0 submit_exceptions=0
batch_sink hf_batched_attention delivered_slices=0 delivered_batches=0 submit_exceptions=0
batch_sink packed_attention_existing_unsupported delivered_slices=0 delivered_batches=0 submit_exceptions=0

Observed output of the separately executed pool block:

pool_cap=1024 retained_bytes=1024 qsize=32 actual_allocator_bytes=2098432 pool_idle=0 drops=0

3. [P2] Pending metadata keeps the new capture lease non-idle after FIFO destruction

Attribution: the raw-pointer ownership leak in TensorMetaFifo predates this PR. This PR does not introduce that original leak. It attaches a new shared CaptureLease to those contexts and promises lifecycle completion, so the existing ownership gap now prevents the new lease contract from completing.

Trigger: metadata is published, then the forward fails before a producer publishes its payload; shutdown destroys the FIFO with an unpopped context still queued.

Expected: once workers stop and pending work is cancelled/dropped, destroying its owning FIFO releases pending contexts and their lease references. An integration must be able to retire a configuration after teardown.

Actual: the new StepContext::capture_lease retains the configuration, while TensorMetaFifo accepts ownership but stores raw pointers without a reclaiming destructor. Destroying the deque does not delete its pointees. Consequently CaptureLease::idle() remains false even after the FIFO is gone.

Executed CPU evidence:

before_fifo idle=1
queued idle=0
fifo_destroyed idle=0

Fix direction: express pending ownership with unique_ptr, or explicitly delete all pending contexts during cancellation/destruction after workers stop. Cover a forward failure between metadata publication and producer execution, followed by teardown/recreation and configuration retirement.

Reproduce finding 3: standard-library-only C++ header test, no ATen/CUDA

In the same pinned checkout, save this as /tmp/pr173_lease_fifo.cpp:

#include <iostream>
#include "ring/tensor_meta.h"
int main() {
  ring_py::CaptureLease lease(17);
  std::cout << "before_fifo idle=" << lease.idle() << '\n';
  {
    ring_py::TensorMetaFifo fifo;
    auto* ctx = new ring_py::StepContext;
    ctx->capture_lease = lease;
    std::vector<ring_py::TensorMeta> metas(1);
    metas[0].last_in_step = true;
    fifo.push_step(ctx, metas);
    std::cout << "queued idle=" << lease.idle() << '\n';
  }
  std::cout << "fifo_destroyed idle=" << lease.idle() << '\n';
}

Compile and run:

g++ -std=c++17 -I native/csrc /tmp/pr173_lease_fifo.cpp -o /tmp/pr173_lease_fifo
/tmp/pr173_lease_fifo

No Torch or CUDA includes/libraries are needed. The observed output is shown above; the expected final value is fifo_destroyed idle=1. This directly demonstrates the ownership problem. The full CUDA forward-failure/engine teardown scenario remains deferred.

4. [P2, source-confirmed; end-to-end runtime deferred] Delivery-barrier timeout starts after unbounded waits

Trigger: drain_to_sink_and_wait(timeout_ms) is called while the producer stream still has work or while staging is saturated behind a slow/stuck consumer.

Expected: a finite timeout bounds the full documented delivery-barrier operation and returns false when that deadline expires.

Actual: the Python-facing wrapper synchronizes the CUDA stream before forwarding timeout_ms. The engine then calls unbounded force_flush_and_wait(); its condition-variable wait has no deadline. Only the final P2P wait receives the timeout. Thus a nominal 1-ms call can spend arbitrarily longer than that before the timeout begins, contrary to the documented timeout behavior.

Evidence status: this is confirmed from the implementation's wait ordering. No CUDA or staging-saturation runtime reproducer has been run; no observed wall-time measurement is claimed.

Fix direction: create one steady-clock deadline on entry and use the remaining budget throughout producer completion, flush, and P2P completion. The drain already has force_flush_and_wait_until. Define and bound the producer-stream completion phase too. If the intention is instead a timeout only for P2P after an unbounded flush, rename/document that narrower contract explicitly so callers cannot mistake it for a bounded delivery barrier.

Inspect finding 4 without GPU execution

These commands only read the pinned source:

git show 7f984280231ef21c51be4d58c8562b353a649e98:native/csrc/ring/ring_engine_py.cu | sed -n '1318,1322p'
git show 7f984280231ef21c51be4d58c8562b353a649e98:native/csrc/ring/ring_engine.cu | sed -n '39,49p'
git show 7f984280231ef21c51be4d58c8562b353a649e98:native/csrc/ring/drain_thread.cpp | sed -n '143,184p'

The relevant chain is:

// RingEnginePy::drain_to_sink_and_wait(timeout_ms)
cudaStreamSynchronize(ms);
return impl_->engine.drain_to_sink_and_wait(timeout_ms);

// RingEngine::drain_to_sink_and_wait(timeout_ms)
drain_->force_flush_and_wait();
drain_->rethrow_drain_failure();
return p2p_->wait_until_processed(drain_->tasks_enqueued(), timeout_ms);

// DrainThread::force_flush_and_wait()
flush_done_cv_.wait(lk, [this, generation] {
    return flush_completed_generation_ >= generation;
});

The later GPU validation should saturate staging with a blocked/slow consumer and measure the full call against a finite deadline; the existing small-queue timeout test does not exercise the preceding flush under saturation. This is a pending test specification, not an executed reproduction.

Additional clean-slate/performance observations, not measured performance regressions: the new buffer pool and per-task timing/barrier bookkeeping apply to existing legacy consumers even without BatchQueue or capture leases. Setting p2p_pool_cap_bytes=0 removes idle retention but still rounds every pooled allocation to at least 64 KiB; it does not restore the previous exact-size allocation path. Default-off legacy-consumer latency/memory and saturated-queue behavior remain on the deferred GPU/benchmark checklist.

@zhaoyboo
zhaoyboo force-pushed the feat/per-hook-gather-and-inline-heads branch from 4362f6f to dc06342 Compare November 4, 2026 21:41
@zhaoyboo
zhaoyboo force-pushed the feat/streaming-batch-handoff branch from 7f98428 to c1d242b Compare November 4, 2026 21:41
@zhaoyboo

zhaoyboo commented Nov 4, 2026

Copy link
Copy Markdown
Collaborator Author

@Samfisheryu thanks for the audit and the self-contained harnesses. All four findings are fixed in c1d242b, rebased onto the revised #172. The default-off pool observation is addressed as well.

1. BatchQueue delivers HF batched captures

  • With only a batch sink installed, HF token [batch, q_len, ...] and HF attention [batch, heads, q, kv] hooks are delivered as one-request batch items, one per request slice.
  • Each item carries the segments row (0, db_start, db_end, orig_position) and is_attn set for attention. Exceptions raised by the consumer are counted in submit_exceptions.
  • Packed attention is still unsupported in both modes, as before.
  • Your extended delivery-control harness, run against this commit (the only harness change is the new trailing bool is_attn in the BatchSink signature):
    batch_sink packed_token delivered_slices=0 delivered_batches=1 submit_exceptions=0
    batch_sink hf_batched_token delivered_slices=0 delivered_batches=1 submit_exceptions=0
    batch_sink hf_batched_attention delivered_slices=0 delivered_batches=1 submit_exceptions=0
    batch_sink packed_attention_existing_unsupported delivered_slices=0 delivered_batches=0 submit_exceptions=0
    
    The legacy_submit_fn lines are unchanged.

2. Retention is charged by backing allocation

  • Pool blocks now carry their real capacity in the DataPtr context. BatchQueue and SliceQueue charge each distinct allocation once, at max(logical bytes, allocated bytes). The 64-KiB size class is therefore what counts against max_bytes, and two items sharing one block are not double-charged.
  • SliceQueue gains an optional max_bytes cap and retained_bytes.
  • Your pool block now gives:
    pool_cap=1024 retained_bytes=65536 qsize=1 actual_allocator_bytes=2098432 pool_idle=65536 drops=31
    
    The first item is admitted under the documented rule that an empty queue always admits one item. The other 31 are shed and counted.
  • On your clean-slate note: p2p_pool_cap_bytes == 0 now allocates exact-size buffers with no pool, which restores the previous allocation path.

3. TensorMetaFifo owns its pending contexts

  • The FIFO holds one owned StepContext per pending step. Clearing or destroying it releases them, and with them their capture-lease references.
  • Your pr173_lease_fifo.cpp, unchanged, now prints before_fifo idle=1, queued idle=0, fifo_destroyed idle=1.

4. One deadline bounds the whole delivery barrier

  • drain_to_sink_and_wait(timeout_ms) takes one steady-clock deadline at entry. It spends the remaining budget on producer-stream completion (polling an event instead of cudaStreamSynchronize), then on the flush (force_flush_and_wait_until), then on the P2P delivery wait. It returns false when the deadline expires in any phase.
  • docs/ring-consumers.md states this contract.

Tests

  • tests/test_native_ring_consumer_fixes.py compiles the real p2p_thread.cpp and queue/pool headers on the CPU. It covers HF token and attention fallback, counted fallback consumer errors, unchanged packed paths, pooled retention charged by allocation, exact-size allocation at cap 0, FIFO ownership of pending contexts, and a deadline-bounded delivery wait.
  • tests/test_ring_consumer_queues.py adds test_queues_charge_each_backing_allocation_once and test_slice_queue_reports_retention_and_honours_an_optional_byte_cap.

Verification for this revision, CPU only: the native extension builds, and the -m cpu suite passes (2382 passed, 459 skipped). The one failure, test_cpu_native_build.py::test_every_torch_including_compile_requests_cxx20[host], also fails on main. Your harnesses were run as shown above, without GPU access. No GPU tests or benchmarks were run. Waiting for the owner's go-ahead: the CUDA forward-failure/teardown scenario, staging saturation measured against a finite deadline, and default-off legacy-consumer latency/memory.

@Samfisheryu

Copy link
Copy Markdown
Collaborator

Independent re-review of c1d242b9e034b06a36df2f69c1df23e8a929560d: the four original findings have targeted fixes, but two new correctness regressions and one CI blocker remain.

The previous HF fallback, allocation-capacity charging, pending-context ownership, and end-to-end deadline findings are addressed at the CPU/source evidence level used in the original review. The three new CPU native regression tests pass locally. The pool=0 path now bypasses pooling and uses exact-sized allocation. Real CUDA teardown, staging-saturation timing and performance remain unverified; no GPU work was run.

[P1] New pool context is deleted twice when tensor construction throws

Location: BufferPool::acquire, lines 88–98.

TensorMaker.context(ctx, deleter) takes ownership of ctx. If make_tensor() throws, that owner destroys the context. The new catch block then calls delete_pooled_block(ctx) again on the freed object. This is a regression in the revised pool implementation, not the prior FIFO ownership issue.

With the actual CPU ATen implementation, a single injected allocation failure at the 96-byte StorageImpl allocation produced:

injected std::bad_alloc size=96
ERROR: AddressSanitizer: heap-use-after-free
READ of size 8
  ring::delete_pooled_block(void*)                 pooled_block.h:39
  ring::BufferPool::acquire(unsigned long)         buffer_pool.h:97
freed by:
  ring::BufferPool::Block::~Block()
  ring::delete_pooled_block(void*)
  at::TensorMaker::make_tensor() [clone .cold]
  ring::BufferPool::acquire(unsigned long)         buffer_pool.h:95

Without ASan this run exited with SIGSEGV. The old #173 revision, tested at the same 96-byte ATen allocation, safely propagated std::bad_alloc. Allocation order differs between implementations: this is allocation 2 in the revised implementation and allocation 3 in the old implementation; the comparison did not conflate different fault sites.

Fix: make context ownership singular through the entire construction. Do not manually delete a context already transferred to TensorMaker. Also cover the pre-transfer allocation with RAII. Add a construction-failure test that proves one release and safe exception propagation.

Actual CPU allocation-failure reproducer (Torch 2.13.0+cu130, no GPU initialization)

Save as pool_failure.cpp in a separate directory and compile from the PR checkout using the common command below. at::zeros allocates a CPU tensor; the pool uses CPU memory. Only one allocation is forced to fail, after ATen initialization. Disable core dumps before running.

#include <ATen/ATen.h>
#include <atomic>
#include <cstdio>
#include <cstdlib>
#include <new>
#include "ring/buffer_pool.h"

// Deterministically simulate one CPU allocation failure in the real ATen
// TensorMaker. No replacement/stub of TensorMaker or pool code is used.
static std::atomic<int> fail_after{-1};
void* operator new(std::size_t n) {
    const int count = fail_after.load();
    if (count > 0 && fail_after.fetch_sub(1) == 1) {
        fail_after.store(-1);
        std::fprintf(stderr, "injected std::bad_alloc size=%zu\n", n);
        throw std::bad_alloc();
    }
    if (void* p = std::malloc(n)) return p;
    throw std::bad_alloc();
}
void operator delete(void* p) noexcept { std::free(p); }
void operator delete(void* p, std::size_t) noexcept { std::free(p); }

int main(int argc, char** argv) {
    at::zeros({1}, at::kByte);  // initialize ATen before failure injection
    ring::BufferPool pool(0);  // released block is freed, never cached
    fail_after.store(argc > 1 ? std::atoi(argv[1]) : 2);
    try {
        auto t = pool.acquire(32);
        fail_after.store(-1);
        std::puts("acquire succeeded");
    } catch (const std::bad_alloc&) {
        fail_after.store(-1);
        std::puts("std::bad_alloc propagated safely");
    }
}

Use ./pool_failure 2 for this revision. Against the old 7f984280... source, ./pool_failure 3 selects the same 96-byte ATen allocation. The numbered allocation site is specific to the stated Torch build; the ownership proof is the .context transfer followed by the catch's second delete. This is a CPU AddressSanitizer result, not a GPU sanitizer run.

[P2] The allocation ledger's first-owner rank breaks eviction decisions for shared views

Locations: fixed rank stored in the ledger, rank-byte accounting, and eviction preflight.

The allocation is charged once, but its entire capacity stays assigned to the rank of the first queued view until the last view leaves. That rank can cease to describe which items keep the allocation alive. bytes_by_rank_ therefore does not tell the eviction preflight which allocations would actually be released:

  1. Enqueue an urgent 40-byte view and low-priority 40-byte view of one 80-byte allocation; drain the urgent view. Add a separate urgent 10-byte tensor. With a 100-byte cap, evicting the sole remaining low-priority view would admit an incoming urgent 80-byte tensor (final retained capacity 90). It is instead refused, because the shared allocation is still charged to the departed urgent rank.
  2. Enqueue the low-priority view first, then an urgent alias of the same allocation. An incoming urgent 80-byte allocation cannot fit even after removing low, because urgent still retains the original 80 bytes. The preflight nevertheless evicts low and then rejects the newcomer, violating the documented “evict nothing unless the newcomer can be admitted” contract.

Executed unchanged C++ queue code:

# revised #173
underestimated_evictable_capacity size=2 retained=90 dropped=1 evicted=0 items=low,urgent,
overestimated_evictable_capacity size=1 retained=80 dropped=2 evicted=1 items=urgent,

# old #173 control with the same inputs
underestimated_evictable_capacity size=2 retained=90 dropped=1 evicted=1 items=urgent,newurgent,
overestimated_evictable_capacity size=2 retained=80 dropped=1 evicted=0 items=low,urgent,

Fix: base admission/preflight on the allocation references that would remain after the proposed evictions, including surviving higher-priority aliases and the incoming item's reference. A permanent first-owner bucket cannot support the required decision. Compute the victim set without mutating the queue; commit it only if it admits the new item.

Complete CPU shared-view eviction reproducer

Save as shared_eviction.cpp; no CUDA headers or CUDA API calls are needed.

#include <ATen/ATen.h>
#include <iostream>
#include "ring/batch_queue.h"

void push(ring_py::BatchQueue& q, const char* act, const at::Tensor& t) {
    ring_py::HookBatchItem item;
    item.act_name = act;
    item.layer_no = 0;
    item.tensor = t;
    item.nbytes = t.nbytes();
    item.req_ids = {act};
    q.push(std::move(item));
}

void show(ring_py::BatchQueue& q, const char* label) {
    std::cout << label << " size=" << q.size() << " retained=" << q.retained_bytes()
              << " dropped=" << q.dropped() << " evicted=" << q.evicted() << " items=";
    for (const auto& item : q.drain(100, 0)) std::cout << item.act_name << ",";
    std::cout << " retained_after_drain=" << q.retained_bytes() << '\n';
}

int main() {
    // A formerly urgent view drains; its shared allocation is now pinned only
    // by low priority. Evicting that low item would admit the new urgent item.
    {
        ring_py::BatchQueue q(100);
        q.set_layer_class_map({{"urgent", 0, 0}, {"newurgent", 0, 0}}, 2);
        auto shared = at::zeros({80}, at::kByte);
        push(q, "urgent", shared.narrow(0, 0, 40));
        push(q, "low", shared.narrow(0, 40, 40));
        q.drain(1, 0);
        push(q, "urgent", at::zeros({10}, at::kByte));
        push(q, "newurgent", at::zeros({80}, at::kByte));
        show(q, "underestimated_evictable_capacity");
    }
    // Conversely, low and urgent items share one allocation charged to low.
    // Evicting low cannot free it. The contract requires no futile eviction.
    {
        ring_py::BatchQueue q(100);
        q.set_layer_class_map({{"urgent", 0, 0}, {"newurgent", 0, 0}}, 2);
        auto shared = at::zeros({80}, at::kByte);
        push(q, "low", shared.narrow(0, 0, 40));
        push(q, "urgent", shared.narrow(0, 40, 40));
        push(q, "newurgent", at::zeros({80}, at::kByte));
        show(q, "overestimated_evictable_capacity");
    }
}
Portable CPU compile/run commands for both reproducers

From the fixed PR checkout, with the two source files in the current directory and a Torch-enabled Python environment:

export CUDA_VISIBLE_DEVICES=''
export OMP_NUM_THREADS=1 MKL_NUM_THREADS=1 OPENBLAS_NUM_THREADS=1
ulimit -c 0
python - <<'PY'
from pathlib import Path
import subprocess
import torch

root = Path(torch.__file__).resolve().parent
includes = [Path("native/csrc"), root/"include",
            root/"include/torch/csrc/api/include"]
libs = root/"lib"
for name in ("shared_eviction", "pool_failure"):
    flags = (["-fsanitize=address", "-fno-omit-frame-pointer", "-g"]
             if name == "pool_failure" else [])
    cmd = ["g++", "-std=c++20", "-O0", "-pthread", *flags,
           "-D_GLIBCXX_USE_CXX11_ABI=" + str(int(torch.compiled_with_cxx11_abi())),
           *["-I" + str(p) for p in includes], name + ".cpp",
           "-L" + str(libs), "-Wl,-rpath," + str(libs),
           "-Wl,--disable-new-dtags", "-ltorch_cpu", "-lc10", "-o", name]
    subprocess.run(cmd, check=True)
PY
./shared_eviction
ASAN_OPTIONS=detect_leaks=0 ./pool_failure 2

The fault run is expected to fail on this revision; use the ASan diagnostic, not the process exit code alone, to identify the second deletion. The audit ran the same source and compiler/library options through an isolated runner with core dumps disabled. No production source was edited.

[P2] The new CPU test cannot run on the configured CPU CI image

Location: test_p2p_batch_handoff_and_pooled_retention.

This test is marked CPU but skips without CUDA headers. The existing CPU CI image does not install those headers, and its skip gate intentionally permits absent hardware only. This is now an observed CI failure on the cumulative head that inherits this test:

CPU CI job 126884423571

2846 passed, 2 skipped, 513 deselected
SKIPPED tests/test_native_ring_consumer_fixes.py:71:
  CUDA headers are required for the ring type declarations
1 cpu test(s) skipped for a reason that is not absent hardware; a skip is not a pass
2848 collected, 1 skipped for absent hardware, 1 unexplained
Process completed with exit code 1

Fix: make the host-only harness build without a CUDA toolkit (for example with the required declaration-only boundary), or explicitly provision its header dependency in the CPU job. Do not merely whitelist the skip and silently lose this coverage. Keep the compiler standard aligned with the project's Torch-header requirements.

These new failures mean this PR still needs changes, despite the original four findings having targeted fixes. The actual CUDA deadline/teardown and end-to-end performance checks remain pending.

@zhaoyboo
zhaoyboo force-pushed the feat/streaming-batch-handoff branch from c1d242b to 21f58b1 Compare November 5, 2026 05:59
@zhaoyboo
zhaoyboo force-pushed the feat/per-hook-gather-and-inline-heads branch from dc06342 to 22a6482 Compare November 5, 2026 05:59
@zhaoyboo

zhaoyboo commented Nov 5, 2026

Copy link
Copy Markdown
Collaborator Author

@Samfisheryu thanks for the re-review and the harnesses. All three items are fixed in 21f58b1, rebased on the revised #172.

[P1] BufferPool::acquire double delete

  • The block has exactly one owner at every step. The raw block is returned if new Block throws. ctx is a std::unique_ptr<Block> until TensorMaker::context() takes it; that call is noexcept in this Torch and only stores the pointer in its unique_ptr. After that, the TensorMaker owns it and its destructor releases it if make_tensor() throws. The manual catch and second delete are gone.
  • Your pool_failure.cpp, unchanged, built with -fsanitize=address -fno-omit-frame-pointer -g and run with ASAN_OPTIONS=detect_leaks=0:
    n=1 injected size=40  -> std::bad_alloc propagated safely, 0 ASan errors   (the Block itself)
    n=2 injected size=96  -> std::bad_alloc propagated safely, 0 ASan errors   (StorageImpl)
    n=3 injected size=176 -> std::bad_alloc propagated safely, 0 ASan errors
    n=4 injected size=24  -> std::bad_alloc propagated safely, 0 ASan errors
    n=5, n=6              -> acquire succeeded, 0 ASan errors
    
    The previous header still reproduces your heap-use-after-free at n=2.
  • New test test_buffer_pool_acquire_failure_releases_the_block_once (tests/native/test_buffer_pool_failure.cpp) fails each of acquire's allocations in turn. The pool keeps an idle cap, so idle_bytes() counts returns: each failure must propagate std::bad_alloc and return the block exactly once, and the pool must then reuse it. The previous code leaks the block at allocation 1 and crashes at allocation 2.

[P2] Shared-allocation eviction

  • The ledger no longer files an allocation's charge under a rank. make_room_locked() now picks victims without touching the queue, worst rank first and oldest first within a rank. For each candidate it counts the bytes released only when the candidate drops the allocation's last remaining queued reference. It never counts the arriving item's own allocation, which that item keeps alive. It stops at the first prefix where retained - freed + cost(arrival) <= max_bytes, where cost(arrival) is 0 when the arrival aliases a queued allocation. Only then does it evict. Otherwise nothing is evicted and the arrival is refused. The existing rule that an emptied queue always admits is kept.
  • Items keep the rank they were enqueued with, so re-installing the class map while items are queued leaves the per-rank depths consistent.
  • Your shared_eviction.cpp, unchanged:
    underestimated_evictable_capacity size=2 retained=90 dropped=1 evicted=1 items=urgent,newurgent, retained_after_drain=0
    overestimated_evictable_capacity size=2 retained=80 dropped=1 evicted=0 items=low,urgent, retained_after_drain=0
    
    (dropped includes evictions, as before.)
  • New test test_batch_queue_eviction_with_shared_allocations (tests/native/test_batch_queue_shared_eviction.cpp) runs both of your cases plus four more:
    • a shared allocation freed only when its last view, of any rank, is evicted;
    • an arrival that aliases a queued allocation, admitted at the cap with no eviction;
    • an alias arriving at a queue that is over its cap;
    • reclassification while items are queued.
  • docs/ring-consumers.md states the rule.

[P2] CPU harness without CUDA headers

  • test_p2p_batch_handoff_and_pooled_retention now always compiles against tests/native/ring/cpu_cuda_stub/, even where a toolkit is installed, so this machine builds exactly what the CUDA-free runner builds. The stub gained cudaHostAllocDefault, plus declarations of cudaHostAlloc and cudaFreeHost with the runtime's C linkage. The harness keeps its own aborting cudaFreeHost. All of these host-only builds use -std=c++20.
  • New test test_cpu_cuda_stub_stands_in_for_the_toolkit walks the harness include graph and fails if any source reaches a CUDA header the stub does not provide. A g++ -H of the build reads no toolkit header.
  • On the stacked head, pytest -m cpu --junit-xml=cpu-results.xml followed by the workflow's skip gate (run verbatim) reports 2833 collected, 1 skipped for absent hardware, 0 unexplained. The one skip is no CUDA device for the non-CPU case, and the p2p test now runs and passes. The test that leaves SIGPIPE at its default, test_native_catalog_connection.py, was run in a separate process; its 68 tests pass with no skips.

Local caveat: on this host, test_native_catalog_connection.py leaves SIGPIPE at its default and kills the process in the next file. Fifty store tests (S3, uploader, spool) also fail because of the local libcurl 7.68 and a mismatched botocore. Both are identical on origin/main and none of those files is touched by the stack, so the suite was split around that file.

On this branch, test_native_ring_consumer_fixes.py, test_ring_consumer_queues.py and the toggle and op host tests pass, and all ring test binaries link. CPU and CPU-ASan only: no GPU tests, CUDA initialization, GPU sanitizers or benchmarks were run. The CUDA deadline, teardown and performance checks stay deferred until the owner gives the go-ahead.

@Samfisheryu

Copy link
Copy Markdown
Collaborator

Independent third review of 21f58b10a3a84952f53fed1f18ae67adaf08749c: all three round-two items are resolved in the CPU/source review.

R2-173-pool — allocation cleanup: the context now has one owner through construction and transfer to TensorMaker. Against the real CPU ATen implementation, injected allocation failures 1–4 safely propagate the exception; positions 5–6 succeed. The positive test confirms one block return and subsequent reuse at every failed construction stage. The previous revision remains the contrary control at the same 96-byte allocation point. This also passed CPU allocation checking.

R2-173-eviction — shared backing references: both original examples now produce the intended results:

underestimated_evictable_capacity:
  size=2 retained=90 dropped=1 evicted=1 items=urgent,newurgent,
overestimated_evictable_capacity:
  size=2 retained=80 dropped=1 evicted=0 items=low,urgent,
retained_after_drain=0 in both cases

An independent reference model recomputing allocation sets, without using BackingLedger, passed 50,000 mixed operations: shared views, partial drain, rank-map changes, eviction enabled/disabled, zero caps and oversized first items. It compared item order, retained bytes and the admission/drop/eviction counters after each operation.

R2-173-ci — host build dependency: the harness now always uses the checked-in declaration stubs and C++20. Independent compiler dependency tracing confirmed that both P2P translation units build without toolkit headers. All 6 CPU consumer-fix tests passed, without skips. The latest CI is fully successful, including the CPU skip gate and native test link step.

The original HF handoff, lease cleanup and deadline changes remain intact. No additional confirmed issue was found in the revised ownership and admission logic. GPU teardown, staging-saturation timing and end-to-end performance remain deferred; no GPU work was performed.

@Samfisheryu Samfisheryu left a comment

Copy link
Copy Markdown
Collaborator

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Approved at 21f58b10a3a84952f53fed1f18ae67adaf08749c following the completed independent code and CPU/host review. The issues raised for this PR have been addressed and independently rechecked. The reviewed commit's CI is successful.

Review record: #173 (comment)

GPU execution and end-to-end performance benchmarks remain a separate follow-up, pending the owner's explicit instruction as agreed. This approval records the completed code/CPU review and does not claim those measurements.

@zhaoyboo
zhaoyboo force-pushed the feat/per-hook-gather-and-inline-heads branch from 22a6482 to 6d7e741 Compare November 5, 2026 23:13
… pool, and capture leases

Deliver legacy-ring captures to in-process consumers without the P2P thread
taking the GIL, account for every capture a consumer does not receive, and
give integrations an end-to-end delivery barrier.

Native:
- SliceQueue (one item per request x hook x step) and BatchQueue (one item
  per hook x step: the whole packed tensor, an int32 segments table and the
  request ids), both pushed in pure C++. BatchQueue is bounded by retained
  bytes and evicts strictly less urgent queued items first under an optional
  (act_name, layer) -> rank map; without a map it is FIFO drop-the-newcomer.
- Drop notices: every shed item names the requests that lost data; a
  saturated notice buffer raises a sticky flag for one poll window.
- Capture leases: an integration-defined identity attached to a step and
  delivered with every slice, batch and drop notice; its reference count
  tells when a configuration's payloads are all consumed. A lease may carry
  a shared drop-rank map.
- The P2P thread copies staging into pooled pageable buffers (warm pages);
  pool state is reference-counted by the delivered tensors, so a tensor that
  outlives the engine stays valid and frees its block when dropped.
- Batch handoff in the P2P thread for packed steps (RingConfig.batch_submit),
  with one step_seq per step and final_logits rows at the original position.
- drain_to_sink_and_wait(): flush, then wait until every flushed task has
  been through the consumer. Consumer exceptions are counted
  (submit_exceptions, last_sink_error); set_abort_on_sink_error makes the
  next prepare_step raise.
- capture_ts_us on every delivery: when the drain first saw the task
  published. request_flush(): flush the tasks reserved so far as soon as
  they are published, whatever the thresholds say.

Python:
- MonitoringEngine(ring_consumer=...) for a callable, SliceQueue or
  BatchQueue (switching batch_submit on); a ring built with no consumer
  warns (RingWithoutConsumerWarning) instead of silently dropping captures.
- RingTransport/StepContext carry a capture lease to the next metadata
  publication only.

Delivery, retention and lifecycle details:
- BatchQueue consumers receive HF batched captures: with only a batch sink
  installed, HF token [batch, q_len, ...] and attention [batch, heads, q, kv]
  hooks are delivered as one-request items (one per request slice, segments
  row (0, db_start, db_end, orig_position), is_attn for attention) instead
  of being dropped silently; consumer exceptions are counted.
- Queue retention charges the allocation an item keeps alive: pool blocks
  carry their capacity in the DataPtr context, and BatchQueue/SliceQueue
  charge each distinct allocation once at max(logical, allocated) bytes.
  SliceQueue gains an optional max_bytes cap and retained_bytes.
  p2p_pool_cap_bytes == 0 now allocates exact-size buffers with no pool.
- TensorMetaFifo owns its pending StepContexts (one per step): clearing or
  destroying it releases them, so their capture leases turn idle.
- drain_to_sink_and_wait(timeout_ms) takes one deadline at entry and bounds
  producer-stream completion (event polling), the flush and the delivery
  wait with it.
@zhaoyboo
zhaoyboo changed the base branch from feat/per-hook-gather-and-inline-heads to main November 5, 2026 23:13
@zhaoyboo
zhaoyboo force-pushed the feat/streaming-batch-handoff branch from 21f58b1 to d91a545 Compare November 5, 2026 23:13
@zhaoyboo
zhaoyboo merged commit 198ce0e into main Nov 5, 2026
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Labels

None yet

Projects

None yet

Development

Successfully merging this pull request may close these issues.

2 participants