CHAPTER 06

Keep the host and device busy

Chapter 1 found two oddities in DeepInfer’s loop. One call to the worker can handle work from two steps, and the outer loop must flush a pending step after the scheduler reports no unfinished requests. Chapter 5 found a third: a result can arrive for a request that the scheduler has already reset. All three come from one design choice. The host plans the next step before it has read the result of the current one.

The question

The host must choose step $N + 1$ before it has read step $N$ ’s result. What can it decide without that result? What must the device supply instead? When may storage or metadata from step $N$ be reused?

6.1 Overlap has a limit

Two parties take part. The host is the CPU process that schedules requests, prepares metadata and submits work. The device executes submitted work in order on a queue, called a stream. Submitting work and executing it are different events. The host can queue a forward pass and move on before the device starts it. Chapter 4 introduced events: markers in the queue that the host can wait for.

A decode step has two device operations here: a forward pass that evaluates the model, and sampling that chooses tokens from its output. The host has four kinds of work: scheduling, submission, post-processing of sampled tokens, and engine callbacks. The next forward needs the tokens chosen by the previous sample. The question is whether the host must read them first.

BASH

python3 labs/13_overlap.py

The experiment assigns durations to these operations and simulates both orders. In the synchronous order, each call submits forward and sampling, waits for sampling to finish, then post-processes its own result. The ZOS order follows Worker.execute_model_step (worker.py, lines 113–189). During decode call $N + 1$ , the host:

prepares inputs and queues forward $N + 1$ ;
waits for sample_end_evt of step $N$ , then post-processes step $N$ ;
queues sampling for step $N + 1$ and runs the engine callbacks in pre_sampler_sync;
records sample_end_evt for step $N + 1$ and keeps the step’s state for the next call.

The lab puts forward and sampling on one device queue, so forward $N + 1$ cannot start before sample $N$ finishes even though the host queued it earlier. The worker’s comment at line 231 states that ordering for the correction work queued before forward; this chapter’s release table returns to the cases where a single queue is assumed rather than shown [7].

OUTPUT

Assigned units. interval = repeating step; idle = device idle per step;
gap = device wait between forward and sample; wait = host wait for sample 3;
first = first sample end; finish = end of final post-processing.
case                       mode  interval idle gap wait first finish
baseline                   sync     20      8      0     10     16    120
baseline                   ZOS      12      0      0      2     16     80
slow scheduling            sync     29     17      0     10     25    174
slow scheduling            ZOS      19      7      0      0     25    120
long post-processing       sync     30     18      0     10     16    180
long post-processing       ZOS      20      8      5      0     16    127
correction before forward  ZOS      13      0      0      3     16     85

Timer view of the baseline: a host timer around 'launch' records 1 unit; the device forward it queued lasts 10.

Saved output · results/13_overlap.txt

Figure 6.1: Three baseline decode steps from labs/13_overlap.py , drawn
from data the script writes. Numbers name the step. In the synchronous
run the host waits for each sample before post-processing it, so the device
idles during that work and the next scheduling pass. Under ZOS, forward N + 1 is queued before the host waits for sample N .
Durations are assigned units, not measurements. — Figure 6.1: Three baseline decode steps from `labs/13_overlap.py`, drawn from data the script writes. Numbers name the step. In the synchronous run the host waits for each sample before post-processing it, so the device idles during that work and the next scheduling pass. Under ZOS, forward $N + 1$ is queued before the host waits for sample $N$ . Durations are assigned units, not measurements.

In the baseline, the device does 12 units of work per step: forward 10 and sampling 2. The synchronous run takes 20. Its device is idle for 8 units while the host post-processes (4), schedules the next step (3) and submits the next forward (1). The host also waits 10 units for each sample. The ZOS run takes 12, the device’s own work, and the host waits only 2. Its ten units of host work fit behind the next forward. Figure 6.1 draws the first three steps of both runs.

The other two cases show the limit. When scheduling takes 12 units, the host needs 19 units per step, more than the device’s 12. ZOS still helps, from 29 to 19 units, but the device now idles for 7 units every step. Overlap can hide host work only behind device work that is already queued. In the third case, post-processing takes 14 units. The host cannot queue sample $N + 1$ until it has post-processed step $N$ , so the device finishes forward $N + 1$ and then waits. The gap is $14 + 1 - 10 = 5$ units: post-processing plus the sampling submission, less the forward. Overlap moved the waiting; it did not remove the dependency.

