Skip to content

record(ORACLE-EXLLAMAV3-GATEABILITY): the pin does not build on aarch64, so the EXL3 arm still has no token oracle (#1901) - #1967

Merged
localai-bot merged 4 commits into
mainfrom
row/ORACLE-EXLLAMAV3-GATEABILITY
Aug 26, 2026
Merged

record(ORACLE-EXLLAMAV3-GATEABILITY): the pin does not build on aarch64, so the EXL3 arm still has no token oracle (#1901)#1967
localai-bot merged 4 commits into
mainfrom
row/ORACLE-EXLLAMAV3-GATEABILITY

Conversation

@localai-bot

@localai-bot localai-bot commented Aug 26, 2026

Copy link
Copy Markdown
Collaborator

exllamav3 was registered gateable = no because nobody had built or run it.
#1901 asked for the measurement. It is now measured, and the verdict stands for
a structural reason rather than an unattempted one: the pin is an x86_64
project and every fleet device is aarch64.

Three rc leases on orin:gpu0, Ubuntu 24.04 aarch64, gcc 13, CUDA 13.0.88,
torch 2.13.0+cu130, a fresh clone at 2398c056 (__version__ = "1.4.3"),
driven with ninja -k 0 so nothing stopped at the first error:

edges 129 · objects produced 122 · FAILED 7
  immintrin.h: No such file or directory            x3
  '__builtin_cpu_supports' was not declared         x3
  identifier "__builtin_ia32_pause" is undefined    x2

122 of 129 translation units compile, including every CUDA kernel this
campaign mirrors — all of quant/ with its 77 comp_units, dsv4_compress.cu,
dsa_topk.cu, hc_mix.cu, routing.cu, attention.cu. The GPU side is not
the obstacle. All seven failures are x86-only host code, and one mistake
explains them: #ifdef __linux__ where the author meant "x86"
(avx2_target.h:10-13, avx512_target.h:9-11, and the
__builtin_ia32_pause() spin loops at cpu/moe_handoff.cu:151-156,
parallel/all_reduce_cpu.cu:110-114). Upstream treats this as a platform fact,
not a defect: its CI declares oses = ["ubuntu-22.04", "windows-2022"] and its
arch list 8.0 8.6 8.9 9.0 10.0 12.0+PTX contains none of GB10's 12.1, Thor's
11.0 or Orin's 8.7.

Porting it would mean repairing five guards and finding a non-x86 path for
~276 lines of intrinsics across the two all_reduce_cpu_avx*.cpp files
(moe_mul1.cpp already has a scalar tier). That needs its own row — and a
patched exllamav3 would not carry the pin's authority, so gateable stays no
either way.

Consequence for the campaign, stated plainly: the EXL3 arm has no token
oracle.
vLLM implements no EXL3 at our pin (verified: zero files matching
exl3|exllamav3|trellis), and the registered secondary cannot run here.

Two facts established for whoever picks up the run half: the text component
would fit — 13.61 GiB carried + 82.59 GiB trellis = 96.20 GiB against
119.63 GiB, since exllamav3 dequantizes the carried FP8/FP4 to FP16 at load,
while adding the MTP drafter reaches 127.54 GiB and does not — and the
checkpoint's namespace matches deepseek_v4.py key-for-key.

Three corrections to the record beyond the verdict:

  • The ancestry this file left open is settled. git fetch --unshallow
    shows the checkpoint's declared quantizer revision 787d1582 is an
    ancestor of the pin, 161 commits back at upstream's v1.2.1, and only one
    commit between them touches the EXL3 storage path — a compute fusion, not a
    layout change. So the pin reads what v1.2.1 wrote. Note DeepSeek-V4 support is
    absent at that revision.
  • The existing TP1 artifact is not loadable by this oracle. DeepSeek-V4-Flash on one GB10 at 44-47 tok/s decode (MiaAI-Lab SparkInfer, EXL3 3.0bpw REAP-K216): load the same quants and match or beat it #1875's
    coalescer emits {prefix}.rank0.{field} while Linear.is_exl3_storage
    requires {key}.trellis with no rank segment. A native-key variant is staged
    beside the evidence.
  • One error in the first run was the measurement's own — a missing
    libcusparse-dev-13-0 produced three unrelated cusparse.h errors. Recorded
    rather than dropped.

Not measured, and named as inferred where inferred: the GB10 itself.
dgx:gpu0 went unhealthy at ~07:48Z and thor:gpu0 at ~08:01Z — two of three
devices inside an hour, the same failure this file already recorded on
2026-08-25. Nothing is queued to close that gap: an rc run submission dies with
the client that made it, so the job believed to be waiting on dgx does not
exist. The oracle file and the NAS README carry the re-submit command instead
of a promise. The gap is narrow (avx2_target.cpp is host C++ with no CUDA,
and GB10 runs the same aarch64 worker image), but it is an inference and says
so.

#1901 stays open for that confirmation and for the run half. Evidence under
/mnt/nas_share/rc/exllamav3-gateability/ with a README-1901.txt index.

Two process corrections this branch carries. The first push (fbb93d49f)
was written before origin/main moved, so its diff-scoped commit-trailers
and commit-style gates skipped while the run exited zero — the documented
trap of reading an exit code instead of the verdict line, which printed NOT a green preflight: a skipped gate reported nothing about this tree. The base is
merged and both gates now run and pass. The superseded head is merged in rather
than force-pushed away.

FOLLOWING_AGENTS_PROTOCOL

Following-Agents-Protocol: true
AI-Assisted: true
Assisted-by: Claude:claude-fable-5 [Claude Code]

mudler added 4 commits August 26, 2026 08:37
…64 (#1901)

`.agents/oracles/exllamav3.md` carried `gateable = no` because nobody had ever
built or run exllamav3 here. That field is now measured rather than unfilled.

The pin does not build on this fleet. Inside `rc` leases on `orin:gpu0`, on
Ubuntu 24.04 aarch64 with gcc 13, CUDA 13.0.88 and torch 2.13.0+cu130, a
worktree-clean checkout of `2398c056` fails. Driving the build with `ninja -k 0`
so nothing stops at the first error puts a number on it: 122 of 129 translation
units compile, and the 7 that fail are all x86-only host code. Every CUDA kernel
this row mirrors compiles, so the GPU side is not the obstacle.

One mistake explains all seven: `#ifdef __linux__` used where the author meant
"x86". The extension carries no architecture guard anywhere, upstream's CI
declares `oses = ["ubuntu-22.04", "windows-2022"]`, and its CUDA arch list holds
none of this fleet's devices. So the verdict stays `no`, and the reason changes
from "nobody tried" to a named, quotable failure with a measured size.

Reading the source found five of those files. The compiler found two more, both
`__builtin_ia32_pause()` spin loops inside `.cu` files, which is why the
enumeration was worth a second lease.

One error in the first run was the measurement's own fault, and the record says
so: a missing `libcusparse-dev-13-0` produced three `cusparse.h` errors that
have nothing to do with exllamav3. Installing it left `avx2_target.o` as the
only failure. An instrument gap that fails toward a code verdict is the trap
this project keeps paying for, so it is written down rather than dropped.

Three things the same session settled that the file previously left open. The
checkpoint's declared quantizer revision `787d1582` IS an ancestor of the pin,
161 commits behind it, at upstream's v1.2.1 tag; the earlier draft could not say
so because its clone was shallow, and it correctly refused to read that absence
as evidence. Only one commit between them touches the EXL3 storage path, and it
is a compute fusion rather than a layout change. And the TP1 artifact the row's
denominator run already produced is NOT loadable by this oracle, because that
coalescer emits `{prefix}.rank0.{field}` while `Linear.is_exl3_storage` requires
`{key}.trellis` with no rank segment.

What could not be measured is recorded as such. `dgx:gpu0` and `thor:gpu0` both
went `unhealthy (no contact)` during the session, so the GB10 itself was never
reached and the run half was never attempted. The gap is narrow, because the
failing code is a host C++ translation unit with no CUDA in it and GB10 runs the
same aarch64 worker image, but the file says "inferred" where it means inferred.

FOLLOWING_AGENTS_PROTOCOL

Following-Agents-Protocol: true
AI-Assisted: true
Assisted-by: Claude:claude-fable-5 [Claude Code]
…64 (#1901)

`.agents/oracles/exllamav3.md` carried `gateable = no` because nobody had ever
built or run exllamav3 here. That field is now measured rather than unfilled.

The pin does not build on this fleet. Inside `rc` leases on `orin:gpu0`, on
Ubuntu 24.04 aarch64 with gcc 13, CUDA 13.0.88 and torch 2.13.0+cu130, a
worktree-clean checkout of `2398c056` fails. Driving the build with `ninja -k 0`
so nothing stops at the first error puts a number on it: 122 of 129 translation
units compile, and the 7 that fail are all x86-only host code. Every CUDA kernel
this row mirrors compiles, so the GPU side is not the obstacle.

One mistake explains all seven: `#ifdef __linux__` used where the author meant
"x86". The extension carries no architecture guard anywhere, upstream's CI
declares `oses = ["ubuntu-22.04", "windows-2022"]`, and its CUDA arch list holds
none of this fleet's devices. So the verdict stays `no`, and the reason changes
from "nobody tried" to a named, quotable failure with a measured size.

Reading the source found five of those files. The compiler found two more, both
`__builtin_ia32_pause()` spin loops inside `.cu` files, which is why the
enumeration was worth a second lease.

One error in the first run was the measurement's own fault, and the record says
so: a missing `libcusparse-dev-13-0` produced three `cusparse.h` errors that
have nothing to do with exllamav3. Installing it left `avx2_target.o` as the
only failure. An instrument gap that fails toward a code verdict is the trap
this project keeps paying for, so it is written down rather than dropped.

Three things the same session settled that the file previously left open. The
checkpoint's declared quantizer revision `787d1582` IS an ancestor of the pin,
161 commits behind it, at upstream's v1.2.1 tag; the earlier draft could not say
so because its clone was shallow, and it correctly refused to read that absence
as evidence. Only one commit between them touches the EXL3 storage path, and it
is a compute fusion rather than a layout change. And the TP1 artifact the row's
denominator run already produced is NOT loadable by this oracle, because that
coalescer emits `{prefix}.rank0.{field}` while `Linear.is_exl3_storage` requires
`{key}.trellis` with no rank segment.

What could not be measured is recorded as such. All three fleet devices went
`unhealthy (no contact)` during the session, so the GB10 itself was never
reached and the run half was never attempted. The gap is narrow, because the
failing code is a host C++ translation unit with no CUDA in it and GB10 runs the
same aarch64 worker image, but the file says "inferred" where it means inferred.

The file also does not claim a confirmation is waiting. A build-only job was
queued on `dgx:gpu0` and an earlier draft of this record said it would land when
the box returned. An `rc run` submission dies with its client, so it did not
survive the session that made it. The record now carries the command to submit
again instead of a promise nothing will keep.

FOLLOWING_AGENTS_PROTOCOL

Following-Agents-Protocol: true
AI-Assisted: true
Assisted-by: Claude:claude-fable-5 [Claude Code]
`origin/main` moved from `3f4161ee8` to `b62354b5a` while this row was measured,
five commits ahead. That left the branch behind its own base, and the two
diff-scoped commit gates then SKIPPED rather than ran: `agent-preflight.sh
--staged` exited 0 while printing "NOT a green preflight: a skipped gate
reported nothing about this tree". Merging restores the base so those gates walk
this branch instead of declining to.

The merge is textual only. It touches no file this row owns, and
`.agents/oracles/exllamav3.md` is byte-identical across it.

FOLLOWING_AGENTS_PROTOCOL

Following-Agents-Protocol: true
AI-Assisted: true
Assisted-by: Claude:claude-fable-5 [Claude Code]
The first push carried `fbb93d49f`, which was written before `origin/main`
moved and therefore had its diff-scoped `commit-trailers` and `commit-style`
gates SKIP while the run still exited zero. The corrected record supersedes it
byte-for-byte; this merge exists only so the update reaches the pull request
branch without a force push.

FOLLOWING_AGENTS_PROTOCOL

Following-Agents-Protocol: true
AI-Assisted: true
Assisted-by: Claude:claude-fable-5 [Claude Code]
localai-bot pushed a commit that referenced this pull request Aug 26, 2026
…lkan red found while gating W2 (#1979)

Not this branch's defect and not this branch's fix. `windows-msvc-vulkan` stops
on `test_openai_api_server.exe exited with status -1073740791` on this pull
request and on #1975, #1969 and #1967, whose diffs share no file under
src/vllm/entrypoints/. The job is PR-only and never runs on main, so there is no
green reference commit to bisect against, which is why four authors could each
meet this red without any of them being able to attribute it.

It rides here rather than in its own record-only pull request because the
alternative to filing it now is not filing it: AGENTS.md requires a bug found
during other work to get an issue that names an owner, and this one is owned by
ENG-RELEASE-WINDOWS rather than by the row that found it.

FOLLOWING_AGENTS_PROTOCOL

Following-Agents-Protocol: true
AI-Assisted: true
Assisted-by: Claude:claude-fable-5 [Claude Code]
@localai-bot
localai-bot merged commit d9a5285 into main Aug 26, 2026
42 of 47 checks passed
localai-bot added a commit that referenced this pull request Aug 26, 2026
…#699) (#1969)

W4b-2 of the dots3-note row (#699). The 33 `sliding_attention` layers
and the
PADDED physical KV row are on the decode path, reached from
`ModelRegistry::Forward`. W4b-1 landed the sliding maths as host code
with no
production call site and named W4b-2 as the row that would wire it; this
is that
wiring, plus the two windowed kernels it needed.

Upstream re-derived at vLLM `origin/main` = `bc2d63e650`.
`git diff --stat d9fbe526c0 -- vllm/models/dots3_note/` is EMPTY against
`bc2d63e650` and against `origin/main` alike, measured in a local vLLM
checkout
rather than assumed, so every §2.3 anchor W4b-1 corrected holds
unchanged at the
newer head.

## The window is a kernel bound here and a gather upstream, and that is
the same function

`vt::MlaDecodeAttention` and `vt::MlaPrefillAttention` each grow an
optional
`AttentionWindow` — the `(left, right)` pair this tree already uses on
`PagedAttentionArgs`, and literally the pair upstream hands
FlashAttention
(`run_sliding_window(..., causal=True, window_size=(sliding_window - 1,
0))`,
`attention.py:300`). Upstream's decode gathers `[max(seq_len -
GATHER_LEN, 0), …)`
into a workspace with `GATHER_LEN` rounded up to 8 and then masks the
scores
(`:76-79`, `:151-152`); the gather is a SUPERSET and the mask makes it
exact,
because Triton needs a power-of-two tile. Walking the paged block table
directly
over the same key range reaches the identical set with no gather and no
mask.

`std::nullopt` is the ABSENT state and a NOT-TAKEN branch, not a wide
window: the
CPU decode's loop START moves, the CUDA decode's `kv_start` moves in
both split
stages, the FA-2 prefill reuses the paged launcher's own
`is_causal = causal && !is_local` normalization, and the CPU prefill's
lower
bound moves. **Both op gates prove that bit-for-bit** rather than to a
tolerance:
a window at least as wide as the sequence produces byte-identical output
to no
window, which a mask applied afterwards could not.

## The padded row needed zero `vt` changes

`Tensor::Slice(2, 0, logical)` shrinks `shape[2]` and keeps both leading
strides,
every MLA cache op sources its strides from the tensor, and that IS
upstream's
`kv_cache[..., : self.head_size]` (`_logical_cache`,
`attention.py:700-702`). The
narrowing is one line in `Dots3NoteModel::ForwardDevice`. This is the
correction
W4b-1 recorded, executed.

The evidence is the RAW cache bytes after a real forward, not an
argument: lanes
`[6, 10)` of every slot a FULL layer wrote are still zero, with the
CONTROL that
the same lanes on the sliding layers carry 28 non-zero values.

## Two of W4a's three refusals are lifted, the third is narrowed, one is
new

| refusal | now |
|---|---|
| any `sliding_attention` layer | **LIFTED** |
| a PADDED physical latent row | **LIFTED** |
| a KV cache row disagreeing with the config | **KEPT**, now compared
against the PHYSICAL row |
| `seq_len > index_topk` | **KEPT, NARROWED** — asked only of a config
that HAS a full layer, because a sliding layer sets `self.indexer =
None` (`model.py:432-434`) |
| a windowed prefill with chunked CONTEXT | **NEW**, in the seam —
upstream never merges context chunks under a window, so there is nothing
to mirror |
| a NEGATIVE `MlaBlockDims::sliding_window` | **NEW**, and it shipped
untested; the review repair gave it a case (below) |
| MoE / nextn | unchanged (W5 / W10) |

The RELEASED `dots-studio/dots3-note-prev` config still refuses, now at
layer 1's
MoE rather than layer 2's sliding attention, so nothing a user can run
changed.

## The gate

`test_dots3_note_attn` **36 cases / 3028 assertions** (30/2418 at
W4b-1),
CPU-only, no GPU, no checkpoint, no speed claim. A MIXED config
(`{full, sliding, full}`, dense MLPs, physical row 10 against the full
arm's
logical 6) is loaded through the REAL registry and run through
`ModelRegistry::Forward` TWICE against one cache pool — a 6-token
PREFILL then a
DECODE of the seventh, over a SHUFFLED block table — and compared
against a
whole-model double reference that dispatches per layer kind into W3's
`ref::Forward` and W4b-1's `sref::Forward`.

Residue 0.0254 relative; bound 6e-2. Three ratios, kept SEPARATE because
merging
them is the row's own finding F1: bound/residue **2.36x**,
nearest-mechanism/bound **2.63x**, nearest-mechanism/residue **6.22x**.
A port
with no window at all lands at 0.819 — 13.6x the bound.

**The fixture was retuned twice, and a measurement forced both.** The
first draft
(four layers, thetas 41 vs 137) measured a residue of 0.119 with the
nearest
mechanism at 0.106 — the nearest defect sat UNDER the quantisation
floor. The fix
was the fixture, not the bound: thetas orders apart the way the released
5e4 vs
8e7 are, and three layers instead of four.

`test_ops_mla_attn` **15 cases / 246290 assertions** (11 / 197113 at the
base
SHA, measured by checking those two test files out at the base,
rebuilding and
running, then restoring them byte-for-byte); `test_ops_mla_prefill`
**7 cases / 329772 assertions** (6 / 329772 before the review repair; 4
/ 242156
at the base). **That +1 case is +0 assertions on purpose, and the count
is the
statement**: the case the repair added is a CUDA one, doctest scores a
case that
returns before its first assertion as PASSED, and this box has no
device. It is a
skip wearing a pass, it is written down as one, and it must not be read
as new
coverage executed here. It WAS executed later, on a leased GPU — see the
CUDA
section below.

Their oracle is the op itself on a DIFFERENT input — a windowed paged
call
against an unwindowed call over the truncated key list — rather than a
second
copy of the same arithmetic.

## The mutation pass, through the harness this tree ships

27 rows through `scripts/mutation-harness.py`, all re-measured at the
final
36/3028 baseline, plus six more for the review repair. Highlights, with
the full
tables in spec §4.8:

- **Two reachability rows.** Deleting the DECODE call site
  (`impl.sliding_window = dims.sliding_window`) and the PREFILL one
(`…, dims.sliding_window)`) each take the gate RED, so the window comes
from
  the config through the real loader and the shared seam.
- **FIVE of the twenty-seven failed to BUILD on their first attempt**,
each on
`-Werror=unused-variable`. Nearly one row in five would have read as a
passing
test in a hand-driven pass; the compiler exit is printed beside every
row and
  each was re-run behind `((void)x, …)`.
- **M16 found a real GATE GAP.** Deleting the new
windowed-prefill-with-context
refusal left the gate green, because the case asserting it never made it
out
  of the draft. It is in the gate now with two controls, and M16b reds.
- **M19/M20 map which gate sees what.** The same leaked window is
invisible to
`test_deepseek_v2_forward` and RED on `test_mla_attention_block`.
Together
they say `impl.sliding_window` really is on the DeepSeek path and its 0
is
load-bearing — which the byte-identity table needs and cannot supply
alone,
  since identical output could also mean nothing ever read the field.

## What a fresh review changed, and what it corrected in the record

Nine findings, all discharged. Six are code and three are record; the
record half
needed the mutation numbers the code half makes measurable, which is why
it is a
second commit.

**The FA-2 MLA-prefill windowed path had no test on any device.** The
record
named both CUDA files together and then described ONE parity case, which
reads as
covering both halves. It did not: the windowed CUDA case was the DECODE
sibling
in `test_ops_mla_attn`, and `test_ops_mla_prefill`'s only `HasCuda()`
cases are
pre-existing and unwindowed. A later lease would have discharged the
decode half
against a record that read as covering both. `test_ops_mla_prefill` now
carries
`CUDA MLA prefill: the sliding window matches the CPU reference`,
comparing the
windowed device call against the windowed CPU op, against the unwindowed
device
call, and against the file's expanded single-query oracle. It
deliberately does
NOT assert the CPU case's wide-window bit-identity: a finite window sets
`is_causal = false` and dispatches FA-2's LOCAL template, so on the GPU
that
agreement is numerical rather than byte-exact. Both entries in `## Owed`
now name
the two halves separately, because a lease closes one at a time.

**A divergence this brick creates and never recorded.** Upstream's
`MLAAttention.get_kv_cache_spec` branches on the window and returns
`SlidingWindowMLASpec(..., sliding_window=self.sliding_window)`
(`vllm/model_executor/layers/attention/mla_attention.py:1215-1219` @
`bc2d63e650`), and every sliding layer sets one
(`vllm/models/dots3_note/nvidia/model.py:457`). `MakeDots3NoteKVCache`
emits one
uniform `v1::MLAAttentionSpec` for all 46. There is **no correctness
consequence** — the window is applied on READ and the gate above proves
it — but
33 of 46 layers then hold a full-length latent cache where upstream caps
a
windowed layer near the window, 513 blocks against 524288 on the
released
config. It is the single largest memory property of this architecture,
and a
token gate structurally cannot see it, which is exactly why `porting.md`
asks for
the memory format to be compared with the oracle explicitly. THREE
pieces are
missing rather than one, so it is scoped to W4b-3 rather than fixed
here. While this
branch sat in review, `SlidingWindowMLASpec` LANDED on `main`
(KV-DSV4-MULTICACHE
W1, #1960, `c1e6f3fb9`), so the merge commit's own
`include/vllm/v1/kv_cache_interface.h`
DEFINES the type this branch's copy still lists as omitted. The
divergence is
unchanged — we still emit one uniform spec for all 46 layers — but W4b-3
now
inherits the spec type and owes only the per-layer wiring, which is less
than was
owed when the paragraph above was written. Verified at merge time rather
than
assumed: `struct SlidingWindowMLASpec : SlidingWindowSpec` at
`include/vllm/v1/kv_cache_interface.h:396` on `origin/main`.

**A byte-identity claim narrowed to what was actually measured.** The
record said
§4.6's six-arm fingerprint table "is reproducible from outside the
session that
produced it". Two of six arms reproduce; four do not, and the first
draft
generalised from the first two. The diagnosis is measured on three legs
and is
not a behaviour change: exactly ONE commit touches the five shared MLA
files
between the two bases and it is W4a itself; `grep -r` finds the
fingerprints only
in the spec, so neither probe was ever committed; and `q_lora_scale`
**did not
exist** at §4.6's base (`grep -c` is 0 there and 2 at §4.8's), so §4.6's
probe
could not have set the field its own arms 2-3 are said to cover. The
byte COUNTS
agree on all six arms, so the two probes agreed on shapes and differed
on values.
The transferable rule is the reason this is in the record rather than
quietly
corrected: **a fingerprint from an uncommitted, hand-written probe is
not a
cross-session reproducible quantity.** Two probes sharing a prose label
are two
instruments, and comparing their outputs measures the authors, not the
code.
Committing the probe is what fixes it; neither scratch tree survives and
a third
hand-written probe would be a third instrument, so it goes to `## Owed`.

**A mutation's stated mechanism was wrong, and the conclusion survives
it.** The
record explained M19's green by saying `test_deepseek_v2_forward`'s CPU
forward
drives the PREFILL half while `impl.sliding_window` reaches only the
decode MQA —
which implies the prefill half WOULD be caught. Measured: leaking the
window into
the prefill call instead is EXIT 0, 11 cases / 1052 assertions, zero
failures,
SURVIVED. It misses a window leak on BOTH halves, because those CPU
cases have no
value oracle for the attention output at all — each compares the model
against
itself under another configuration, or against nothing.

**A refusal that shipped untested now has a case.**
`MlaBlockDims::sliding_window < 0` is load-bearing, because the ops read
the
window as `> 0`, so a negative value throws nowhere and silently
degrades a
windowed layer to full attention. It has controls on both sides of the
boundary
(0 is ABSENT and legal, 513 is legal) so it cannot pass on an
implementation that
refused every window. `test_mla_attention_block` 2247715 -> **2247718**
assertions, 12 cases unmoved.

**One refusal is DECLINED rather than tested, on a stronger reason than
the
review's.** The per-step `ld.head_size() <= physical_row` check is
UNREACHABLE,
not merely untested: `physical_latent_row()` IS `swa.latent_row()`, so
on a
sliding layer the comparison is an identity, and on a full layer
`ParseDots3NoteParams` has already refused the violating config at load.
R3 is
the measurement that turns that from an argument into a gate — deleting
the
CONFIG-level refusal reds `test_dots3_note_scaffold`, so the closure is
gated
rather than assumed. The review's reading, that `Tensor::Slice`
backstops it,
understates it: the backstop is not reached either. It is kept as the
executable
spelling of upstream's `assert physical_head_size >= self.head_size`,
the site
says so, and `## Owed` carries it as an untested assertion.

**A dead assignment and an uncalled accessor are deleted.**
`step.rope_cache = &rope_full` had no reader in its TU and would have
handed a
future one the full arm's rope on a sliding layer, or an empty `Tensor`
on a pure
SWA config; a null fails at the first read instead.
`MlaBlockDims::has_sliding_window()` had no caller anywhere and no call
site was
invented for it, because every consumer wants the value and
`ForwardMlaAttentionBlock` assigns it unconditionally precisely so a 0
cannot be
skipped. Two stale comments describing the pre-retune four-layer fixture
go with
them.

## The seam is byte-identical on all six arms, measured within this
session

Six arms across the seam's branch space (q_lora present/absent, both
rope
layouts, both dtypes), BASE in a separate `git archive` tree at
`925a4a587` with
a byte-identical probe (`md5sum` equal), its own configure and its own
build.
**Six for six identical, base against head.** That is the claim this
table was
built to make, and it is the claim it supports; the cross-session
comparison it
also attempted is the one narrowed above.

`test_mla_attention_block` **12 / 2247718** (12 / 2247715 before the
review
repair added the negative-window case) and `test_deepseek_v2_forward`
**11 / 1052**; `test_deepseek_v2_decode_graph_seam` 3/230,
`test_ops_mla_cache`
9/2947, `test_dots3_note_scaffold` **26 / 110819**.

NOT run and named: the SACRED DeepSeek-V2-Lite e2e token gate needs a
~29.26 GiB
checkpoint on a CUDA host; this brick ran CPU-only on a box with
neither.

## `## Nothing lands dead`

Everything this PR adds is reached from `ModelRegistry::Forward`, and
two
reachability mutations say so rather than the sentence: deleting the
decode call
site (`impl.sliding_window = dims.sliding_window`) and the prefill one
(`… , dims.sliding_window)`) each take the gate RED.

What is NOT reached, declared: **W4b-1's `dots3_note_attn.{h,cpp}`
sliding
functions** — `ForwardSlidingAttention`, `GatherSwaKv`,
`ApplySwaScoreMask`,
`BuildSlidingWindowMetadata`, `WritePaddedMlaCache`,
`NarrowLogicalCacheRows` —
still have no production call site and did not gain one here, because
the device
path reaches the same key set through the paged block table instead of
upstream's Triton gather-plus-mask. They are the gate's oracle, the
status W3's
`ForwardFullAttention` has had since W4a. Owning row
`MODEL-MM-dots3-note-dots3-note-for-causal-lm`, issue #699, listed under
`## Owed` in the spec.

## The CUDA half is compiled and EXECUTED, on sm_87

The two CUDA files this PR changes — `src/vt/cuda/cuda_mla_attn.cu` and
`src/vt/cuda/cuda_flash_attn_fa2.cu` — were written on a box with no GPU
and no
`nvcc`, and for most of this row's life the record had to say "written,
not
compiled, not run". An `rc run` lease on `orin:gpu0` (Jetson AGX Orin,
sm_87,
~36 minutes of device time) closed both halves of that. The SHA was
PROVEN rather
than asserted: the job cloned in-container and refused to build unless
`git rev-parse HEAD` equalled `53424910dfa31fbd10bcb3296a12401eaed8ee54`
with
`git status --porcelain` empty.

**Compiled, on two toolchains.** Both objects were deleted first,
gencode was
read from `compile_commands.json`, and real per-arch SASS was confirmed
with
`cuobjdump --list-elf` rather than a PTX leg.

| TU | CUDA 12.6 / sm_87 | CUDA 13.0 + the full CI arch list |
|---|---|---|
| `cuda_mla_attn.cu` | `sm_87.cubin`, 1.56 MB | rc=0, **10 cubins** —
80, 86, 87, 89, 90a, 100a, 103a, 110, 120a, 121a |
| `cuda_flash_attn_fa2.cu` | `sm_87.cubin`, 514 KB | rc=0, **6 cubins**
— 80, 86, 87, 89, 120a, 121a |

The second column reproduces what CI's `cuda-fat-build` asks, on the
toolchain it
uses. FA-2 being ON was MEASURED three ways rather than inferred from
the default
— `CUDA feature fa2: ENABLED for [87]`, `VLLM_CPP_FLASH_ATTN:BOOL=ON` in
CMakeCache, and the generated manifest `VLLM_CPP_CUDA_FA2_COMPILED_ARCHS
"87"` —
which matters because `cuda_flash_attn_fa2.cu` enters the archive only
when
`VLLM_CPP_FLASH_ATTN` (default ON), CUTLASS headers and a non-empty
`VT_FA2_ARCHS` all hold, and `cuda-fat-build` supplies the second with
`-DVLLM_CPP_CUTLASS_FETCH=ON` and the third through an arch list that
intersects
the `fa2` feature set at 80, 86, 87, 89, 120a and 121a.

**Executed — and the assertion counts are the proof.** Same binaries in
both
columns; the control is `CUDA_VISIBLE_DEVICES=""`, so the delta is the
device and
not the build.

| run | cases | assertions |
|---|---|---:|
| windowed decode alone, no device | 1 | **0** |
| windowed decode alone, on device | 1 | **49,158** |
| windowed prefill alone, no device | 1 | **0** |
| windowed prefill alone, on device | 1 | **467,010** |

`1 case / 0 assertions / SUCCESS!` is exactly the shape these two cases
have worn
until today, and the right-hand column is **the first execution either
has ever
had**. Whole-binary figures on device: `test_ops_mla_attn` 246,290 ->
2,401,528
and `test_ops_mla_prefill` 329,772 -> 2,931,678. Every filter matched
exactly one
case, so no zero-match false green; all exit codes 0 and no timeout.

**Numerically correct.** Windowed decode `MaxAbsDiff(gpu, cpu)` is
2.38e-07,
2.68e-07 and 2.68e-07 across the three split arms, against a `< 1e-3`
bar.
Windowed prefill, with 475 (query, key) pairs dropped across 57 queries:
`gpu_win`
vs `cpu_win` 0.00294137, `gpu_none` vs `cpu_none` 0.00294137 and
`gpu_win` vs
`expanded` 0.00294137, all against `< 3e-2` — and the decisive one,
**`gpu_win`
vs `gpu_none` = 1.06055** against a `> 1e-2` bar. The window
demonstrably BITES
on the device, so the FA-2 launcher is not dropping `window_size` on the
floor.
That is the exact defect the `is_local` normalization exists to prevent,
now
measured rather than argued.

**The wrinkle, stated rather than tidied away.** Three of those four
prefill
numbers are byte-identical at 0.00294137. The likely reason is that the
global
max-error element sits on an UNWINDOWED row, where `gpu_win`/`gpu_none`
and
`cpu_win`/`cpu_none`/`expanded` each coincide and all three become the
same
subtraction — but that element was not instrumented, so this is
inference and not
a measurement. It is not degeneracy: if the CPU arms agreed wholesale
the second
assertion would read about 1.06, not 0.0029.

**Two limits.** Execution is proven on **sm_87 only**; the ten-arch
result is
COMPILE-ONLY, because CUDA 13 cannot run against that box's NVRM 540.4.0
driver
(`cudaGetDeviceCount rc=35`). And this is KERNEL-level parity on two
ops. It is
not the e2e model gate, which is unrelated and still owed.

## The two red CI lanes are pre-existing and unrelated, with the
evidence to re-derive that

Neither red is attributable to this change, and both are named here
rather than
left for a reader to assume.

- **`build-test-cpu`** fails at the TEST step, not the build:
**627 of 628 passed**, and the single failure is `test_cpu_threadpool`
at
  `tests/vt/test_cpu_threadpool.cpp:539`, `CHECK(ratio < 100.0)` reading
  `2 threads 0.48 us, 5 threads 52.638 us, ratio 109.662`. It is a fixed
wall-clock ratio guard whose denominator is the small-thread leg, so it
fires
on GitHub's 4-vCPU runner while needing roughly a 37x blowup to catch a
real
  regression on a large box. Issue
[#631](#631) owns it and it is
deliberately unfixed, because raising a threshold to green a gate is a
scope
change that needs its own spec. It touches no MLA, attention or model
path.
- **`windows-msvc-cpu` and `windows-msvc-vulkan`** are the baseline-less
lane,
  issues [#503](#503) and
[#584](#584). Controlled rather
than
asserted: they fail identically on open PRs #1967, #1956 and #1935, and
ZERO
of the last 12 windows runs produced a `main` verdict, so there is no
green
  baseline for this branch to have regressed from.

## The merge that built clean and threw

Integrating `origin/main` a second time merged clean, COMPILED clean,
and took
`test_dots3_note_attn` from 36/3028 green to five cases throwing
`resident weight: EMPTY tensor has no host bytes to alias`. The incoming
change
is #1952's review finding #1953: `ResidentWeight` now refuses an empty
weight,
because an empty one aliases a null host pointer no downstream op can
detect.
The refusal is right and this branch was on the wrong side of it — it
made BOTH
rope caches resident up front while the materializer deliberately leaves
the
unused one empty (each is 64 MiB at 524288 positions). One guard per
cache fixes
it, and the comment now says the guard is the CONTRACT rather than an
optimization.

The two branches touched different files, so nothing about this was
visible to
the merge: it is "merge-tree CLEAN is not merge-tree BUILDS" one step
further
along, where it builds too. Re-running the focused gate set AFTER the
merge and
BEFORE the push is the only thing that caught it. Spec §4.8 records it.

## What is owed, and the split I am proposing

The row is NOT finished, and the sm_87 result above does not narrow this
list.

**W4b-3**, and the line is real: the DSA lightning indexer's SELECTION
shares
nothing with the sliding window. The sliding layers have no indexer at
all; the
indexer is on the FULL layers, and lifting `seq_len > index_topk` needs
indexer
weights on device, a logits kernel, a top-k and a SPARSE MLA attention
kernel on
both backends. Bundling it here would have doubled a PR that already
carries two
new kernel bounds and a seam extension. W4b-3 carries it, together with
the
windowed-prefill-with-context refusal and the `SlidingWindowMLASpec`
divergence
above.

**`.agents/specs/dots3-note.md` is STALE against this body at the merge
commit,
in three places, and W4b-3 owns every one.** The spec is the row's
authoritative
record, so a merge that leaves it disagreeing with its own commit
message is a
debt and is written down as one. It is not fixed on this branch because
a push
now would cancel a `cuda-fat-build` this change has already had
destroyed twice,
and the fresh reviewer sanctioned folding record-only corrections into
W4b-3:

- §4.8 reads `test_mla_attention_block` **12 / 2247715** and calls it
"unmoved
  from the numbers §4.6 recorded". The measured value at this SHA is
**12 / 2247718** — the implementer, a fresh reviewer and the operator
each
measured it independently and agree. `2247718` appears nowhere in the
spec.
  This body says 2247718 and the spec contradicts it.
- §4.8's `#### The #1969 REVIEW-REPAIR rows` section says "Five more
rows" and
its table lists **six**: R1, R1-control, R2, R3, R4 and R4b. This body
says
  six.
- **The most misleading of the three.** §4.8 and `## Owed` still say the
CUDA
lease "cannot currently be taken" and name `thor:gpu0` as the host that
discharges it. Both statements are now false. The compile and the
execution
above were discharged on `orin:gpu0`, and `thor` could never have
discharged
the PREFILL half at any point, because `VT_CUDA_FEATURE_TABLE`'s `fa2`
row is
`8.0,8.6,8.7,8.9,12.0a,12.1a` — sm_87 carries FA-2 and thor's sm_110
does not.
Left alone, the record claims blocked work that has in fact been done,
on a
  host that could not have done it.

**Two comment-level corrections also go to W4b-3, named here so the miss
is in
the history rather than only in a review thread.** They are not fixed on
this
branch for the same reason, and neither is worth a gate run of its own:

- `src/vllm/model_executor/models/dots3_note_device.cpp`'s upstream
anchors for
the `deepseek_v2.py` diff are off by three. The lines are **460** and
**498**
at `bc2d63e650`, not 457 and 495 — 457 is `qk_nope_head_dim: int,` and
495 is
a bare `)`. The substance of that correction is right; only the two
numbers
  are wrong.
- The new CUDA prefill case's comment says the CPU fixture is
"unchanged". It is
not literally one experiment: the CUDA case uses `h = kHeadsLite` (16)
where
  the CPU case uses `h = 4`. Lengths, window and seeds do match.

**Still open, and blocked on things this row cannot buy.** The
end-to-end parity
gate against vLLM — token-exact or the ratified near-tie form, chosen by
measurement — plus every throughput, latency and memory axis, and the
quantized
arms. No host this project owns runs the `dots3_note` oracle at any
published
precision, and the row is beyond the parity pin. `## Owed` in the spec
carries
each of these with its owner and #699.

Row stays SPIKE. Under §6.4 option B the CPU gate remains a consistency
gate
against an independent reference, not a correctness gate against an
oracle: no
vLLM instance for `dots3_note` runs on any host this project owns.

FOLLOWING_AGENTS_PROTOCOL

Following-Agents-Protocol: true
AI-Assisted: true
Assisted-by: AGENT:claude-opus-5 [Claude Code]

---------

Co-authored-by: Ettore Di Giacinto <mudler@localai.io>
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