GLM-5.3 Flash performance campaign: 3.21x FP8, 4.96x MXFP4 warm short decode, plus MXFP4 release support - #42
Merged
Merged
Conversation
Adds the observability M0 needs to instantiate the campaign cost model, and repairs the pre-existing race that the first protected prefill arm exposed. Observability, opt-in and off by default: - STRATA_GLM53_PHASE_PROFILE=1 emits structured [glm53-phase-profile] JSON covering per-phase graph time, attention/KDA/MLA/feed-forward, host-expert subphases, useful expert weight bytes, cache, CUDA transfer/call/sync, allocation, RSS, VRAM and token IDs. Default execution enables no CUDA event timing. - STRATA_GLM53_PREFILL_PAGE_TOKENS controls scheduler prefill width for the M0 reuse sweep. The production default remains 64. Every counter update is behind `if (config.phase_profile)`. The only work left on the disabled path is the steady_clock reads that bracket each phase: 348 per decode step at a measured 18.665 ns, or 0.00072% of a 903 ms step. Profiled and unprofiled arms produce identical token IDs. KDA recurrent-state repair: Head workers each called sequence.recurrent(layer) for the first time concurrently, racing that span's lazy copy-on-write allocation and corrupting the heap; the observable was exit 139 on a wide prefill page. The span is now resolved and shape-checked once on the scheduler thread and workers take disjoint head subspans of it. This changes neither arithmetic nor operation order, and it strictly removes per-head work from the always-on path. The two changes ship together because they are one unit of work: the paging control is what exposed the race, and every arm in experiment 0196 requires the repair to run at all. Gates: make check passes; the custom suite is 331/369 passed, 38 skipped, 0 failed, unchanged from before this work. Full evidence, including a discovered pre-existing defect where prefill page width changes generated output, is in experiments/docs/experiments/0196. Co-Authored-By: Claude Opus 5 <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_014jtzY5AZ3KBAjBA5BHgPzH
Campaign contract section 8 requires a load report carrying cold load, warm
load, first token, bytes read, faults, RSS, VRAM and any declared
I/O-dependent tier. This adds it, model-neutral, because nothing in it is
GLM-specific.
Layering: the /proc and /sys readers go in strata_platform, which already
reads sysfs for HardwareProfile and cannot depend upward; assembling a report
needs the checkpoint's backing device from resolve_backing_storage, which
lives in strata_engine, so the report itself goes there. RuntimeSession wires
it because load and first token are model-neutral boundaries it already owns.
The report carries one field the contract does not name and experiment 0197
showed was the missing measurement: average storage queue depth per phase,
from the kernel's weighted in-flight time -- the figure iostat calls aqu-sz.
0197 could not distinguish "the runtime issues one read at a time" from "the
device cannot go faster" from outside the process, and rejected three
candidate mechanisms without finding the real one.
Its first run settles that question and reframes it:
load 1.14 s read 0.00 GB 0 major faults rss 0.45 GB
device queue depth 0.01, busy 0.0%, 0 reads
first-token 71.88 s read 18.90 GB 1,340,638 major faults rss 136.93 GB
device queue depth 4.79, busy 37.9%, 1,514,102 reads of 12.5 KiB
Queue depth 4.79 means the runtime is not serialized, so adding threads or
issuing reads from more places cannot help. What it is doing is assembling
18.90 GB out of 12.5 KiB requests while the device sits 62% idle. Every probe
in 0197 read contiguous 64 KiB to 4 MiB chunks, which is why none reproduced
the production rate.
The report also measures, in-process, the reframing 0197 argued from arms:
`load` is 1.14 s and reads nothing, while `first-token` carries every byte and
every fault. On a lazily mapped checkpoint there is no load time to optimize.
Marks are taken at boundaries only -- two /proc reads and two sysfs reads per
phase -- so nothing samples inside a decode step. Emission is gated on the
existing report_placement_plan switch.
Gates: make check passes; suite 338/376 passed, 38 skipped, 0 failed, with
seven new tests covering queue-depth derivation, the serialized-reader shape,
absent and non-monotonic counters, zero-length intervals, and render before
any phase is marked.
Co-Authored-By: Claude Opus 5 <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_014jtzY5AZ3KBAjBA5BHgPzH
Closes three of the four discovery gaps M1's audit found. Campaign contract section 8 requires discovering huge-page availability, PCIe locality and the peer/P2P/NVLink matrix; none of the three was recorded anywhere. Huge pages, into HardwareProfile: page size, explicit hugetlbfs total and free, and whether transparent huge pages are enabled at all and in which mode. The reference host reports 2 MiB pages, zero explicit pages reserved, and THP in `madvise` rather than `always` -- so a long-lived arena would have to ask with MADV_HUGEPAGE, and nothing does. That is now a visible fact. PCIe locality and peer ranking, into the placement plan. CudaBackend already exposed device_numa_node and high_speed_peer_access_supported and both were queried at runtime, but nothing persisted them, so a cached plan could not be verified against the topology it was built for. The plan now carries per-device NUMA node and a row-major high-speed peer matrix, and the report prints them: devices cuda:1 (23.30 GiB free of 23.56 GiB, numa 1), cuda:2 (..., numa 1) peer no high-speed peer link between admitted devices Both match what M0 recorded by hand in the frozen manifest, which is the point: discovered rather than asserted. Schema is bumped to v3 and v2 plans are discarded rather than upgraded. A v2 plan carries no peer matrix, and reading that as "no high-speed link anywhere" would let a cached plan refuse a fast path the machine actually has -- the failure being avoided is silent and one-directional, so the version gate is the conservative choice. Two encoding details worth naming. numa_node is -1 when the driver declines to say and JsonCursor parses unsigned only, so it is stored plus one rather than widening the JSON layer for a single field. The peer matrix is uint8_t, which streams as a character, so it is written through an explicit unsigned cast; a test covers the round trip of both, including the -1 sentinel and the clear diagonal. Still open, deliberately: the pinned host-memory budget. No cudaHostAlloc or cudaHostRegister exists in the runtime, so every transfer is pageable. It is immaterial at short context but M0 measured 295.77 GB of D2H at long context on the MLA fallback path, and the transfer path it would serve is M4's to change. Gates: make check passes; suite 340/378 passed, 38 skipped, 0 failed. Co-Authored-By: Claude Opus 5 <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_014jtzY5AZ3KBAjBA5BHgPzH
…ol width Three M1 gate items, all small. Admission now refuses a plan whose peer topology differs from the probed machine. Capacity mismatches already failed loudly; a peer mismatch would not have. A plan cached on a peer-linked host and replayed on one without the link would admit a fast path the hardware cannot serve, and the run would simply be slower and wrongly attributed rather than rejected. Placement determinism is now gated by a test rather than assumed: two solves of identical checkpoint, topology and request encode byte for byte, and a plan survives an encode/decode cycle unchanged. Comparing encodings covers the memory accounting as well as the layer assignment, since any field that drifted between two solves would change the encoding. The DeepSeek host MoE pool takes its width from HardwareProfile instead of a hard-coded 48, which was that executor's development box's logical CPU count asserted as a constant -- the class of fact hardware_profile.hpp exists to stop keeping in the type system, and a silent mis-size on every other machine. The policy stays in the executor (one worker per usable CPU across every node, because a both-shards pool spans them); only the measurement moves. This is safe rather than merely plausible: the sole production construction, in deepseek_rank_local_layer_executor.cu, passes both_shards=false with an explicit CPU list, so the constant was never reached on any live path. The accompanying test pins both shapes. Gates: make check passes; suite 343/381 passed, 38 skipped, 0 failed. Decode and prefill were measured back to back against a rebuilt f5b203e under identical host conditions -- decode -1.09%, prefill +0.57%, output byte-identical -- which is M1's 2% clause. See experiment 0199, including why comparing against a recorded figure could not have shown that: MLA wall time swings 2.37x with host state on unchanged code. Co-Authored-By: Claude Opus 5 <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_014jtzY5AZ3KBAjBA5BHgPzH
… figure The load report printed 2.15 GB of VRAM per device where the runtime actually used 19.65 GB, a factor of nine. It had been handed plan.device_resident_bytes, which is the workspace the plan reserves up front, not the device expert cache the runtime demand-fills. Found by M1's own "no unaccounted full-size duplicate representation" audit, which is the one place that error would do real damage: a report understating VRAM by 9x invites exactly the conclusion the audit exists to test. It now samples the driver at every phase mark. Devices are set before the first mark so load and first-token both carry a figure, and the difference between them is informative in itself: load 1.19 s vram 19.41 GB in use, 19.41 GB in use first-token 39.14 s vram 19.62 GB in use, 19.63 GB in use The device is already at 19.41 GB when load returns, so the expert cache is filled by the background warm-up thread rather than during the first forward -- an M1 loading fact the report shows only because the figure is measured rather than planned. An unreadable device contributes zero rather than aborting the report; the storage, fault and RSS figures are still worth having. Gates: make check passes; suite 343/381 passed, 38 skipped, 0 failed. Co-Authored-By: Claude Opus 5 <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_014jtzY5AZ3KBAjBA5BHgPzH
…rvice M2's first mechanism, and it is a scheduling change rather than a layout one. HostWorkerPool::parallel_for hands out consecutive indices from a single counter, and the GLM host expert loops index tasks as (expert * intermediate + row). Twenty-eight runners therefore walked one contiguous FP8 matrix with a 28-row stride each, so the hardware saw 28 interleaved strided streams where the source reads as sequential. Experiment 0198 measured that shape at 0.69 GB/s against 1.96 GB/s for contiguous blocks on identical cold bytes; this is the warm answer. parallel_for_blocked claims `block` consecutive indices per fetch_add. parallel_for now delegates to it with block 1, so no other caller changes. The four GLM expert loops use 64: 256 KiB of contiguous FP8 weight per claim, wide enough for the prefetcher, narrow enough that 28 runners still balance across a 2048-row expert at 32 claims each. Measured on protected short warm decode, one binary, three arms back to back, bracketed by repeats so host drift is bounded rather than assumed: block 1 96.75 s decode 72.70 s expert service 16.75 GB/s block 64 78.76 s decode 50.03 s expert service 24.34 GB/s block 1 97.39 s decode 72.94 s expert service 16.70 GB/s Host expert service -31.2%, useful bandwidth 1.45x, decode 1.323 -> 1.625 tok/s, prefill 34.60 -> 13.53 s. Host drift between the bracketing arms is +0.32%, so the effect is about a hundred times the drift. The down projection carries it at -56.1%: it reads 4,096 rows of 2,048 bytes, so a 28-way stride breaks a shorter run than gate/up's 4,096-byte rows. Exactness is measured, not argued. Each task writes a disjoint output element, so reordering changes no arithmetic; useful expert bytes are identical to the byte across arms (1,217,920,499,712) and output matches M0's canonical hash. STRATA_GLM53_EXPERT_DISPATCH_BLOCK exists so both arms of the A/B run the same binary; 1 reproduces the previous dispatch. Production default is 64. Against M2's >=40 GB/s gate this reaches 24.34 from 16.75. It does not clear the gate alone and was not expected to; 30,208 timed allocations per 128 decode tokens and timed view resolution remain, both already counted by the profiler. Gates: make check passes; suite 346/384 passed, 38 skipped, 0 failed, with three new tests covering coverage, contiguity within a claim, and the zero-block guard. See experiment 0200. Co-Authored-By: Claude Opus 5 <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_014jtzY5AZ3KBAjBA5BHgPzH
M2 requires timed allocations and checkpoint view resolution to be zero. M0 counted both every step: 30,208 timed temporary expert allocations per 128 decode tokens, and view resolution at 0.35-0.72 s. Expert views are now resolved once into a process-lifetime cache indexed by (layer, expert slot), slot 288 being the shared expert. Every MoE call had been rebuilding three module-name strings per expert and performing three manifest lookups -- 27 per call, 5,376 calls per 128 decode tokens. Glm53CheckpointReader::view computes a span into the mapping and touches no pages (read is the one that moves bytes), so caching a view changes nothing about what is read, only how often the manifest is consulted. The cache is mutex-guarded rather than lock-free and the lock is taken once per call around the whole group: the call site is the scheduler thread today, but nothing in the type system says so. The MoE scratch buffers are grown once and reused. Every element is overwritten each call -- gate/up writes every activation slot, down every output slot -- so reuse without clearing is exact, and the byte-identical output gate verifies that rather than the argument doing so. The allocation counter now reports what happened instead of a constant, so the profiler shows the term reaching zero. The first attempt used function-local `static thread_local` scratch and segfaulted: a thread_local resolves to the *worker's* copy inside the dispatch lambda, not the caller's, so all 28 workers indexed into their own empty vectors. The buffers are filled by the worker pool, so they are Impl members, for the same reason host_moe_workers is. Measured three arms back to back against 140781e, bracketed for drift: prev 49.11 s expert service 24.80 GB/s 30,208 allocs 440.6 ms view curr 48.19 s expert service 25.27 GB/s 3 allocs 130.2 ms view prev 49.60 s expert service 24.56 GB/s 30,208 allocs 518.3 ms view Timed allocations 30,208 -> 3, and zero in prefill. View resolution 3.4x cheaper. Expert service improves 1.88% against 0.98% host drift, which is about twice the drift and is reported as exactly that -- a small real effect near the noise floor, not a throughput win. The removed terms were 0.44 s of view time and some allocator work inside a 49 s step, so a few percent was the arithmetic ceiling before the arms ran. Prefill shows no detectable change. Output is byte-identical across all three arms. Useful expert bandwidth stands at 25.27 GB/s against M2's >=40 GB/s gate. The remaining mechanism is the one the gate was written for: a transformed, tiled, NUMA-local expert representation. Gates: make check passes; suite 346/384 passed, 38 skipped, 0 failed. See experiment 0201. Co-Authored-By: Claude Opus 5 <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_014jtzY5AZ3KBAjBA5BHgPzH
M3's first lever: run the ninth expert of every MoE layer -- the shared one, the only one the router does not choose -- on the GPU that owns the layer, while the host works through the other eight. M0 attributed 67.7% of a decode step to host expert service with both GPUs idle for 94% of it, so this screened at 7.5% of the step. The obstacle was exactness, not capacity. The existing FP8 MoE kernels sum with a block reduction while glm53_host_fp8_dot_avx2 sums into eight __m256 accumulators and combines them with a fixed tree, and floating- point addition is not associative. glm53_shared_expert_dot_kernel replicates the association: one thread per output row holding an 8x8 register block that is the AVX2 register file, the same fma(mul(decoded, scale), activation, acc), and the same combine tree. Every rounding and the SwiGLU stay on the host, which keeps expf -- where glibc and CUDA differ by up to 2 ULP -- out of the comparison entirely. It is bit-exact: zero mismatches over 8,192 output rows at sm_86 and at the production compute_75 JIT, and 129-token decode is byte-identical across five alternating arms at the hash M2 has produced since 0200. It does not pay, so it ships off. Removing 8.20% of the expert bytes the host services took 3.7% off the gate/up dispatch and 2.5% off down, then gave 1.0 s back waiting on the device collect: net -0.84% on host expert service against host arms whose own spread was 0.48%, and device arms with ten times that variance. At full coverage that is about 1.4% of a decode step, for 1.06 GB of the VRAM M4's routed hot tier wants. The mechanism is retained default-off because M4's expert-residency work needs exactly these primitives -- an exact device FP8 expert dot, a checkpoint-native resident tier, and an arena reservation taken before the weight cache is built with the capacity. Two measurements that outlive this lever: - Warm decode reads nothing from storage. Sampling /proc/<pid>/io across a full run records zero bytes and zero major faults during the decode window. The load report's "I/O dependent" verdict is a capacity projection, not a property of this operating point. - STRATA_GLM53_EXPERT_BOUND prices an expert directly. Across the resolved range, removing 41.0% of host expert bytes bought 26.3% of host expert service and 57.4% bought 40.7% -- roughly 0.7% of service per 1% of bytes. That is the number M4's hot tier has to clear, and it is now measured rather than assumed. These arms produce incorrect text by construction and are never candidates. Also fixes an inference that would have been wrong: [glm53-shared-expert] reports which layers took which path, and it showed 11 of the 42 MoE layers reaching the MoE through host_moe_page rather than host_moe even at a single decode row -- the MLA fallback layers. Records 0210, 0211. Co-Authored-By: Claude Opus 5 <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_014jtzY5AZ3KBAjBA5BHgPzH
Section 11's first gate item needs sequential decode routes with reuse distance and byte-weighted coverage before any residency policy is designed, and there was no way to get them out of the runtime. STRATA_GLM53_ROUTE_CENSUS=<path> writes one row per routed MoE layer visit: phase, request, position, layer, and the eight experts the router chose in order. 42 layers times 8 experts is 336 uint16 per decode token, so a 128-token census is 86 KB and can be held whole and written once rather than streamed. Recorded before the prefetch guard in observe_route, because prefetch is a separate opt-in that is off whenever the host expert path owns the experts -- which is exactly the configuration whose routes M4 needs. Co-Authored-By: Claude Opus 5 <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_014jtzY5AZ3KBAjBA5BHgPzH
The kernel was exact but slow for a structural reason: one thread per output row streams a whole 4,096-byte weight row alone, which is entirely uncoalesced. It measured 9.8 GB/s on gate and up against a 3090's 936, and 2.149 ms per expert against the roughly 1.04 ms the 28-core host pool takes -- the GPU was twice as slow per expert as the CPU it was meant to relieve. The fix costs nothing in exactness because the association already had the right shape. The host sums into eight __m256 accumulators; assigning one thread per accumulator *group* rather than per row means thread g owns accumulators[g] and all eight of its lanes, and reads block[g*8 + lane] -- so the eight threads of a row read the 64 bytes of a block contiguously, as one uint2 and two float4 each. The width-4/2/1 combine tree becomes eight __shfl_down_sync per step, and thread 0 ends holding exactly the accumulators[0] the intrinsics would have held. 95.6 GB/s on gate and up, 159.0 on down: 9.8x and 8.1x, 0.229 ms per expert. Zero mismatches against glm53_host_fp8_dot_avx2 over all 8,192 output rows, and 129-token decode is still byte-identical at b3deffc5d0f0. This matters for M4 rather than for the shared expert. The route census (0212) puts a VRAM-feasible static tier at 28-37% of routed bytes, and the projection assumed the GPU could serve its share inside the host's window. At 2.149 ms per expert the device was 90% of the critical path and the projection was about to collapse; at 0.229 ms it is 9%, and the constraint moves back to VRAM where the census can reason about it. Co-Authored-By: Claude Opus 5 <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_014jtzY5AZ3KBAjBA5BHgPzH
Correction to the verdict recorded with the kernel two commits back. The lever was measured at -0.84% and held off; that measurement was taken against a device dot running one thread per output row, at 2.149 ms per expert against the host pool's ~1.04 ms. The GPU was slower per expert than the CPU it was relieving, so the collect wait ate the dispatch saving. Coalescing the kernel (2193ee3) changed the input to that decision, not just its margin. Re-measured with the same five-arm alternating protocol: both device arms beat the host arms bracketing them, by 4.49% and 6.33% of host expert service. Mean -4.93% on service, -2.79% on decode wall. gate/up -7.16% and down -5.98%, which is what removing one expert of nine predicts from makespan alone -- 288 blocks over 28 workers falling from eleven rounds to ten is -6.7%. The two device arms agree to 1.6% while this session's host arms span 7.06%, so the reading rests on the bracketing comparison rather than on means. Output is byte-identical at b3deffc5d0f0 across all five arms, the hash this campaign has produced since 0200. The tier costs 1.06 GB, about one expert per layer of the K=11 that the route census measures as VRAM-feasible, and returns several times that. Automatic admission declines silently when the capacity is not there; an explicit opt-in that cannot fit is an error. Still open: 11 of the 42 MoE layers reach the MoE through host_moe_page at a single decode row and do not take the device path. Covering them should take -2.79% to roughly -3.8%. Records 0210, 0211, 0212. Co-Authored-By: Claude Opus 5 <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_014jtzY5AZ3KBAjBA5BHgPzH
… defect Continues 0214 from where the previous session stopped. Its closure engineering is done and its exactness ladder now has a verdict. Workspace reservation. The rewritten resident path sized its workspace from `history * expanded_width`, which grows by 16,384 floats per token, so it would cudaFree and cudaMalloc inside a timed decode step -- which 0214's gate forbids outright and which would also make a step's cost depend on when the buffer last grew. It now reserves max(history, maximum_context), 134 MB at 2,048 tokens on this model, inside the workspace reserve the planner already withholds. Offsets use the reserved extent so a workspace pointer means the same thing at every position. Two real arithmetic mismatches against the host fallback, both fixed: the score dot used a separately rounded __fmul_rn/__fadd_rn where the host writes `score += q * kv` and GCC contracts it to an FMA on this FMA3 host; and the scale used rsqrtf, a 2 ULP approximation, where the host computes 1.0f / sqrtf once. A second oracle on the latent cache itself, comparing the device state buffer's row against the host cache's after the fallback appends it. Verdict: the candidate fails and is not timed. Two-token smoke and the 4-token layer-boundary discriminator pass; 129-token repeated decode state evolution diverges at token ~35. Localized to layer 15, position 45, one BF16 ULP. The latent oracle never fires, so the carried state is bit-identical and the defect is in the published branch. Neither arithmetic fix moved the failing values at all, which rules out the score path. The cause is expf. CUDA's differs from glibc's in 30.38% of f32 results and 0.000375% after BF16 rounding; GLM-5.3 makes 64 heads x history x 11 layers exponential calls per token, so at history 46 a token has an 11.4% chance of a visible divergence and the nine steps from position 36 to 45 survive with probability 0.34. That predicts the observed failure and explains why a four-token oracle passes. This is not a kernel defect. Two independently implemented exponentials disagree in the last bits and the exact-mode contract makes the host's the definition. The fix follows the accepted shared-expert precedent: the device does the linear algebra, the host does the transcendental. Return the 64 x history scores for exp and normalization -- 12 KB per layer at history 46 against the 13.4 MB per layer per token the chain removes. Record 0214. Co-Authored-By: Claude Opus 5 <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_014jtzY5AZ3KBAjBA5BHgPzH
Confirms 0214's expf diagnosis by fixing it. CUDA's expf and glibc's disagree in 30.38% of f32 results and 0.000375% after the BF16 rounding the model applies to the coefficient; at 64 heads x history x 11 layers that is an 11.4% chance of a visible divergence per token at history 46. glm53_mla_decode_to_mhc now stops at the raw scores and returns heads * history of them. The runtime applies the maximum, std::exp, the sum and the BF16 coefficient rounding using the accepted fallback's own arithmetic. glm53_mla_decode_finish uploads the coefficients and completes the value-weighted sum and the output projection. The fused attention kernel is split into glm53_mla_scores_kernel and glm53_mla_weighted_kernel, both keeping the host's sequential association and FMA contraction, and the workspace carries heads * reserved_history extra floats for the exchange. The round trip is 12 KB per layer at history 46 and 512 KB at the 2,048-token ceiling, against the 13.4 MB per layer per token of activation D2H the resident chain exists to remove. The 48-token layer-boundary oracle, which previously died at layer 15 position 45, now completes with no mismatch at any MLA layer or position. The 129-token gate still fails, and the two facts together locate what is left: the oracle certifies every MLA layer's branch through 48 tokens, yet the text diverges at byte 184, inside that window. So the remaining divergence is not in the MLA attention. It is in the rest of the layer. forward_layer_resident is admitted by (glm53_kda_layer(layer) || resident_mla_enabled()), so enabling resident MLA moves those 11 layers' mHC pre, input layer norm and mHC post onto the device too -- the control runs them on the host for exactly those layers. The oracle downloads `normalized`, which is the output of the device mHC and norm, and feeds it to the host attention, so it is blind by construction to a difference in what produced it. Candidate remains rolled back and untimed. Next is a direct host/device comparison of mHC pre, the input norm and mHC post on those 11 layers. Also corrects the reserved workspace figure: 268 MB at 2,048 tokens, not 134 -- expanded_width is 32,768 because it holds key and value together. Record 0214. Co-Authored-By: Claude Opus 5 <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_014jtzY5AZ3KBAjBA5BHgPzH
0214's attention oracle downloads the layer input that the device mHC and norm produced and feeds it to the host attention, so it compares attention given identical input and is blind to whatever produced that input. This adds the comparison it could not make: download the same layer input and check it against a host mhc_pre plus norm computed from the same streams and the same two tensors the device was given. It runs on every resident layer, not only the MLA ones, because the KDA layers take the device path in both arms and a difference there would mean the control is itself a mixture. It is. The oracle fires at layer 0 -- a KDA layer -- at the first decode position: host -0.003372 against device -0.004150, 7.8e-4 absolute. The resident MLA work did not introduce this and nothing in this campaign had compared the two implementations before. A plausible mechanism is precision: the host mhc_pre accumulates the stream RMS and all 24 projection rows in double before a 20-iteration Sinkhorn split with a 1e-6 tolerance, and an f32 device projection would feed slightly different values into an algorithm that amplifies them. Recorded as an unresolved discrepancy rather than a device defect: it is not ruled out that this host reconstruction is incomplete. If it holds, M4's MLA exactness gate cannot be met as written. The control runs device mHC on 34 layers and host mHC on 11, and the campaign's reference hash encodes that mixture, so a candidate that moves 11 layers between the two implementations cannot be byte-identical to it however exact its attention is -- which is exactly what today measured: the MLA branch is provably exact at every layer and position through 48 tokens and the text still diverges. Record 0214. Co-Authored-By: Claude Opus 5 <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_014jtzY5AZ3KBAjBA5BHgPzH
Settles the mHC question and fixes the 129-token gate.
The two mHC implementations are different functions. The survey oracle
disagrees on essentially every element of every layer, KDA and MLA
alike -- 4096/4096 at layers 0 and 4, worst 0.16 absolute. Confirmed
without the oracle: STRATA_GLM53_FUSED_KDA=0, which puts all 45 layers
on the host path, yields bf2016cd3e39 against the production
b3deffc5d0f0, diverging at byte 391, with three repeats confirming the
production configuration is deterministic.
So the campaign's reference hash encodes a mixture: device mHC on the 34
KDA layers, host mHC on the 11 MLA fallback layers. Host mhc_pre
accumulates the stream RMS and all 24 projection rows in double before a
20-iteration Sinkhorn split; the device does the same work in f32 with
split partial reductions, and since collapsed is linear in the Sinkhorn's
pre weights a precision difference becomes a percent-level one.
That is why the gate could never pass. Enabling resident MLA moves 11
layers between the two implementations, so its output must change however
exact its attention is -- which is exactly why the layer oracle could
certify every MLA branch through 48 tokens while the text diverged.
dsv4_mhc_publish_branch and STRATA_GLM53_RESIDENT_MLA_HOST_ATTENTION give
the control the candidate actually needs: device mHC on all 45 layers in
both arms, differing only in where the MLA attention runs.
The gate passes. Three control and three candidate observations,
alternating, 129 tokens, all six arms one hash fe74dc4ec7ab:
decode wall 68.75 -> 52.26 s -24.0% (control spread 3.3%,
candidate within 1.2%)
MLA 18.60 -> 3.93 s -78.9%
activation D2H 18.99 -> 0.32 GB -98.3%
synchronizations 14,208 -> 7,168 -49.5%
synchronization s 5.59 -> 0.22 -96.0%
This is the largest measured win in the campaign and it is exact.
It does not ship yet. fe74dc4ec7ab is a third configuration -- device mHC
on all 45 layers -- so landing it changes production output, not because
the attention is inexact but because it makes the model compute mHC one
way instead of two. Nothing in the architecture justifies the split and
the campaign had never compared them, so this is arguably a latent defect
being repaired; but it re-baselines every figure anchored to
b3deffc5d0f0. Recorded as an owner decision in the contract's blockers.
Record 0214.
Co-Authored-By: Claude Opus 5 <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_014jtzY5AZ3KBAjBA5BHgPzH
Owner ruled that the model computing mHC two different ways by layer type is a latent defect being repaired, not a regression introduced, and authorized the re-baseline. Resident MLA is now on by default and all 45 layers use one mHC. Exactness rests on the isolated gate from the previous commit: with the mHC held fixed on both arms and only the attention varying, three control and three candidate observations are byte-identical. Shipped before/after, one protected session, new/old/new: new fe74dc4ec7ab 52.37 s old b3deffc5d0f0 77.40 s new fe74dc4ec7ab 52.57 s Decode 77.40 -> 52.47 s median, -32.2%, 1.65 -> 2.44 tok/s, 1.48x. That is larger than the isolated -24.0% because the 11 MLA layers also stop paying host mHC. The campaign reference hash moves from b3deffc5d0f0 to fe74dc4ec7ab. Every figure anchored to the former describes the pre-repair mixture. STRATA_GLM53_RESIDENT_MLA=0 restores the old path. STRATA_GLM53_RESIDENT_MLA_HOST_ATTENTION=1 keeps the new mHC while running the attention on the host, which is the control that closed the gate and is retained for future exactness work. Record 0214. Co-Authored-By: Claude Opus 5 <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_014jtzY5AZ3KBAjBA5BHgPzH
This reverts commit dba3dca.
GLM-5.3-Flash is published in two releases Strata can execute: the pinned FP8 E4M3 block-128 checkpoint the campaign targets, and a quark MXFP4 export whose routed experts are E2M1 nibble pairs with one E8M0 scale byte per 32 columns. Only the routed experts are quantized in either; the MXFP4 release leaves its whole spine, its shared experts, and the routed experts of layers 3, 5 and 6 in BF16 under the publisher's mixed-precision correction, where the FP8 release stores that spine E4M3. The release is resolved once, at open, from pinned index extents, and every layer below branches on the manifest rather than on a build flag or a flag the caller has to remember. The host MoE gains an MXFP4 decoder and a BF16 decoder alongside the existing FP8 one, all three sharing the eight-chain accumulator structure so the formats differ only in how a weight is decoded. Sixty-four columns span exactly two group-32 scales, which is what keeps that structure intact for MXFP4. The FP8 path is untouched: same dot, same activation quantization, same device shared-expert tier. Both of those are FP8-specific and are now gated on the release, because MXFP4 is a weights-only format with no device counterpart and quantizing its activations to E4M3 would discard precision the format never asked to lose. An explicit shared-expert opt-in against an MXFP4 checkpoint says so rather than silently running the host path. The nibble order is measured, not assumed: dequantizing layer 20 expert 0 from both releases and correlating gives 0.993 for column 2b in the low nibble of byte b and 0.000002 for the swap. Co-Authored-By: Claude Opus 5 <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_014jtzY5AZ3KBAjBA5BHgPzH
Both host expert dots run at 2.7-2.8 GB/s at the production gate/up shape, which says the row dot is uop-bound, not byte-bound: halving the checkpoint buys nothing by itself, and what separates the formats is the instruction sequence that turns a stored code into a float. `vgatherdps` is microcoded on this Broadwell host. The FP8 dot has to pay it -- a 256-entry E4M3 table has no in-register form -- but E2M1 does not: its sixteen values are eight magnitudes and a sign bit, so one `vpermps` on an eight-float magnitude table plus a shifted XOR decodes eight nibbles. The result is bit-identical to indexing the value table, sign of zero included, and it is what the format's structure allows rather than a numerics trade. Measured single-threaded, 4096 columns, 2048 rows, three runs: fp8 1484-1529 ns/row mxfp4 782- 808 ns/row 1.87x bf16 930- 946 ns/row 1.60x The gather form it replaces measured 2390 ns/row against FP8's 2529 under the same load -- 1.06x, which is why the first integrated arm showed no win and is worth recording as the reason. Co-Authored-By: Claude Opus 5 <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_014jtzY5AZ3KBAjBA5BHgPzH
This reverts commit ccf5897.
0213 turned this off: three device arms measured +1.71% decode while the tier removed 11.1% of host expert bytes, and host expert service moved only -0.21%. A saving worth about 5.4 s materialized as about 0.1 s and no record said where it went. It went into the round trip, and the round trip was expensive because the device was busy. 0213 predates the resident MLA chain landed in 0214, which took activation D2H from 2.265 to 0.002 GB per token, kernel time down 98.3% and synchronizations down 49.5%. The device is now idle at 0.77% of a decode step and the round trip costs 70 us. The mechanism did not change; its dependency did. Re-measured as 3+3 alternating pairs with an ordering repeat, twice, since a shipped default must not rest on a benchmark-only memory policy: interleaved decode 50.130 -> 46.650 s -6.94%, service -7.77% default policy decode 50.360 -> 46.890 s -6.89%, service -7.50% 4.1-4.3 s of host dispatch work removed against 0.76 s of collect wait added. No range overlap in either pair, ordering repeats +/-0.20%, all sixteen outputs hash fe74dc4ec7ab. The byte-to-service conversion is 0.70 per 1%, which is the curve 0211 measured. The default and an explicit request are now distinct values, because they must differ where a checkpoint cannot host the tier: the kernel decodes E4M3, so an MXFP4 checkpoint declines silently under the default and errors under an explicit opt-in rather than quietly running the host path that opt-in asked to replace. Co-Authored-By: Claude Opus 5 <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_014jtzY5AZ3KBAjBA5BHgPzH
Keep the release's BF16 shared-expert weights in checkpoint-native form and reproduce the host AVX2 eight-accumulator dot association exactly on CUDA. The host still owns every BF16 rounding, clamp, SwiGLU, and reduction boundary.\n\nThe format-tagged descriptor preserves the existing FP8 path while automatic admission now selects FP8 or BF16 from the opened checkpoint. The BF16 tier is 2.11 GB across all 42 MoE layers and reduces the cache arena by the same physical bytes before allocation.\n\nA row-complete CUDA test compares all three BF16 projections bit for bit against the host dot. The real MXFP4 two-arm smoke is output-identical and make check passes.
Document the two pinned checkpoint representations, manifest-driven runtime selection, exact native shared-expert tiers, current defaults, and the protected 3.88 tok/s MXFP4 operating point.
This reverts commit f7df715.
This reverts commit 62debab.
…context limit The chat and server commands omitted two things the published decode rates depend on. CUDA_DEVICE_ORDER=PCI_BUS_ID is required because many shells export FASTEST_FIRST, under which --devices 1,2 silently selects a different card and the run still succeeds with different output. numactl --interleave=all is worth about 1.13x and removes a 6.7% run-to-run placement lottery measured in 0218. Also documents why context is capped at 2,048: the checkpoint's k-pool sparse indexer is unimplemented, and the dense MLA workspace would need 47.2 GiB at 32,768 tokens against a 2 GiB per-device reserve. It is a missing feature, not a knob, and it is refused at admission rather than silently approximated. Co-Authored-By: Claude Opus 5 <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_014jtzY5AZ3KBAjBA5BHgPzH
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
1 participant
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.
Campaign branch for GLM-5.3 Flash text performance. 38 commits; every mechanism
is measured, and every rejection is preserved with its evidence rather than
quietly dropped.
Headline, measured on the composed stack
Canonical protected short decode — 36-token prompt, 129 generated tokens,
context 2,048, two RTX 3090s, greedy, seed 33,377,335,
numactl --interleave=all,three fresh processes per release after a discarded warm-up:
M0's canonical warm short decode was 1.107 tok/s on FP8. These are integrated
figures, not a product of per-record deltas.
Long context (1,536-token prompt) went 1.423 → 4.132 tok/s from the MLA kernel
repair alone; it has not been re-measured since the static expert tier landed,
so no combined long-context figure is claimed.
What landed
block-128 checkpoint or the quark MXFP4 export, selected at open from pinned
index extents. Host MoE gained MXFP4 and BF16 decoders beside the FP8 one.
The E2M1 nibble order was measured, not assumed — 0.993 cosine against the FP8
quantization of the same tensor versus 0.000002 for the swap.
glm53_mla_scores_kernelwas launching 64blocks on an 82-SM GPU at 2.08% achieved occupancy and 0.64% of DRAM
bandwidth. Long decode 1.423 → 4.132 tok/s, MLA kernel service −96.7%, exact.
kernel_nanosecondswas accumulated in onefile only, so the resident MLA, KDA and expert kernels reported ~zero. Every
"the GPU is idle" conclusion in this campaign rested on that; it is repaired.
MXFP4 / −19.1% FP8 decode at 35–40% measured route coverage.
MXFP4 (−25.8% decode).
recurrence, exact concurrent decode cohorts, bounded admission derived from
host and per-sequence CUDA state.
vpermpsrather thanvgatherdps: 1.87x the FP8 row dot,bit-identical.
What was rejected, with evidence
llama.cpp PR #27342; worth 1.085x here against a 25–30% bar. Speculation
amortizes fixed per-step cost, and after this campaign's repairs the step is
72.7% expert weight streaming at the hardware ceiling. The MLA repair moved
break-even from 2.32 to 4.58 — the largest win made the speculative milestone
unreachable.
expert bytes/token fall 17%, but expert-major sharing converts memory traffic
into arithmetic this host has none spare of. M7's paging, admission, prefix
reuse and cancellation are retained on their own merits.
reverted; no residue in the tree.
The transferable result: every "reuse a weight read across more rows"
mechanism — speculation, wider pages, larger cohorts — is bounded by a ~55 GB/s
compute roofline against a 33 GB/s memory ceiling, already ~90% consumed at one
row per traversal.
Verification
Clean CPU and CUDA builds; 358/396 tests with 38 fixture skips;
make check4/4with a Gemma fixture-only skip. No rejected mechanism remains in the tree. No
WIP or fixup commits.
README.mdanddocs/models/glm53.mdupdated.Known gaps
prefill shapes, the server concurrency/prefix matrix — has not been re-run
on the composed stack. Only decode has.
argument for why is recorded rather than argued around.
landed.