The first sample ends at time 16 in both baseline runs. Overlap does not shorten the first step, because nothing is queued before it. At the end, the ZOS run must still wait for and post-process its pending step, the flush from chapter 1. The step interval describes repeating throughput; the fill and the drain are separate costs.

One more detail is visible in the code order. The worker records sample_end_evt after the callbacks, not immediately after queuing sampling (lines 161–168). The lab models that event as complete at the later of the sampling end and its own submission. Waiting for the event can therefore take longer than waiting for sampling itself would.

The loop below produces the timeline. work advances the host clock; enqueue starts device work at the later of that clock and the end of previously queued work. wait advances the host clock only if the event is still pending. The cost keys name scheduling (S), submission (L), post-processing (P), callbacks (C), forward (F), sampling (M) and optional correction (X):

Listing 6.1. Queue the next forward before collecting the previous result.

PYTHON

38    for n in range(1, steps + 1):
39        work("schedule", n, costs["S"])
40        work("launch", n, costs["L"])
41        d2h_done = enqueue("correct", n, costs["X"]) if zos and costs["X"] and n > 1 else 0
42        enqueue("forward", n, costs["F"])
43        if zos and n > 1:
44            wait(sample_end[n - 1], n - 1)  # sample_end_evt(N).synchronize()
45            work("post", n - 1, costs["P"])
46            wait(d2h_done, n)  # d2h_evt.synchronize() before sampler metadata
47        work("launch", n, costs["L"])
48        enqueue("sample", n, costs["M"])
49        work("callbacks", n, costs["C"])
50        sample_end[n] = max(free, t)  # the event is recorded after callbacks
51        if not zos:
52            wait(sample_end[n], n)
53            work("post", n, costs["P"])
54    if zos:
55        wait(sample_end[steps], steps)  # final flush of the pending step
56        work("post", steps, costs["P"])

Book lab · labs/13_overlap.py · lines 38–56

6.1.1 Which interval does a timer see?

The last line of the output repeats chapter 2’s warning in a new form. A host timer around a submission measures the submission: one unit, while the forward it queued lasts ten. CPUTimer.record reads the host clock before and after the caller’s block (timer.py, lines 50–63). The worker uses it around the wait in post-processing:

Listing 6.2. Time the wait, then read only completed graph timers.

PYTHON

262with cpu_timer.record("worker.post_process.sync"):
263    step_state.sample_end_evt.synchronize()
264with cpu_timer.record("worker.post_process.graph_timer_flush"):
265    # flush the time of completed graph, no effect on the running graph.
266    cuda_graph_timer.flush_out()

DeepInfer · deepinfer/worker/worker.py · lines 262–266

The first block measures how long the host waited for the device. In the lab’s terms, that is the wait column: 2 units per ZOS baseline step and 0 when scheduling is slow. A large wait means the host was ahead; a small wait with idle device time means the host was behind. Device timing comes from recorded events instead. CUDAGraphTimer.flush_out waits for each recorded pair before reading its elapsed time (lines 196–212), and the comment notes that only completed graphs are flushed. Neither kind of timer includes a request’s queueing or delivery; chapter 2’s boundaries still apply [46].

6.2 One request crosses two steps

Chapter 1 identified the asynchronous-scheduling precedents acknowledged by the deployment Q&A. The idea carried forward is to overlap host preparation with device execution. The question here is how DeepInfer makes that overlap work when the host’s information is a step behind. Its implementation matches rows by request identity, supplies tokens and length corrections on the device, and retains results and metadata until their readers are finished. The next sections trace those obligations. They explain this implementation’s choices without claiming that row remapping or deferred release originated here.

Overlap has a price in information. When the host schedules step $N + 1$ , step $N$ is still running. The host does not yet know which token step $N$ will sample. If step $N$ verified draft tokens, it does not know how many were accepted. The device knows both soon after sampling, and the next forward runs on the device. DeepInfer therefore lets the device complete the host’s one-step-old metadata.

BASH

python3 labs/14_lagging_metadata.py

