[original-el8]

Problem

Storage lowering (_lower_storage in ignition/preparation/compilation.py) turns a complete, contiguous copy into a persistent region into Inductor's post-graph writeback when no later write invalidates the copy's source. That rule came from 317c1ec and was amended by cbb7d9d. It does not consider a source that reads the destination's own region.

Two row expansions from 022c33b have that shape:

• ignition/components/vocabulary.py: when the vocabulary shards cover the output width exactly, gathered = output[:rows], and __call__ runs output.copy_(gathered[row_indices, :V]).
• ignition/recipes/deepseek41/decoder.py: head_hidden = scratch.hidden[:head_rows], then s.hidden.copy_(head_hidden[samples]).

Inductor inlines each gather into the in-place writeback. In a compiled 12-row DeepSeek V4.1 speculative forward (vocabulary 129280), one Triton kernel loads from and stores to the same logits buffer:

tmp20 = tl.load(in_ptr4 + (x0 + 129280*tmp18), xmask)   # in_ptr4 is the logits buffer; tmp18 in [0, 6)
tl.store(out_ptr3 + (x4), tmp149, xmask)                # sampling logits derived from tmp20
tl.store(out_ptr4 + (x4), tmp20, xmask)                 # out_ptr4 is the same logits buffer

Program instances overwrite rows that other instances still read, and token selection uses those rows. The head hidden buffer has the same pattern.

With speculative decoding, DeepSeek V4.1 Flash (TP4, four DGX Spark hosts) returns wrong tokens when two or more requests are verified in one round: duplicated or skipped tokens, and text from other concurrent requests. Draft acceptance falls at four or more concurrent requests. A single request is unaffected, and --mtp-draft-tokens 0 is correct at every concurrency. The rule is recipe-independent, so any forward with this copy shape is exposed.

Change

A complete copy becomes a writeback only when its source reads no storage region the copy writes (mutable_keys.isdisjoint(accesses[source])). Otherwise it stays an ordered flare::storage_call, so the gather completes before the destination is written. Every other complete copy keeps fusing into Inductor's writeback. The package format is unchanged; existing packages must be re-exported to compile with this rule.

Tests:

• tests/test_compilation.py::test_row_expansion_reads_overlapping_rows_before_writing (CUDA): compiles and replays a prefix-row expansion into the same buffer and compares exactly with eager.
• tests/preparation/test_invocation_views.py::test_copy_lowering_preserves_destination_for_sources_gathered_from_it (CPU): an overlapping gather stays an ordered copy; a gather into a separate buffer remains a writeback.

Validation

Check	Without this change	With this change
CUDA row-expansion test, GB10 (SM121)	58% of elements mismatched after one replay	passes, three runs
CPU copy-lowering test	overlapping case fails	passes
test_compilation.py, test_forward_compile_options.py, test_invocation_views.py on GB10	—	51 passed; bf16_add_out native replay fails with and without this change because the installed Heliopause lacks the qualification feature
DeepSeek V4.1 Flash, temperature-0 lookups and counting prompts, serial	32/32	32/32
Same, concurrency 8 (shared prompts / distinct prompts)	10-18/32	32/32 / 32/32
Same, concurrency 4 / concurrency 2	18/24 / 11/12	24/24 / 12/12
Complete copies kept as writebacks per compiled forward (115 forwards)	38.8	37.4
In-place gather-store kernels on the logits or head hidden buffers	present in nearly every speculative forward	none
Draft acceptance per offered draft, concurrency 4 / 8	0.41-0.50 / 0.33-0.44	0.73-0.77 / 0.55
llm_inference_bench decode, concurrency 4 (8K / 32K / 128K context)	115.4 / 119.9 / 128.3 tok/s	150.9 / 145.0 / 145.2 and 147.9 / 143.9 / 146.2 tok/s (two passes)

Package validation conditions: TP4 across four DGX Spark hosts over RoCE, 8192-token prefill chunk, speculation enabled with adaptive verification. Other benchmark cells (concurrency 1, 2 and 8; prefill; coding) stay within run-to-run variation of the package compiled without this change.

Limitations

• The exported package used the same condition written as a set intersection. Its source tree also carried an uncommitted preparation allocator change that the 8192-token prefill chunk needs to fit in GB10 memory; that change does not touch storage lowering.
• Only DeepSeek V4.1 Flash packages were exported and served. Other recipes compile through the same rule but were not re-exported.
• The first launch of the validated package, right after rank packages were distributed, exited after two minutes on a RoCE collective timeout while rank 0 had about 8 GiB in swap. Relaunches served 244 checked requests and two benchmark passes without failure. The cause was not traced.