feat(ENG-EXPERT-STREAM-DEVICE): W0f — the dense weights were resident twice; --device cuda now decodes, and the token gate fails on a near-tie the alias is measured not to cause (#1299) - #1326
Open
localai-bot wants to merge 8 commits into
Open
Conversation
… twice, and that is what ran the box out (#1299) With W0's lane on, `Qwen3.8-2.4T-A95B UD-Q1_0` LOADS on `--device cuda` on a 119.631 GiB GB10 and then exhausts the machine inside its first forward: zero decode steps, seven attempts, every one identical. The lane was doing its job. The DENSE weights were the cost. The measurements name it rather than a reading of the code. A **0.15 GiB** slot arena died exactly where an 18.55 GiB one did, so the arena is not it. A 1-token prompt, whose protected set fits with no in-place fallback at all, behaved identically to a 5-token one, so prefill protection is not it. Growth was ANONYMOUS — `RssAnon` 8.1 -> 61.4 GB — while file-backed stayed flat, so nothing was pinning the mapping. Host anon plus swap reached ~65 GB against system `used` ~119 GB, and the ~42 GB difference is device memory that unified memory does not charge to RSS. So every non-expert weight was resident twice: once as the host `OwnedTensor`, once as `ResidentWeight`'s device staging copy. About 39 GiB of the 61.20 is `attn_qkv` (21.56) and `ssm_out` (17.25), which the GDN V-head reorder makes `kTransformedWeight` and therefore expands to bf16 in OWNED host buffers. The CPU arm pays that once and serves; the CUDA arm paid it twice and could not. `ResidentWeight` now takes the same branch W0c gave `KqExpertSlice`, on the same probed predicate: where `Platform::host_memory_is_device_addressable()`, it returns a tensor over `w.bytes.data()` instead of `Alloc` + `Copy` into `w.d_dev`. A DISCRETE device answers false, falls through, and is byte-identical to before — asserted by its own case, not by inspection. SAFETY IS BY ALIGNMENT, NOT BY SURVEYING KERNELS, and that is the design decision worth arguing. The staging branch is a verbatim byte copy that returns the same dtype, the same shape and the same dropped marker set, so the only thing any consumer can notice about the substitution is the pointer's alignment. `kDeviceAliasAlignment` is 256 because that is what `cudaMalloc` returns, which makes the two pointers indistinguishable and makes the per-kernel question go away. Deriving a smaller floor does not close: the widest hand-written dereference is a 16-byte `cp.async` granule whose gate checks the SHAPE and assumes the base, while cuBLASLt is separately PROMISED 256 by a preference default this tree never sets. A plain `std::vector<uint8_t>` gives 16 and no more (a large glibc block is an mmap chunk landing at page+16), which is exactly what the transformed weights arrive as, so `MakeHostBytesDeviceAliasable` re-homes an OWNED misaligned buffer once into an aligned block — one memcpy that REPLACES the host-to-device copy it removes. A misaligned BORROW declines and stages instead, because copying a clean file-backed GGUF mapping into anonymous memory would create the residency this change exists to remove, and would break a tied `token_embd`/`lm_head` pair's single keep-alive. Four preconditions were established before the branch was written, because each one could have invalidated it. LIFETIME: `w.bytes` is owned by `Qwen3_5MoeLoadedModel::owned_weights_`, declared before the runner so it outlives every graph, and no reachable site frees, re-points or madvises-away a dense weight's bytes after the GGUF loader returns — `ReleaseHost`'s two callers name only `expert_*_fp4` and `expert_*[se]`, and `AdoptDeviceBytesAsHost` is inert on CUDA because the BACKEND predicate is false. WHICH WEIGHTS: the transformed ones are OWNED anonymous buffers, which is why re-homing is what makes the change move any bytes at all. LAYOUT: no weight on this path needs a device layout different from its host bytes, because the function never produced one; the layout-bearing markers are refused by name instead. DISCRETE: gated. Also fixed in flow, found while establishing the third precondition: #1320. `VT_CPU_QUANT_REPACK` rewrites a Q8_0 weight into the `block_q8_0x4` i8mm interleave at load, only the CPU `MatmulBTKernel` understands that layout, and unlike its sibling `elem_kn_repack` it had NEITHER a CPU-platform gate in the loader policy NOR a refusal here — it rides a HOST-CPU i8mm probe that says nothing about where the weight executes, so an aarch64 box doing `--device cuda` satisfies it. That is wrong tokens, not a crash, and it is precisely what a CUDA-versus-CPU token gate would have reported as a W0f defect. It gets the tripwire its sibling already has, covering both branches. Currently silent on this checkpoint and that is measured, not assumed: one Q8_0 tensor at 0.01% of parameters, and the instrumented load recorded `quant_repack = 0`. Red first: `tests/vllm/model_executor/test_resident_weight_host_addressable.cpp` failed 3 of 6 cases and 7 of 25 assertions on the unchanged tree, for the intended reason — `d_dev` populated, allocs incremented, the tensor pointing at the staged copy. Green after at 9 cases / 45 assertions, over a fake kXPU platform whose backend answers `UnifiedMemory() == true` and `DeviceMemoryIsHostAddressable() == false`, which is the GB10 CUDA backend's own pair rather than an arbitrary one. Nine mutations, each reported with `applied`, `compiled` and a non-zero case count: delete the aliasing branch, make the predicate unconditional, delete each of the three refusals, drop the `borrowed()` guard, claim alignment without providing it, re-home without copying the bytes — all RED. The ninth is the reachability link: corrupting `ResidentWeight`'s host-aliasing arm reds `test_expert_stream_wiring`, which enters through `Qwen3_5Model::Forward`, so the production forward's numbers demonstrably flow through this function. `test_expert_stream_device_slot`'s "an unclaimed tower still stages normally" case moved with the behaviour it describes. It asserted `d_dev != nullptr`; on a host-addressable platform nothing is staged any more, so it now asserts the property it was always about — the refusal did not fire and a usable tensor came back — and gains the discrete arm, where "normally" still means a staged copy. No GPU number is claimed here. G0-CORRECT, G0-LIVE and G0-SPEED remain W0e's, still PENDING on a lease. FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: claude-code:claude-opus-5-1m [Claude Code]
…that can throw, and include what W0f uses (#1299) Three corrections to W0f, none of which changes what any test observes. `MakeHostBytesDeviceAliasable` held a raw `p` from `::operator new` across the `shared_ptr` construction that takes ownership of it, and that construction allocates a control block — so the one throwing step sat inside the one window where nothing owned the allocation. The keep-alive is now built immediately after the allocation and before the memcpy. The madvise still runs against the OLD buffer while it is mapped, which is the ordering that matters and is unchanged. `<new>` and `<cstddef>` are now included where the over-aligned `operator new` / `operator delete` and `size_t` are used, rather than arriving transitively. FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: claude-code:claude-opus-5-1m [Claude Code]
…release, and the first device run could not say why it failed (#1299) Two repairs and one instrument, all found by running W0f rather than reading it. **The use-after-free a fresh review caught.** `MoeBlockBf16Cuda` captures `ResidentWeight(...).data` for all E experts into a DEVICE-resident pointer table, uploads the table once, and then releases the host mirrors. It justifies that with a premise stated in its own comment: "once the device copy exists it is authoritative and nothing reads the host bytes again: every consumer of an expert weight goes through `ResidentWeight`, which returns `d_dev` when populated." That was true while the function had two behaviours. W0f gave it a third: on a host-addressable platform it ALIASES, `d_dev` is never populated, and the captured pointers ARE `w.bytes.data()`. The release then frees the memory the resident table points at while the grouped GEMM keeps reading it for the model's lifetime, including from inside a captured graph. The reviewer demonstrated it with a scratch case replaying the two lines in order: exit code `-11`, `test case CRASHED: SIGSEGV`. Not hypothetical hardware — Qwen3-Coder-30B-A3B BF16 is recorded token-exact 6/6 on `dgx:gpu0`, which is the one GPU this project reaches and the one the predicate answers true on. The repair asks the question the block actually needs: not "did we upload" but "is there a device copy to be authoritative", per weight, which `d_dev` already answers. It is `nullptr` on exactly the arm that aliases and non-null on every arm that staged, so the discrete behaviour the paragraph was written for is unchanged. **The A/B knob the house convention requires.** `VT_ADOPT_DEVICE_BYTES` and `VT_MOE_HOST_FREE` both exist because a default-on residency change needs a same-binary control. W0f shipped without one, and it needs one more than they did: `laguna.cpp` records a MEASURED GB10 penalty for reading system-allocated memory from the GPU instead of a `cudaMalloc` allocation, worst on a long-K low-parallelism GEMV, and `VT_LAGUNA_RESIDENT_BF16W` exists to escape exactly that. W0f installs that retag by default. `VT_QWEN35_ALIAS_HOST_WEIGHTS=0` makes every call decline, so one build measures both arms. **The instrument, and why an RSS curve was not one.** The first device attempt died the same way the pre-W0f runs did — the memory guard tripped at 3.1 GiB available, zero decode steps — and the only evidence was `free -m` every 15 seconds. It shows about 47 GB appearing in 30 seconds at the first forward, and that reading is equally consistent with three different failures: the branch declined and staged as before; the branch re-homed and the old pages did not come back; something else allocated. Those want three different changes, and no amount of staring at the curve chooses between them. So `MakeHostBytesDeviceAliasable` now reports WHICH of its outcomes each weight took and counts the bytes, and `ResidentWeight` prints the split every 4 GiB it has seen, on the existing `VT_LOAD_STATS` switch. Periodic and not at exit, because `[vt load] bytes@exit` is an `std::atexit` handler and the run being measured is one a memory guard SIGKILLs — the one number that would explain the run is the one number the run cannot print. No behaviour changes for a platform that answers the predicate false, and the focused gate is unchanged at 9 cases / 45 assertions. FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: claude-code:claude-opus-5-1m [Claude Code]
…r release depends on (#1299) `MoeBlockBf16Cuda` asked `d_dev != nullptr` inline before releasing each expert's host mirror. That is the right question — it is "is there a device copy to be authoritative", which is exactly what the aliasing arm makes false — but written inline it is a pointer test that reads like a null check, and the next person to add a residency to `ResidentWeight` has nothing to notice. `HostMirrorIsRedundant` gives it a name and a paragraph, and the paragraph is the use-after-free that taught it: the expert pointer table captured `ResidentWeight(...).data`, and on a host-addressable platform those pointers ARE the host bytes the release was about to free, for the model's lifetime and from inside captured graphs. Behaviour is identical on every platform; this is the same test with a name a gate can mutate. FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: claude-code:claude-opus-5-1m [Claude Code]
… the device run did not touch (#1299) **The device run happened, and the row's own stop condition stops it.** One `rc hold` on `dgx:gpu0`, source `9c783a8be`, 4000 slots, greedy, 32 tokens, both arms interleaved on the same lease and the same binary. G0-LIVE PASSES. `--device cuda` produced 32/32 steps where seven previous attempts produced zero, with a decode-phase `exhausted` delta of 0 (6077 at step 1 and at step 32, the structural prefill number this spec predicted), a clean `W0E_DOCKER_RC=0`, peak RSS 97.75 GiB and swap untouched. The instrument added for the run says why: 60.793 GiB of dense weight aliased instead of duplicated into device memory, against ~9.2 GiB that declined (misaligned GGUF borrows) and still stages. G0-CORRECT FAILS, and what the failure MEANS is now measured rather than guessed. The 32 ids match the CPU arm for six tokens and diverge at the seventh. An instrumented CPU run shows that at that exact step the CPU arm's own top-2 is `303` — the token CUDA emitted — behind `7172` by 0.264709 logits on 18.78, or 1.4 %; one step later the margin is 0.022802, about 0.1 %. The CPU arm on the same binary and lease reproduced its recorded ids byte for byte, and the instrument counted `w0f-alias` calls 0 on that arm, so the divergence is the two arms' GEMM arithmetic and W0f cannot reach it. The declared gate still fails and the wave still stops, which is correct; whether a token-exact cross-arm gate is the right instrument for a greedy path this finely balanced is an operator decision and is carried under `## Owed`. G0-SPEED is therefore VOID and NOT claimed, though it was taken: steady-state 4.09-5.69 s/token on CUDA against 8.04-9.27 on CPU. A speed number behind a failing correctness gate is exactly the shape #912 F1 was. **Three review findings the run did not cover.** A direct-upload borrow that happens to be 256-aligned took the alias branch and so skipped issue #150's windowed page release — a third path past a release whose own comment insists it happens on every path, data-dependent at roughly one borrow in eight; it now calls `ReleaseDirectUploadSource` before returning. The sentence "a borrow owns no anonymous pages", which three separate places reason from, is no longer universally true now that re-homing creates one, and the exception is written down beside the code that creates it. The re-pointing has no memo and is therefore unsynchronised, which is safe only because first touch happens inside one forward on one thread; that precondition is stated rather than left to be rediscovered. Also corrected: the header claimed 256 "because that is what `cudaMalloc` returns". CUDA guarantees only "suitably aligned", so the honest basis is cuBLASLt's `CUBLASLT_MATMUL_PREF_MIN_ALIGNMENT_A_BYTES` default, which this tree never sets and which dominates the strictest explicit in-tree gate (32). And the claim that no consumer can tell the two pointers apart is now scoped: alignment makes the substitution CORRECT, but `laguna.cpp` records a measured GB10 bandwidth penalty for system-allocated memory, and the Vulkan and Metal backends distinguish pointers by identity. Both are why `VT_QWEN35_ALIAS_HOST_WEIGHTS` exists. The test binary's three arm switches became a scope guard, so a `REQUIRE` that aborts a case body can no longer leak the discrete arm into every later case. FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: claude-code:claude-opus-5-1m [Claude Code]
…case, so their mutations can fail (#1299) Both repairs in the previous commit were correct and neither was gated, which means both would have survived a mutation that deleted them. Two cases close that, and each one exists because a specific mutation came back GREEN against me until it did. `HostMirrorIsRedundant` is now asserted on both arms: false for an aliased weight, where the host bytes are the only copy and releasing them frees what the kernel reads, and true for a staged one, where the pre-W0f release stays correct. The second half is not decoration -- without it the invariant is satisfied by refusing every release, which would silently undo a measured host-memory lever. The case deliberately does NOT dereference the freed buffer: a segfault is a red that also destroys the rest of the binary's report, so it asserts the decision the production site now asks instead. The direct-upload page release on the alias branch gets a case in the file that already owns that behaviour. It builds an `mmap` borrow, asserts that the borrow really is 256-aligned so the case cannot pass by declining, aliases it, and checks the consumed source pages went away. `mmap` always returns page-aligned memory, so this is every direct-upload borrow on such a platform rather than a corner. Focused gate 9 cases / 45 assertions to 10 / 51; `test_load_direct_upload` 14 / 187 to 15 / 193. FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: claude-code:claude-opus-5-1m [Claude Code]
…not the one I remembered (#1299) The header justified `kDeviceAliasAlignment = 256` partly by saying it dominates "every explicit pointer gate in the tree (the strictest is 32, in `src/vt/cuda/cuda_nvfp4_sm12x.cu`)". Both halves are wrong, and a number quoted in a rationale is exactly the kind that later gets treated as measured. Checked rather than recalled. `grep -rn MIN_ALIGNMENT src/vt/` returns nothing, so cuBLASLt's documented 256-byte `MIN_ALIGNMENT_A_BYTES` default really does apply to every matmul this tree issues, which is the load-bearing half and survives. The other half does not: the only genuine POINTER-alignment gate in the CUDA kernels is `cuda_matmul_nvfp4.cu`'s `reinterpret_cast<uintptr_t>(prow) & 0xf`, asking for 16. The `% 32` and `% 64` tests that read like alignment gates are dimension checks on `d` and `dv`, not on an address. The conclusion is unchanged and now rests on what is there: 256 is at least what every consumer is promised. FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: claude-code:claude-opus-5-1m [Claude Code]
…he ends of the range (#1299) The W0e figures went in as "steady-state 4.09-5.69 s/token" against "8.04-9.27", read off the last six steps of each arm. That is a range whose ends are the least representative numbers in it, chosen by where the log happened to be tailed. Recomputed over all 31 DECODE steps of each arm, excluding step 1 because it is prefill and not a decode step: CUDA min 3.012, median 4.598, max 126.456; CPU min 7.857, median 9.055, max 23.174. Both maxima are the first decode step, with the slot cache cold, which is why the medians are the figures and why the earlier six-step window flattered both arms by starting after that. Two cautions ride with the numbers, because they will outlive this commit. The implied 1.97x is NOT a result: it rests on a token comparison that FAILED, and the row's stop condition voids it. And this CPU arm is faster than the 11.05 s/token previously recorded at 4000 slots, so the same-lease interleaved denominator taken here and that earlier figure are different measurements and must not be mixed into one ratio. FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: claude-code:claude-opus-5-1m [Claude Code]
localai-bot
changed the title
localai-bot
changed the title
Collaborator
Author
Fresh-implementer repair of the review FAIL — the new head is on a separate refThe repaired history is pushed as Why the history had to change. What that costs. Moving this PR onto the repaired history needs a force-push
The PR body above has been rewritten and is the squash commit message. The finding that decided itThe G0-CORRECT attribution rested on a control that is true by construction. The discriminating experiment ran on Identical algo and bit-exact output, so W0f cannot move a logit. It ran on Gate output
One earlier full-gate attempt is recorded VOID rather than quietly rerun: a |
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
2 participants
Add this suggestion to a batch that can be applied as a single commit.This suggestion is invalid because no changes were made to the code.Suggestions cannot be applied while the pull request is closed.Suggestions cannot be applied while viewing a subset of changes.Only one suggestion per line can be applied in a batch.Add this suggestion to a batch that can be applied as a single commit.Applying suggestions on deleted lines is not supported.You must change the existing code in this line in order to create a valid suggestion.Outdated suggestions cannot be applied.This suggestion has been applied or marked resolved.Suggestions cannot be applied from pending reviews.Suggestions cannot be applied on multi-line comments.Suggestions cannot be applied while the pull request is queued to merge.Suggestion cannot be applied right now. Please check back later.
With W0's lane on,
Qwen3.8-2.4T-A95B UD-Q1_0(369.97 GiB) LOADED on--device cudaon a 119.631 GiB GB10 and then exhausted the machine inside itsfirst forward: zero decode steps, seven attempts, every one identical (#1299).
It now decodes. 32/32 steps, and the token gate that would let us publish a
rate does not pass.
What the defect was
The measurements in #1299 name it rather than a reading of the code. A
0.15 GiB slot arena died exactly where an 18.55 GiB one did, so the arena is
not the cost. A 1-token prompt, whose protected set fits with no in-place
fallback at all, behaved identically to a 5-token one, so prefill protection is
not it. Growth was ANONYMOUS (
RssAnon8.1 to 61.4 GB) while file-backed stayedflat, so nothing was pinning the mapping. Host anon plus swap reached ~65 GB
against system
used~119 GB, and the ~42 GB difference is device memory thatunified memory does not charge to RSS.
So every non-expert weight was resident twice: once as the host
OwnedTensor,once as
ResidentWeight's device staging copy. On a part where device memory IShost memory, the second copy buys nothing and costs everything.
The change, and the one decision worth arguing
ResidentWeightnow takes the same branch W0c gaveKqExpertSlice, on the sameprobed predicate: where
Platform::host_memory_is_device_addressable(), itreturns a tensor over
w.bytes.data()instead ofAlloc+Copyintow.d_dev. A DISCRETE device answers false, falls through, and is byte-identicalto before, asserted by its own case rather than by inspection.
Safety is by alignment, not by surveying kernels. The staging branch is a
verbatim byte copy returning the same dtype, shape and dropped marker set, so
the only thing a consumer can notice about the substitution is the pointer's
alignment.
kDeviceAliasAlignmentis 256 because cuBLASLt'sCUBLASLT_MATMUL_PREF_MIN_ALIGNMENT_A_BYTESdefaults to 256 and this tree neversets it (
grep -rn MIN_ALIGNMENT src/vt/is empty), which dominates everyexplicit pointer gate in the CUDA kernels — there are at least seven across four
files and the strictest asks 32, not the single 16-byte gate two earlier
drafts of that comment claimed. A plain
std::vector<uint8_t>gives 16 and nomore, so
MakeHostBytesDeviceAliasablere-homes an OWNED misaligned buffer onceinto an aligned block: one memcpy that REPLACES the host-to-device copy it
removes. A misaligned BORROW declines and stages instead, because copying a clean
file-backed GGUF mapping into anonymous memory would create the residency this
change exists to remove.
Can the substitution move a logit? Measured, and the answer is no
The first version of this pull request attributed the step-7 token divergence to
CPU-versus-CUDA GEMM arithmetic on two grounds, and a fresh review showed both
are vacuous. "The instrument counted
w0f-aliascalls 0 on the CPU arm" istrue by construction for every possible state of W0f, correct or corrupt:
ResidentWeighttakes anis_cpu()early return about ninety lines above thealias branch, so the CPU arm can never reach it. "The CPU arm reproduces its own
reference" constrains only the arm W0f cannot reach. Both show the branch is
platform-gated. Neither discriminates. They are withdrawn.
The only thing a consumer can notice about the substitution is the pointer, so
that is what was measured, on
thor:gpu0(NVIDIA Thorsm_110, driver 13020,cudart 13000, cuBLASLt 130101). Thor answers this branch's own predicate TRUE
—
cudaDevAttrPageableMemoryAccess = 1,cudaDevAttrIntegrated = 1— so it is amember of the population W0f serves, and a
cudaMallocpointer there is a realdevice pointer exactly as on GB10.
The probe transcribes both cuBLASLt formulations out of
src/vt/cuda/cuda_matmul.cuat this branch — row-major NN
MatmulKernelCuda, where the weight is operand B,and column-major TN
MatmulBTKernelCuda, where it is operand A — over six shapesoff the checkpoint's own
embedding_length = 8192at M = 1, 5 and 32.Twelve measurements,
PROBE_FAILURES=0.default == MIN_ALIGNMENT 256cublasLtMatmul, weight fromcudaMallocvs a 256-aligned HOST blockSUCCESSThe selection is reported whole rather than by id. At M=1 N=8192 K=8192, both
layouts:
id=66 tile=573 stages=35 splitK=5 reduction=2 swizzle=0 custom=1 inner=0 ws=163856 waves=0.8000, identical across all four queries. Theinstrument discriminates: five DIFFERENT configurations appear across the six
shapes (tiles 393, 537, 573, 576; workspaces 0 through 5,242,896), so a uniform
answer is not a probe that reports one thing regardless of its input.
The structural reason needs no lease and is checkable by reading:
cublasLtMatmulAlgoGetHeuristictakes(handle, operationDesc, Adesc, Bdesc, Cdesc, Ddesc, preference, count, results, returned)and no operand pointers,so alignment can reach the heuristic only through
CUBLASLT_MATMUL_PREF_MIN_ALIGNMENT_*_BYTES, which this tree never sets.Identical algo AND bit-exact output, so W0f cannot move a logit, and the
step-7 divergence is the two arms' GEMM arithmetic. Two things this does not
establish. It ran on Thor, not on the GB10 the token gate ran on — same
predicate class, not the same silicon; the GB10 leg is queued on
dgx:gpu0,costs seconds, needs no model load, and is carried under
## Owed. And the16-aligned arm also came back bit-exact, which is NOT a licence to lower
kDeviceAliasAlignment: twelve shapes is not the enumeration, and cuBLASLt isstill promised 256.
Fixed in flow
#1320.
VT_CPU_QUANT_REPACKrewrites a Q8_0 weight into theblock_q8_0x4i8mm interleave, only the CPU
MatmulBTKernelunderstands that layout, andunlike its sibling
elem_kn_repackit had NEITHER a CPU-platform gate in theloader policy NOR a refusal here. It rides a HOST-CPU i8mm probe that says
nothing about where the weight executes, so an aarch64 box doing
--device cudasatisfies it. That is wrong tokens, not a crash, and it is precisely what a
CUDA-versus-CPU token gate would report as a W0f defect. It gets the tripwire
its sibling already has, on both branches.
What two fresh reviews found, and what each repair is
A use-after-free, first instance.
MoeBlockBf16CudacapturesResidentWeight(...).datafor every expert into a device-resident pointer tableand then releases the host mirrors, justified by "once the device copy exists it
is authoritative". W0f made that false: the aliasing arm never populates
d_dev, so the captured pointers ARE the bytes being freed, for the model'slifetime and from inside captured graphs. The first reviewer demonstrated it with
a scratch case taking SIGSEGV.
HostMirrorIsRedundantnames the invariant therelease actually needs.
The same use-after-free, second instance, and the gate was blind to it.
ReleaseResidentQwen3_5DenseHostWeightsguarded ond_dev || d_dev_f32.d_dev_f32is a bf16-to-f32 UPCAST into a separate allocation; it is not a copyof those bytes and can never stand in for them.
PrepareBf16Residentpassesexactly four weights to BOTH
raw()andf32()—gdn.conv1d_weight,gdn.norm_weight,attn.q_norm,attn.k_norm— and on the aliasing armraw()leaves
d_devnull whilef32()setsd_dev_f32, so the disjunction passed andfreed the bytes the aliased raw tensor points at. The second reviewer's mutation
M10 weakened
HostMirrorIsRedundantto exactly that disjunctive form and thesuite stayed 10/10 GREEN. The site now asks the invariant, and a new case reds
that mutation. Nothing reaches the combination today, and the reason is an
accident worth writing down:
DirectDeviceLoadEligiblerequires!platform.is_unified_memory(), and on CUDAis_unified_memory()andhost_memory_is_device_addressable()are the SAMEpageable && integratedconjunction computed independently in
cuda_backend.cu:363andplatforms/cuda.cpp:215— an equality nothing states, documents or gates.A refusal that fired above the memo it should have deferred to. The
"no host bytes"
VT_CHECKsat aboveif (!w.d_dev). Pre-W0f a weight with apopulated
d_devand a released host mirror returned fine fromd_dev; afterW0f it threw. Its justification is true of the dense weights and false of the
expert weights the same function serves, whose misaligned GGUF borrows decline
the alias, stage, get a
d_dev, and are then released by the guarded loop besidethe pointer capture — so this change created the population. The condition is now
"nothing to serve", not "no host bytes".
A page release that repeated once per forward step.
MakeHostBytesDeviceAliasable's aligned-in-place branch callsReleaseDirectUploadSource, andResidentWeightcalls it with no memo — about1,361 times per decode step.
AdoptDeviceBytesAsHostcalled it exactly once,behind
if (!w.d_dev). WithLoadWindowedReleaseEnabled()default-ON, everydecode step would
madvise(MADV_DONTNEED)the weight the GPU is about to read,which then re-faults. Correctness survives; throughput would not. Consuming the
mmap_srcrecord IS the memo, and the new case asserts the COUNT by restoringthe pattern and looking for it again — the first version asserted only that the
release happened, which a release that happens every time also satisfies.
A skipped page release (found by the first review). A direct-upload borrow
that happens to be 256-aligned took the alias branch and so skipped issue #150's
windowed release, a third path past a release whose own comment insists it
happens on every path.
mmapalways returns page-aligned memory, so that wasevery such borrow, not a corner.
Gates, all three reported as they fell
Run on
dgx:gpu0inside onerc hold, source9c783a8be, 4000 slots, greedy,32 tokens, both arms interleaved on the SAME binary and lease.
G0-LIVE: PASS. 32/32 steps where seven previous attempts gave zero.
Decode-phase
exhausteddelta 0 (6077 at step 1 and at step 32; the total isthe structural prefill number this spec predicted, and gating the total would
report a red for a healthy lane).
W0E_DOCKER_RC=0, no guard trip, peak RSS97.75 GiB, swap untouched. The instrument added for the run counts
60.793 GiB of dense weight aliased instead of duplicated — first-forward
totals, at the point re-homing plateaus, call 1361 — against ~9.2 GiB that
declined and still stages. The qualifier is part of the number: the counters are
per CALL and there is no memo on the alias branch, so quoted bare the same figure
is a traffic count and not a residency measurement. Peak RSS is the independent
corroboration.
G0-CORRECT: FAIL. The ids match the CPU arm for six tokens and diverge at the
seventh: CPU
7172, CUDA303. At that step the CPU arm's own top-2 is303, the token CUDA chose, behind by 0.264709 logits on 18.78 (1.4 %);one step later the margin is 0.022802. The arms rank the same candidates and
disagree about a coin flip.
G0-SPEED: VOID, by this row's own stop condition, and deliberately not led
with. Taken for the record only, over the 31 decode steps of each arm: CUDA
median 4.598 s/token, CPU median 9.055. The implied ratio rests on a token
comparison that failed, and this CPU arm is faster than the 11.05 s/token
previously recorded, so the two are different measurements and must not be mixed.
Owed, stated rather than elided
thor:gpu0,which is in the same predicate class, but the token gate ran on
dgx:gpu0.The job is queued there behind a four-hour render; it needs no model load.
test_expert_stream_wiringentersQwen3_5Model::Forwardand M-R9b/M-R9c redit, but on the CPU device, which returns above the alias branch. In CI the
branch is reached only through
detail::StageWeightForTest, a test-only seam.The device evidence is real (43,501 entries through that entry point on
dgx:gpu0) and is not repeatable in CI. Listed under## Owedin the spec, as## Nothing lands deadrequires.certainly a non-host pointer". Nothing printed the pointer and nothing called
cudaPointerGetAttributeson it. In a change whose central risk is handingdevice kernels host pointers, that is the finding not to dismiss.
stacked on
row/ENG-EXPERT-STREAM-DEVICE-W0at95883dcae, which is not onmainand whose PR was closed while this repair was in flight. That is alanding blocker for the operator, not a defect in this change.
Evidence
Red first: the three repairs failed 3 of 12 cases and 4 assertions on the
unchanged tree, and the page-release repair failed 2 of 16 cases and 3
assertions, each for the intended reason. Green after at 12 cases / 71
assertions (
test_resident_weight_host_addressable) and 16 / 203(
test_load_direct_upload). The row's earlier claim of "9 cases / 45" wasalready wrong at the head it described; the count now comes from the binary's own
last
test cases:line.Mutations
Eleven, each reported with the four facts a mutation result is worthless without:
that the edit applied (
git diff --statnon-empty), that it compiled (anon-building mutation is INVALID, not a pass), a non-zero case count (a
filter matching nothing prints SUCCESS), and the binary's exit code captured
directly rather than through a pipe. Whole binaries, no
-tcfilter. Every onerestored and verified byte-identical by sha256.
d_dev_f32disjunct back at the dense release siteHostMirrorIsRedundantto the disjunctive form (the reviewer's M10)mmap_src, so the release repeats every callResidentWeight's CPU branchResidentWeight's first lineM-R2 is the one the second review asked for: weakening the invariant to
d_dev != nullptr || d_dev_f32 != nullptrleft the suite fully green before thischange and reds it now.
M-R9 is reported as GREEN because it is, and it is a finding rather than a
pass. A mutation that corrupts the weight's VALUE does not move
test_expert_stream_wiring, because that test asserts lane behaviour and nottokens. M-R9b and M-R9c disambiguate it: making the same function THROW reds
4/4 cases, so a production forward does reach
ResidentWeight— the call isproven, the value is not checked. The earlier claim that "the reachability
mutation corrupts
ResidentWeightand redstest_expert_stream_wiring" is trueonly for a mutation that throws, and is corrected here.
Issue: #1299
Also fixes: #1320
Spec:
.agents/specs/expert-stream-device-slots.mdBase:
row/ENG-EXPERT-STREAM-DEVICE-W0at95883dcae, because the W0 lane thisbuilds on is not on
main.FOLLOWING_AGENTS_PROTOCOL
Following-Agents-Protocol: true
AI-Assisted: true
Assisted-by: claude-code:claude-opus-5-1m [Claude Code]