The experiment translates the inspected matching code and correction formula into lists. It builds a batch of four requests and checks the corrected inputs against a synchronous reference that already knows step $N$ ’s result. We add one missing fact at a time.

6.2.1 The missing token

Begin without draft tokens (nextn=0). Each request evaluates one position per step, so the scheduler knows the length exactly. It lacks only the input token, which step $N$ is still sampling.

OUTPUT

A. missing token, same order: nextn=0, total_q=4, rows with padding=4, same-order path
  id remap q k_cpu base corrected true  inputs
  11     0 1    33   33        33   33  [501]
  12     1 1    34   34        34   34  [502]
  13     2 1    35   35        35   35  [503]
  14     3 1    36   36        36   36  [504]

B. reordered survivors and a new request: nextn=0, total_q=4, rows with padding=4, seq_id map
  id remap q k_cpu base corrected true  inputs
  13     2 1    35   35        35   35  [503]
  11     0 1    33   33        33   33  [501]
  15    -1 1    41   41        41   41  [515]
  12     1 1    34   34        34   34  [502]

Saved output · results/14_lagging_metadata.txt

Rows are matched by seq_id, a stable request identifier, not by position in the batch. The remap column gives the matching row of the previous batch, or $- 1$ when there is none. A matched row reads its input token from the previous sample’s output on the device. Request 15 was not in step $N$ ; the host already knows its token, 515, and supplies it. In the second batch the survivors have been reordered, as chapter 5’s length sort can do, and request 14 has left. The order changed, but every matched row still finds its token. The corrected and synchronous key lengths agree because nothing about length was missing.

6.2.2 The missing accepted length

Now let each request verify two draft tokens per step (nextn=2), so each step evaluates three positions. After verification, a request keeps its main token plus the accepted drafts. The host cannot know that number while scheduling step $N + 1$ .

The scheduler reserves the maximum. When it schedules a request, it records the step’s query length as in flight (scheduler.hpp, line 726). The next step’s key length is the committed count plus in-flight positions plus the new query length (line 697). That is an upper bound [37].

OUTPUT

C. missing accepted length, EOS and graph padding: nextn=2, total_q=12, rows with padding=8, same-order path
  id remap q k_cpu base corrected true  inputs
  11     0 3    37   34        35   35  [601, 611, 612]
  12     1 3    38   35        37   37  [2, 621, 622]
  13     2 3    39   36        39   39  [603, 631, 632]
  14     3 3    40   37        40   40  [604, 641, 642]
  evaluated although step N sampled EOS: [12]; their step N+1 output is not published

Saved output · results/14_lagging_metadata.txt

Follow request 11. It had 31 committed positions before step $N$ . Step $N$ reserved 3 and step $N + 1$ asks for 3, so the host’s key length is 37. With drafts enabled, the C++ preparation subtracts the current query length, giving a base of 34 (zos.cpp, line 188). The device then applies the reference formula. For real row $i$ with remap $r_{i}$ and query length $q_{i}$ ,

k_{i}^{'} = k_{base, i} + q_{i} - {\begin{matrix} nextn - a_{r_{i}}, & r_{i} \geq 0, \end{matrix}

where $a_{r_{i}}$ counts the drafts accepted in the matched previous row, excluding the main token. Request 11 accepted none: $34 + 3 - (2 - 0) = 35$ . The synchronous reference agrees: $31 + 1 + 0 = 32$ committed positions after step $N$ , plus three new queries. Request 13 accepted both drafts, so its reserved length was already right: $36 + 3 - 0 = 39$ . The base exists so that the formula does not add $q$ twice [40, 43].

The input tokens change too. A matched row’s first input is the token sampled in step $N$ ; the remaining inputs are the new drafts that step $N$ proposed. Request 12 shows a further consequence. Its step $N$ sample is the end-of-sequence token, number 2 here. The host learns this only when it post-processes step $N$ , after forward $N + 1$ has been queued. Step $N + 1$ therefore evaluates request 12 anyway. Evaluation and publication are separate: the row can be computed, but its result will not be returned to the user. Section 6.4 follows that row to the end of its life.

The batch has 12 query positions, so the worker selects a captured graph for 16 positions and adds four padding rows. There are now eight metadata rows: four real requests with three positions each, and four padding rows with one position each. The script checks that the padding rows keep $q = k = 1$ and remap $- 1$ : the correction applies only to the first n_seq real rows. Section 6.5 explains the padding.

The lab computes the inputs and lengths with this list calculation. It visits only the real rows, leaving padding unchanged, and checks the results against a separate reference based on the known accepted counts:

Listing 6.3. Translate token selection and length correction into lists.

PYTHON

63def correct(rows, prev, remap, q, k, nextn):
64    inputs = []
65    for i, row in enumerate(rows):
66        old = prev[remap[i]] if remap[i] >= 0 else None
67        inputs.append([row.fresh[j] if old is None else
68                       old.sampled if j == 0 else old.drafts[j - 1] for j in range(q[i])])
69        if nextn:  # worker.py:220-227; padded rows are left unchanged
70            k[i] += q[i] - (nextn - old.accepted if old else 0)
71    return inputs

Book lab · labs/14_lagging_metadata.py · lines 63–71

6.2.3 A query length chosen on the device

Chapter 8 adds a third unknown. With variable-length verification, the number of drafts each request verifies depends on confidence estimates from the previous sample. The host fixes only the step’s total query length; the device divides it among requests. Chapter 8 describes that allocation. Here only its dependency matters. The sampler’s host metadata needs the per-request lengths, which now exist only on the device. A non-blocking copy returns them to pinned host memory, and the host must wait for the copy’s event, d2h_evt, before updating that metadata (worker.py, lines 146–152).

The last row of lab 13 adds this case as a one-unit correction queued before forward. It runs after sample $N$ and before forward $N + 1$ on the same stream. Its event is therefore complete soon after the host’s existing wait for sample $N$ ; the step lengthens by the correction’s device time, from 12 to 13 units. That result depends on the assigned durations. The code establishes only the order.

6.3 Read the correction path

The correction happens in Worker._correct_seqlens_qk, called before forward when the worker has a previous decode step. It first builds the host-side inputs:

Listing 6.4. Match the batches, then copy the host arrays to the device.

PYTHON

209remap_cpu, seqlens_q_cpu, seqlens_k_cpu, has_overlap = prepare_zos_inputs(
210    seq_metadata_list, prev.seq_metadata_list, n_seq_with_padding, self.nextn
211)
212if not has_overlap:
213    return None
214remap = remap_cpu.cuda(non_blocking=True)
215seqlens_q = seqlens_q_cpu.cuda(non_blocking=True)
216seqlens_k = seqlens_k_cpu.cuda(non_blocking=True)

DeepInfer · deepinfer/worker/worker.py · lines 209–216

prepare_zos_inputs first tries the cheap case. If both batches have the same length and the same seq_id at every position, row $i$ matches row $i$ (zos.cpp, lines 148–157). Otherwise it builds a map:

Listing 6.5. Match by request identity when the order has changed.

CPP

159bool has_overlap = same_order && !seq_metadata_list.empty();
160if (!same_order) {
161    std::unordered_map<int64_t, int32_t> prev_indices;
162    prev_indices.reserve(prev_seq_metadata_list.size());
163    for (int32_t i = 0; i < static_cast<int32_t>(prev_seq_metadata_list.size()); ++i) {
164        if (!prev_seq_metadata_list[i]->is_placeholder) {
165            prev_indices.insert_or_assign(prev_seq_metadata_list[i]->seq_id, i);
166        }
167    }
168    for (int32_t i = 0; i < static_cast<int32_t>(seq_metadata_list.size()); ++i) {
169        auto prev = prev_indices.find(seq_metadata_list[i]->seq_id);
170        remap_indices[i] = prev == prev_indices.end() ? -1 : prev->second;
171        has_overlap |= remap_indices[i] >= 0;
172    }
173}
174if (!has_overlap) {
175    return {std::nullopt, std::nullopt, std::nullopt, false};
176}

DeepInfer · csrc/engine/zos.cpp · lines 159–176

Placeholder rows from the previous step are excluded; they carry no request. If no row matches, the function returns nothing and the worker skips correction. The host then has complete information for every row in the batch. The function next writes the arrays that the device will correct:

Listing 6.6. Write base lengths for real rows and fixed values for padding.

CPP

184for (int32_t i = 0; i < static_cast<int32_t>(seq_metadata_list.size()); ++i) {
185    const auto *seq = seq_metadata_list[i];
186    remap_ptr[i] = remap_indices[i];
187    seqlens_q_ptr[i] = seq->seqlen_q;
188    seqlens_k_ptr[i] = seq->seqlen_k - (nextn > 0 ? seq->seqlen_q : 0);
189}
190for (int64_t i = seq_metadata_list.size(); i < n_seq_with_padding; ++i) {
191    remap_ptr[i] = -1;
192    seqlens_q_ptr[i] = 1;
193    seqlens_k_ptr[i] = 1;
194}

DeepInfer · csrc/engine/zos.cpp · lines 184–194

Line 188 is the subtraction that defines the base. Lines 190–194 fill the padding rows that the lab checked. Back in _correct_seqlens_qk, the worker queues length correction before calling execute_main_model:

Listing 6.7. Correct lengths on the device; copy query lengths back if they vary.

PYTHON

218if envs.DEEPINFER_ENABLE_MTP_VARLEN:
219    compute_seqlens_q_by_confidence_rates(total_q, remap[:n_seq], prev.sampler_output.confidence_rates_gpu, seqlens_q[:n_seq])
220if self.nextn > 0:
221    compute_seqlens_k_zos(
222        seqlens_k[:n_seq],
223        seqlens_q[:n_seq],
224        remap[:n_seq],
225        prev.sampler_output.num_accepted_tokens_gpu,
226        nextn=self.nextn,
227    )
228
229d2h_evt = None
230if envs.DEEPINFER_ENABLE_MTP_VARLEN:
231    # H2D, correction, and D2H are ordered on the current stream; only varlen rewrites this pinned source.
232    seqlens_q_cpu[:n_seq].copy_(seqlens_q[:n_seq], non_blocking=True)
233    d2h_evt = torch.cuda.Event()
234    d2h_evt.record()

DeepInfer · deepinfer/worker/worker.py · lines 218–234

The first branch is chapter 8’s allocation; the second is the key-length formula. Both slice to n_seq, leaving padding untouched. The comment states the ordering that lab 13 assumed: the host-to-device copy, the correction and the return copy are all queued on the current stream, ahead of forward. The kernel’s arithmetic has a reference implementation in the same module:

Listing 6.8. The reference arithmetic behind the fused kernel.

PYTHON

44seqlens_k += seqlens_q
45mask = remap != -1  # [n_seq] bool
46# Clamp remap to valid index (use 0 for fresh seqs, will be masked out)
47remap_clamped = remap.clamp(min=0)  # [n_seq]
48# Gather prev_num_accepted_tokens for all seqs (invalid ones get tokens[0], doesn't matter)
49accepted = prev_num_accepted_tokens[remap_clamped.long()]  # [n_seq]
50# Compute adjustment: only apply where mask is True
51adjustment = torch.where(mask, nextn - accepted, torch.zeros_like(seqlens_k))
52seqlens_k -= adjustment

DeepInfer · deepinfer/model_executor/layers/tilelang_ops/compute_seqlens_k_zos.py · lines 44–52

Lab 14 translates these lines. Scenarios A and B check matching and fresh inputs without drafts; scenario C checks accepted-length correction for matched requests and leaves padding unchanged. Fresh requests with drafts are not exercised. These checks do not run the kernel or test its handling of real tensor layouts [43].

Model input preparation then uses the corrected query lengths to decide which token positions to gather. Fresh rows get their tokens from the host (zos.cpp, lines 99–113); matched rows read the previous outputs on the device. The gather’s docstring states both branches:

Listing 6.9. Matched rows read the previous outputs; fresh rows read host tokens.

TEXT

21"""Gather input_ids from fresh_tokens or prev sampled/predict tokens.
22
23For each (seq i, token j) with j < seqlens_q[i], let q_base = cu_seqlens_q[i-1]
24(0 for i == 0):
25    if remap[i] >= 0 (old seq):
26        j == 0: out[q_base+j] = prev_sampled_token_ids[remap[i]]
27        j >  0: out[q_base+j] = prev_predict_token_ids[remap[i], j-1]
28    else (fresh seq):
29        out[q_base+j] = fresh_tokens[i, j]
30"""

DeepInfer · deepinfer/model_executor/layers/tilelang_ops/gather_input_ids_zos.py · lines 21–30

prepare_token_ids_zos calls this gather when a remap exists, passing the corrected seqlens_q (prepare_tensor_input_zos.py, lines 235–241). For matched rows, the previous step’s sampled token and new drafts reach the next forward without a round trip through the host [40, 42].

6.3.1 Two sets of metadata describe one step

The corrected lengths feed forward $N + 1$ . Sampling for step $N + 1$ needs host-side metadata too, and it is prepared later. By the time the worker reaches it, it has post-processed step $N$ , so each sequence’s committed count now includes step $N$ ’s accepted tokens. update_metadata_for_sampling_zos uses that count:

Listing 6.10. Rebuild sampling lengths from the newly committed count.

CPP

299metadata->seqlen_k = metadata->seqlen_q + seq.cache.num_tokens;
300if (remap_ptr[i] >= 0) {
301    metadata->refresh_for_sampling(seq, nextn);
302}
303scheduler_outputs.seqlens_q[i] = metadata->seqlen_q;

DeepInfer · csrc/engine/zos.cpp · lines 299–303

The comment above, line 298, explains the choice in the author’s translation: an extra forward after end-of-sequence may not have appended its token to the sequence, so sampling uses the cached length. Request 12 in lab 14 is such a case. Matched rows also refresh their draft and sampling inputs from the sequence. Forward metadata and sampling metadata both describe step $N + 1$ , but they become valid at different times and from different sources: the device correction, then the host’s committed count [40].

6.3.2 When the correction does not run

ZOS is on by default only for some engine types: NORMAL and COMPLETION, not PROMPT (engine_type.py, lines 16–17). EngineArgs applies that default with os.environ.setdefault, so an existing DEEPINFER_ZERO_OVERHEAD_SCHEDULING setting wins (arg_utils.py, lines 51–55). Prompt steps always post-process their own result in the same call; chapter 1 printed that branch. A prompt step can still post-process a preceding decode step, because that happens earlier in the call. Correction returns early when there is no previous state, when the previous step was a placeholder, or when no row matches (worker.py, lines 198–213) [44, 39, 7].

6.4 Last use, reuse and the cost of end-of-sequence

Request 12 sampled its end-of-sequence token in step $N$ . Follow it through three consecutive decode calls:

Call $N$ queues the forward and the sample that choose the end-of-sequence token. The host has not read it.
Call $N + 1$ schedules request 12 again and queues forward $N + 1$ with its row. It then post-processes step $N$ , which marks the request finished and publishes its last token. Sampling for step $N + 1$ still includes the row.
Call $N + 2$ begins with scheduling. free_finished_seqs removes the finished request from the running list and hands it on as finished (scheduler.hpp, lines 562 and 752–763). Forward $N + 2$ no longer includes it. Post-processing step $N + 1$ waits for its sampling event and skips the finished row (lines 905–913). pre_sampler_sync then releases its storage:

Listing 6.11. Finished requests return their storage after the wait.

PYTHON

844if self.engine_type in [EngineType.COMPLETION, EngineType.NORMAL]:
845    for seq in scheduler_outputs.finished_seqs:
846        self.finalize_seq(seq)

DeepInfer · deepinfer/engine/llm_engine.py · lines 844–846

finalize_seq calls the scheduler’s free_seq (lines 804–806), which chapter 4 followed into the block allocator. The last device access to request 12’s storage belongs to step $N + 1$ . The worker waited for that step’s sampling event earlier in call $N + 2$ , so the release follows the last use [5, 33].

The trace assumes consecutive decode calls. A prefill step in between changes when the decode cleanup pass runs, so the release point needs a separate trace.

The extra row is the price of not waiting. Each finished request is evaluated once more after its last published token. As an illustration, suppose there are no drafts and each decode step produces one useful token. If answers averaged 200 tokens, one extra evaluation per answer would add about 0.5 per cent to decode rows, ignoring the first token produced by prefill. That figure is arithmetic, not a measurement. With drafts, both the number of normal decode steps and the number of positions in the extra step change. Chapter 5’s stale-result guard is the same idea applied to a request preempted while its step was in flight, and chapter 1’s final flush drains the one step that no later call will collect.

6.4.1 What makes a release safe

Chapter 4’s second lab released a block while a queued copy still needed it, and a new writer corrupted the copy. That lab did not prove that every release after submission is unsafe. The governing condition is narrower: a conflicting use by a new owner must follow the last access by the old owner. A host can meet it by waiting for an event. It can also rely on an order that is guaranteed some other way, such as two operations on the same stream. Allocator reuse, the lifetime of host objects and the publication of a result are different questions, and DeepInfer answers them differently. Table 6.1 lists the release points met so far.

Resource or release	Condition that permits reuse
Preempted KV blocks	`_preempt`returns allocations at once, while an earlier step may still be in flight. Same-stream order is a plausible explanation for safe reuse, not a complete proof here; it needs the audit below.
Finished request	In consecutive decode calls, the worker has waited for the extra step’s sampling event before `finalize_seq`releases storage.
Fork copy source	`clone_done_evt`completes before the source’s pin is dropped (chapter 4).
Output pool	Returned addresses stay valid for the caller until the next release phase; the caller must clone what it keeps (chapter 1). This is a host-facing contract, not a device event condition.
Sequence metadata	`WorkerStepState`keeps it until post-processing. `LLMEngine.post_process`then collects it for deferred release (line 1013); `DeferredRelease.release`returns it to the pool in `pre_sampler_sync`(lines 100–103).
Borrowed scheduler lists	The copy and zero lists must be consumed before the next scheduling call clears and reuses their storage (`scheduler.hpp`, lines 279–284).

Table 6.1: Release points and the condition that permits reuse. Only the first row lacks a complete argument in the inspected source.

The first row remains open. A complete argument must account for every access to the reused cache: eager forward, graph replay, draft execution, copies and zeroing, and any storage hooks. Graph capture uses its own stream with explicit waits (cuda_graph_runner.py, lines 100–137); that fact alone neither proves nor disproves safe runtime reuse. The release’s storage hooks are empty, so they say nothing about the unreleased service. Until that audit is done, the first row is a scoped inference from the inspected control flow [5, 41, 45].

6.5 Graphs constrain the input buffers

A CUDA graph records a fixed sequence of device operations, with fixed shapes and buffer addresses, so that the host can later replay it with one submission. Fewer submissions mean less host work per step, which lab 13 showed to be the limit on overlap. The cost is rigidity.

DeepInfer captures graphs for a list of batch sizes: powers of two up to an interval, then multiples of that interval, with a default interval of 8 (cuda_graph_runner.py, lines 41–51; environ.py, line 155). The largest size is max_num_seqs times $1 + nextn$ , and PROMPT engines do not use decode graphs (model_runner.py, lines 72–79). Selection uses total_q, the number of query positions, not the number of requests (cuda_graph_runner.py, lines 143–148). Chapter 2’s distinction between requests and positions decides which graph runs. The worker adds padding rows to reach the selected size (worker.py, lines 202–207). Lab 14’s 12 positions became 16.

The graph’s inputs must also be where it expects them. Under ZOS, the runner is built with device-side input storage (model_runner.py, line 78), and the graph data’s docstring says that source tensors share device storage in that mode (cuda_graph_runner.py, lines 26–31). The corrections in this chapter therefore write into buffers that replay reads. Correcting a host copy instead would require the host to read the previous result first, which is the wait ZOS was built to avoid [45].

We can now label each piece of state by the step that produced it, the event or stream order that allows the next read, and the condition that allows reuse. The next chapter spreads this work across devices, where a step must also wait for other devices.

6.6 Experiments to try next

In labs/13_overlap.py, find the smallest post-processing cost that creates a gap between forward and sample in the ZOS run. Derive it from the other durations before running the script.
Set the forward duration to 2 and keep the baseline host costs. Which side limits the ZOS interval? What would graph replay need to reduce to help?
In labs/14_lagging_metadata.py, give request 14 an accepted count of 0 in scenario C. Predict its base, corrected and true key lengths.
Mark one previous row as a placeholder in scenario B. Translate the placeholder check from zos.cpp, lines 151 and 164, into the lab, then predict which rows match.
For one row of the release table, list every access you would need to find to turn its condition into a proof. Separate device accesses from host object lifetimes.