[Analysis] The DRAM bandwidth wall: decode performance ceiling for streaming MoE, measured paths forward

[ZacharyZcR]

Summary

Streaming MoE decode (S=1) on GLM-5.2 744B is DRAM-bandwidth-bound, not PCIe-bound and not compute-bound. This analysis quantifies the physical ceiling, shows we are already at 79–96% of it on a 6x RTX 5090 rig, identifies where the remaining 4–21% is lost, and proposes research directions — both software and hardware — for the community to discuss.

All numbers below are from controlled measurements on a 6x RTX 5090 + dual Xeon Silver 4510 (24 physical cores) + 251 GB DDR5 system, with GLM-5.2 int4 (gs=64), DRAFT=0, usage-snapshot-controlled placement, and PROF=1 instrumentation. PRs #429 (NUMA arena), #432 (device router), #433 (grouped kernels), #434 (resident accumulation), and #445 (pin budget fix) are all applied.

---

1. The bandwidth hierarchy — where the wall actually sits

For decode S=1, every expert evaluation is a GEMV: pure memory-bound, zero compute reuse. The only thing that matters is how fast weight bytes reach the arithmetic units.

Tier	Theoretical peak	Measured	Time to read one expert (~20 MB)
GPU VRAM (GDDR7, per card)	~1,792 GB/s	—	0.011 ms
PCIe 5.0 x16 (per card)	~63 GB/s	~50 GB/s	0.4 ms
CPU DRAM (DDR5, dual-socket)	~120 GB/s	62–70 GB/s	0.3 ms
NVMe SSD (PCIe 4.0 x4)	~7 GB/s	host-dependent	~3 ms

The critical ratio: VRAM bandwidth is 25–28x DRAM bandwidth. Once an expert lives in VRAM, it is effectively free. Once it lives in RAM, it costs 0.3 ms regardless of who computes it.

2. PCIe is not the bottleneck (and NVLink would not fix it)

A common intuition is that PCIe is the chokepoint between GPU and CPU memory. The data says otherwise:

PCIe 5.0 x16 throughput:   ~63 GB/s (measured ~50 GB/s)
DDR5 DRAM throughput:      ~120 GB/s theoretical, 62-70 GB/s measured

These are the same order of magnitude. Sending a RAM-resident expert to the GPU over PCIe (~0.4 ms) costs about the same as the CPU reading it directly from DRAM (~0.3 ms). The H2D transfer buys nothing — the GPU's 1,792 GB/s internal bandwidth advantage is gated by the DRAM egress rate, not by the interconnect.

Even replacing PCIe with NVLink-C2C (900 GB/s, as in Grace Hopper) only moves the bottleneck from the interconnect to the DRAM exit:

Interconnect	GPU reads RAM expert at	Bottleneck
PCIe 5.0 x16	~63 GB/s	PCIe ≈ DRAM (both limit)
NVLink-C2C	min(900, 120) = 120 GB/s	DRAM egress

That is a ~2x improvement, not a regime change. The wall is the DRAM itself.

We validated this experimentally: a "stream RAM experts to GPU" prototype (H2D + GPU compute per expert) ran 4x slower than CPU-direct (4.24 → 1.16 tok/s) due to per-expert H2D launch fragmentation — each of the ~4 RAM experts per layer is only ~20 MB, and the transfer setup overhead dominates. The branch is preserved at experiment/stream-ram-experts-FAILED.

3. Physical ceiling calculation

GLM-5.2 decode reads per token:

78 layers x topk=8 experts/layer x ~20 MB/expert = 12,480 MB ≈ 12.2 GB/token

Theoretical ceiling at various VRAM residency ratios (assuming 35 ms fixed attention + dense cost, measured):

Scenario	CPU read volume	At DDR5 measured (70 GB/s)	At DDR5 theoretical (120 GB/s)	Ceiling tok/s
100% RAM, 0% VRAM	12.2 GB	174 ms	102 ms	4.8 – 7.3
50% VRAM (our 6x5090)	6.1 GB	87 ms	51 ms	8.2 – 11.6
100% VRAM	~0	~2 ms	~2 ms	~27

Our measurements vs. ceiling

Configuration	Measured tok/s	Ceiling (measured BW)	Utilization
General workload, 6x5090	6.4–6.9	~8.2	79–84%
Prompt-tuned placement, 6x5090	7.9	~8.2	96%

The prompt-tuned result is essentially at the DRAM-bandwidth ceiling for this hardware. General workloads lose ~20% to placement inefficiency.

4. Why DDR5 utilization is only 58% (70 of 120 GB/s)

The ~50 GB/s gap between theoretical and measured DRAM bandwidth is the single largest source of lost performance. Root causes, in order of impact:

a) GEMV granularity is too fine

Each expert is ~20 MB. At S=1, the CPU scans each expert's weight matrix exactly once with a single-row dot product. The working set per OMP parallel region is ~200 KB per thread — too short for the memory controller to establish a sustained streaming pattern.

b) OMP fork/join fragmentation

The current execution model opens ~2 OpenMP parallel regions per expert (gate+up fused, then down projection). With ~4 CPU-tier experts per layer × 78 layers = ~312 expert evaluations per token, that is ~624 OMP fork/join cycles per token. Each cycle costs ~5–10 μs in barrier overhead and, more critically, invalidates the per-core prefetch state.

The XEXP=1 path (cross-expert single OMP region, colibri.c:3180) demonstrates the fix: one region across all experts in a 64-expert block, two internal barriers, verified byte-identical. It is currently CPU-only and gated on specific conditions (S==1, all int4, all resident, no CUDA).

c) NUMA non-uniformity

Dual-socket Xeon with DDR5 channels split across two NUMA nodes. Expert slabs allocated on node 0 and accessed by cores on node 1 pay a ~30% bandwidth penalty on the cross-UPI hop. PR #429's per-layer arenas stabilized bandwidth from ±21% variance to ±3%, but the mean is still ~70 GB/s, not 120 GB/s, because slab-to-core affinity is not yet optimized.

d) Memory controller occupancy

DDR5's peak bandwidth requires high outstanding-request counts to keep all banks busy. The short, sequential scan pattern of a single-expert GEMV does not generate enough concurrent requests — the memory controller's bank-level parallelism is underutilized.

5. Proposed research directions

Software optimizations (actionable now)

Direction	Mechanism	Expected impact	Relevant work
Cross-expert single OMP region	Generalize the XEXP=1 path to CUDA-enabled builds; one fork/join per block of 64, not per expert	Reduce OMP overhead ~10x; improve prefetch continuity	colibri.c:3180, #442
NUMA-aware expert scheduling	Pin experts to arenas on the NUMA node whose cores will compute them; process same-node experts consecutively	Eliminate ~30% cross-UPI bandwidth penalty	#429 (arena foundation), #82
Larger effective GEMV granularity	Batch adjacent-token expert sets (lookahead overlap is 36–41% at d=1–8); or fuse gate+up+down into a single streaming pass	Better memory controller utilization	#441, #442
E8/fmt=6 byte reduction	int3-g64-e8-rot containers: −25% expert bytes, +3.3pp quality vs shipped int4	25% less DRAM traffic per token; more experts fit in VRAM	#452 (ladder complete through converter)
Prefetch pipeline depth	Increase PILOT queue depth beyond 1; overlap next-layer prefetch with current-layer compute	Hide remaining disk latency; improve page-cache warmth	#441, #480

Hardware paths (for discussion and community benchmarking)

Platform	Unified BW	Expert read time (20 MB)	Projected ceiling (50% resident)	Status
Our rig (DDR5 + 6x5090)	70 GB/s (DRAM)	0.3 ms	~8 tok/s	Measured
NVIDIA Grace Hopper (GH200)	~500 GB/s (HBM3e CPU-side)	0.04 ms	~25 tok/s	Available; needs benchmark
NVIDIA DGX Spark (Grace Blackwell)	~200 GB/s (unified)	0.1 ms	~15 tok/s	Shipping; #161 has partial data
Apple M5 Max (unified)	~546 GB/s	0.037 ms	~27 tok/s	Available; needs benchmark
CXL-attached HBM	~250–400 GB/s (CXL 3.0 pool)	0.05–0.08 ms	~18–25 tok/s	Early availability (Samsung, SK Hynix CMM-H)
Processing-in-Memory (PIM)	Eliminates data movement	N/A	Architectural change	Research (Samsung HBM-PIM, UPMEM); not production-ready for MoE
Silicon photonics / optical I/O	>1 TB/s chip-to-chip	<0.02 ms	>30 tok/s	Research (Ayar Labs, Intel); 2026–2028 timeline

The key insight: unified high-bandwidth memory architectures (Grace Hopper, Apple Silicon) bypass the DRAM wall entirely by placing the expert weights in the same memory technology the compute units read at full speed. This is the only path to a regime change — software optimizations on DDR5 systems can recover the 58% → ~80–90% utilization gap, but cannot break the 120 GB/s ceiling.

6. Call for community contributions

We are opening this analysis to invite collaboration on three fronts:

1. Cross-platform bandwidth benchmarks. If you have access to Grace Hopper, DGX Spark, Apple M-series (M4 Max/Ultra, M5 Max), or CXL-attached memory systems: run GLM-5.2 with PROF=1 and report the expert-matmul and bandwidth numbers. The ceiling calculation above can be validated or refined with real data from these platforms.

2. Software optimization PRs. The XEXP cross-expert region and NUMA-aware scheduling directions above are concrete, measurable, and independently implementable. We would be happy to review and A/B any PR that targets the 70 → 120 GB/s gap.

3. Architecture discussion. If you have domain expertise in memory systems, CXL, optical interconnects, or PIM: what is realistic on a 1–2 year timeline for a system that can hold ~300 GB of expert weights at >200 GB/s sustained random-read bandwidth? The streaming MoE use case (many small sequential reads, each ~20 MB, to different addresses) is not the typical HPC streaming pattern, and we would value informed perspective on which technologies actually fit.

The fundamental physical constraint is clear: DRAM bandwidth is the wall, and the wall is 15–28x below GPU memory bandwidth. The question is whether we optimize within it or build around it.

---

Measured on 6x RTX 5090 + dual Xeon Silver 4510, GLM-5.2 int4-gs64, DRAFT=0. Related: #431 (CUDA choreography analysis), #442 (scalar reduction), #441 (PILOT QD), #452 (E8 container), #429 (NUMA arena). All data from PRs #432/#433/#434/#445 applied.

[ZacharyZcR]

Follow-up: we built the dual-lane prototype — the physics is real, and here is exactly where it breaks on a 251 GB box

The analysis above predicted that the DDR controller can serve two concurrent readers — the CPU and a GPU DMA engine — at a higher aggregate than either alone. We prototyped this end-to-end ("dual-lane": COLI_DUAL_LANE=N sends every Nth RAM-tier expert to a GPU that reads it out of DDR while the CPU loop computes the rest). Six iterations later we have a validated mechanism, a parity result, and three hard walls worth documenting so nobody re-derives them. Branch: feat/dual-lane-expert on our fork; everything env-gated, zero impact when off.

1. The microbenchmark says yes

CPU-only OMP reduction vs. PCIe H2D DMA (pinned source) vs. both concurrently, on the 6x5090 / dual Xeon Silver 4510 / 251 GB DDR5 rig:

readers	CPU	DMA	total DDR pull
CPU alone	70 GB/s	—	70 GB/s
DMA alone	—	56 GB/s	56 GB/s
concurrent 50:50	70	37	108 GB/s
concurrent 75:25	77	42	119 GB/s

Two readers genuinely lift total draw toward the 120 GB/s theoretical ceiling: +54–70% over the single-reader baseline. The lane is not about the GPU computing faster — it is a second concurrent pull out of the memory controller.

2. The end-to-end arc (GLM-5.2 int4, 128-token greedy, hit 100%, DRAFT=0, COLI_NUMA=1)

version	mechanism	tok/s	what it taught us
baseline	CPU single-lane	8.25	CPU leg runs at 70.9 GB/s
v1	sync cudaMemcpy from pageable slabs	4.57	pageable H2D is a CPU-mediated staging copy — it burns DDR twice and fights the CPU loop for the very bandwidth it was meant to add
v2	zero-copy kernel reads (registered + mapped)	1.71	demand-fetched PCIe reads are latency-bound: ~2.8 GB/s effective, not the 56 GB/s of bulk DMA. Zero-copy GEMV over PCIe is dead on arrival
v3	double-buffered DMA pipeline (copy/compute streams), arenas registered at load	8.16	parity. Per-expert cost now ~free, but 48 GB register budget covers only 19/78 layers
v4	persistent lane worker (kills pthread churn), 64 GB budget	8.01	the cancellation result: 13% of rows offloaded, CPU contended 71→62 GB/s — the two effects cancel to the millisecond
v6	arenas allocated as cudaHostAlloc pinned from birth, 76 GB	6.48	pinned arenas lose NUMA interleave (mbind is a no-op on driver mappings): CPU leg collapses to 38 GB/s

The v4 cancellation is the cleanest statement of the regime: offloading f of the bytes only wins if f exceeds the CPU's contention loss. On this box that crossover is around f≈0.35–0.45, and we could not reach it — see the walls.

3. Three walls, precisely characterized

Wall 1 — pageable staging (v1). cudaMemcpyAsync from non-pinned memory stages through a driver bounce buffer with a CPU memcpy: the "second lane" reads DDR at 3 passes per byte and steals CPU cycles. Any dual-lane design must DMA from page-locked memory, full stop.

Wall 2 — registration inflates RSS ~0.57x (v3–v5). cudaHostRegister on the pin arenas added ~0.55–0.59 bytes of anon-RSS per byte registered (48 GB → +24 GB, 64 GB → +38 GB; same with cudaHostRegisterDefault, so it is not the Mapped flag; mechanism unidentified — IOMMU/driver double-mapping suspected). A 251 GB box holding a 191 GB warm store can afford to register barely a third of it → offload fraction stuck at ~13% → parity forever. Full-store registration drew the OOM killer at anon-RSS 245 GB.

Wall 3 — pinned allocation breaks NUMA (v6). Allocating the arenas with cudaHostAlloc avoids the registration pass, but mbind(MPOL_INTERLEAVE) silently fails on driver-owned mappings, so the arenas land on one node and the CPU's own reads of them drop from 70.9 to 38 GB/s. The dual-socket interleave (#429/#384) is load-bearing for the CPU lane; any pinning scheme that forfeits it loses more than the lane gains.

4. What would cash the physics

The mechanism is validated and the engineering chain (DMA pipeline, persistent worker, per-layer batch handoff) now runs at zero marginal cost. What is missing is coverage — being able to pin most of the warm store without paying walls 2 or 3:

1. More RAM headroom (384–512 GB hosts): register the full store, keep interleave, predicted aggregate ~95–100 GB/s → ~9.5–10 tok/s (+15–21%) on this workload. If anyone runs a big-RAM dual-socket box with a spare GPU, the branch is ready to test.
2. Understand the 0.57x registration inflation. If it is IOMMU bounce-buffering or a driver setting, eliminating it makes wall 2 vanish on existing hardware. Datapoints from other kernels/drivers welcome — cudaHostRegister 48 GB and watch anon-RSS.
3. E8/fmt=6 (fmt=5: a deployable E8-lattice grouped container — int3 that beats shipped int4 quality at 25% fewer bytes (design + plan) #452) shrinks the warm store 25% (191 → ~143 GB), which alone frees enough headroom to register a majority of it. The compression line and the bandwidth line compose multiplicatively — this is one more reason the int3 container matters beyond quality-per-byte.
4. NUMA-aware pinning: interleave-place pages first (first-touch or numa_alloc_interleaved), then register — preserves wall-3's interleave if wall-2's inflation is solved.

5. Corrections to the opening analysis

Two numbers in the opening post need refinement in light of the measurements: the CPU's GEMV bandwidth is ~40 GB/s without NUMA interleave and ~71 GB/s with it (the 70 GB/s "measured DDR5" figure was the best case, not the typical one), and the "GPU lane at wire speed" assumption only holds for bulk DMA from pinned memory — zero-copy kernel reads run 20x slower than the link. Both corrections strengthen the original conclusion: the wall is DDR effective throughput, and software placement (NUMA, pinning, granularity) moves it by 2x before any hardware change does.

Full logs and the microbenchmark source are on the bench host; happy to share either. If someone wants to reproduce the microbenchmark on Grace Hopper / Strix Halo / Apple silicon (where the "two lanes" collapse into one unified controller), that comparison would directly test the unified-memory thesis from section 6 of the opening post.

[DavidePapero]

Very thorough analysis, thank you.

I make my usual premise that there is no polemical intent in me and I hope I'm not being annoying.

This analysis focuses on the bandwidth, but due to my (profound) ignorance I don't know a detail that I consider important to make myself understood. At a given instant, for all instants of the calculation, is the memory always occupied "exactly" 100% in counts? Maybe for "monster" computers this is true, although on "normal" computers it obviously isn't.

For cases where the answer is no then, I believe, the length of time a single block remains in memory becomes important. Not the average time, the distribution of the duration of permanence of the generic block, preferably together with the distribution of the size of the blocks. I know I was incomprehensible so I'll give an example.

Let's assume two levels of memory, one slow and one fast.

Ideal case, at any moment the fast memory would be enough to contain all the blocks to do the calculations. But over time it's not enough. If the residence time of a block is less than the transport time between slow and fast memory, it stays where it happens. If the permanence time of a block of data in memory is greater than the transport time and it is not possible to saturate the fast memory at all times, then if a block of memory "lands" on the slow memory it is still better to move it "while" it is being used and finish the calculations on the fast memory. This is not necessarily possible and I expect that the cost of managing synchrony will make this idea ridiculous.

Normal case of a low-end computer, but with two levels of memory (effectively with a graphics card). Memory not enough to do the math. The computing unit spends some of its time waiting. I expect performance to improve the larger the gap in megabytes of blocks carried in fast memory versus slow memory. This discussion makes sense only if the residence time of the data blocks in memory has a distribution and it is possible to know in advance which data blocks must go at high frequency, i.e. with low residence time in memory during a calculation batch.

That is, the bandwidth measurement of the various types of memory is important because it indicates bottlenecks, but, I believe, it is not sufficient. I believe that to optimize the process, the quantity of memory that is moved in the various types of memory is also necessary; the distribution of bytes over time in the various memories would be better, and not just the aggregate data. But to start with, even the aggregate data would be interesting

If in an LLM, in fact, there is simply a linear scan of the network, everything written is a load of nonsense. Don't kill me.

Davide

[ZacharyZcR]

@DavidePapero Not annoying at all — you've asked exactly the right question, and the answer is what makes streaming MoE decode such a strange workload.

Your framing (block residence time vs transport time, and their distributions) is the correct lens for most memory hierarchies. If a block gets reused while it's resident, a fast tier pays for itself; if its reuse interval is shorter than the transfer cost, moving it was a mistake. Classic cache economics.

The strange property of MoE decode at S=1: within one token, every expert weight block is read exactly once. The GEMV sweeps 20 MB of weights, produces one row, and never touches those bytes again for that token. Within-token reuse is zero — there is nothing for a cache to exploit. That is why the analysis above collapses to pure bandwidth: when every byte is used once per pass, the only thing that matters is how fast you can stream bytes, and "occupancy" is simply the fraction of time the memory controller has requests in flight (the 58% DDR utilization figure is exactly that number falling short).

Where your residence-time distribution comes back is across tokens: expert popularity is highly skewed (Gini ≈0.67 on GLM-5.2's routing traces), so a hot expert is re-read every few tokens while a cold one sleeps for thousands. The reuse-interval distribution you're describing is precisely what the engine's learned placement consumes: .coli_usage records the popularity, the pin/VRAM tiers keep the short-reuse-interval experts in fast memory, and the measured hit rates (50–100% depending on RAM) are that distribution being cashed in. Your "if residence is shorter than transport, leave it where it is" rule is also literally the engine's decision boundary — cold-tail experts stay on disk because promoting them costs more transfer than their reuse justifies.

So both views are right, at different timescales: within a token it's a pure streaming problem (bandwidth wall), across tokens it's the tiering problem your distributions describe — and the second one is already what the placement layer optimizes. The one place where within-token reuse could be created is batching multiple requests so one expert serves several rows per read — and that's blocked by the physics in section 4: concurrent requests' expert sets barely overlap (~3%), so the reads don't merge. That non-overlap is, in your language, a reuse-interval distribution with almost no mass at "immediately".

[RDouglasSharp]

I am on a 128Gb M5 Maxed out Max. I am happy to benchmark, if you specify what the exact benchmark you are looking for is.

Here is the Profile from the short benchmark I am typically running (this includes unmerged #587 ). I have a different benchmark for focusing on long prefill.

PROF=1 PROMPT="Compare the myths of Lucifer and Prometheus" NGEN=64 COLI_METAL=1 DIRECT=1 MTP=0 TEMP=0 COLI_METAL_RESSET=1 SNAP=~/glm52_i4_g128 ./colibri 64

prefill 9 tokens in 8.22s | decode 64 tokens in 30.30s (2.11 tok/s) | expert hit rate 73.0% (pin 61.1% + lru 11.9%) | RSS 92.43 GB
experts loaded/token: 590.6 (per-layer 7.88 across 75; baseline topk=8) | TOPK=0 TOPP=0.00
speculation: 1.02 tokens/forward (63 forwards per 64 tokens) | MTP acceptance 0% (0/0)
PROFILE: expert-disk 54.579s service / 18.177s wait | expert-matmul 4.981s | attention 5.003s (including kvb 0.000s) | lm_head 0.000s | other 2.143s
ATTENTION: projection/RoPE 0.000s | score-softmax-value 0.000s | output projection 0.000s
P0-EXEC: routed CPU 0.000s / 0.00 GB/s (0 row) | routed GPU critical 0.000s | router 0.006s | residual P2P 0.000s / 0 hop | orchestration 2.137s
METAL-ATTN: layer GPU 4914 | gpu-wall 4.81s (kernel 3.83s | cpu-sched 0.04s gpu-sched 0.55s)
METAL: blocchi GPU 9010 | fallback CPU 0 | expert su GPU 41454 | setup 0.87s gpu-wall 4.42s (kernel 6.33s) scatter 0.10s
METAL-RESSET: flush 0.32s (residency-set commit in moe_submit, outside setup/gpu-wall)
[PROF] decode forwards: 63 | latency p50 468.3 ms | p90 550.6 ms | p99 580.6 ms | max 632.0 ms | 1.02 tok/forward
[PROF] expert I/O: 204.972 GB fetched (3202.7 MB/token, 6.76 GB/s over the run) | hit 73.0% (23078 pin + 4501 lru / 10221 load) | 590.6 loads/token | 54.6s read service / 18.2s felt wait
[PROF] DISK-CLASS (recency split): cold 10127 x 203.09 GB @ 3.76 GB/s-thread, wall 18.0s @ 11.27 GB/s-wall, avg-conc 3.0 (direct 10091/10127) | warm 94 x 1.89 GB @ 3.33 GB/s-thread, wall 0.5s @ 3.84 GB/s-wall, avg-conc 1.2 (direct 94/94) | disk-busy 18.0s (60% of window)
[PROF] resident experts: 2184 pinned (43.8 GB) + 2175 in LRU (43.6 GB, cap 29/layer)
[PROF] time shares: expert-I/O 60% | expert-matmul 16% | attention 17% | lm_head 0% | other 7%
[PROF] P0 execution: routed CPU 0.000s / 0.00 GB/s (0 row) | routed GPU critical 0.000s | tier straggler 0.00x | router 0.006s | residual P2P 0.000s (0 hop, 0.000 ms/hop) | orchestration 2.137s
[PROF] verdict: I/O-bound — 60% of the time waits on expert reads (hit 73%). More cache is the lever: raise RAM_GB (or add RAM). Try PIPE=1 (overlap reads with matmul).

[ZacharyZcR]

Thank you — an M5 Max is exactly the platform this analysis needs, because it's where the "two lanes collapse into one unified controller" question from section 6 gets tested for real. But your profile shows the machine isn't hitting the wall this issue is about yet, and why is the interesting part.

Your run is disk-bound, not bandwidth-bound. The engine's own verdict says it: hit 73%, 6.76 GB/s expert I/O, 60% of decode waiting on SSD reads. That 6.76 GB/s is your SSD, not your unified memory. The 546 GB/s unified-memory advantage this issue is about never comes into play here — because 128 GB can't hold the ~300 GB g128 container, most experts stream from disk every token. So your box is currently measuring section 4's disk wall, not the DRAM wall.

And that's itself a datapoint worth stating in the thread: on Apple silicon the capacity wall gates the bandwidth advantage. A unified 546 GB/s controller is wasted if the model can't fit in it — the experts are on SSD at 7 GB/s regardless. This is one more reason the E8/int3 container (#452) matters specifically for Apple: shrink GLM under ~128 GB and the M5 Max flips from disk-bound to a pure unified-memory-bandwidth test, which is the regime where its architecture should actually win.

Two benchmarks that would be genuinely valuable, both runnable today:

(1) The unified-memory bandwidth number itself — the one cell in the table above that's a guess for Apple, not a measurement. A 20-line STREAM-style copy is enough:

// membw.c — clang -O3 membw.c -o membw && ./membw
#include <stdio.h>
#include <stdlib.h>
#include <string.h>
#include <time.h>
int main(void){
  size_t n = 8UL<<30;                 // 8 GB buffers
  char *a = malloc(n), *b = malloc(n);
  memset(a,1,n); memset(b,2,n);       // fault both in
  struct timespec t0,t1; clock_gettime(CLOCK_MONOTONIC,&t0);
  for(int i=0;i<10;i++) memcpy(b,a,n);
  clock_gettime(CLOCK_MONOTONIC,&t1);
  double s=(t1.tv_sec-t0.tv_sec)+(t1.tv_nsec-t0.tv_nsec)*1e-9;
  printf("memcpy: %.0f GB/s (read+write)\n", 2.0*n*10/s/1e9);
  return 0;
}

That gives the sustained effective bandwidth to compare against the 546 GB/s spec — the direct analogue of the "DDR5 120 theoretical / 70 measured" line for x86.

(2) A hit-rate sweep — RAM_GB from small to as large as your 128 GB allows, logging hit% and tok/s at each. That maps where your box crosses from disk-bound to bandwidth-bound, and how close to 100% hit you can push before capacity runs out. The [PROF] verdict line flips from "I/O-bound" to "compute/bandwidth-bound" somewhere on that curve — finding that crossover on Apple is the datapoint no x86 box can provide.

The COLI_METAL_RESSET=1 residency-set work you're doing (#587) is directly relevant to (2), since it changes how much stays resident. If you run either, the numbers go straight into the table. Much appreciated.

[RDouglasSharp]

int main(void){
size_t n = 2UL<<30;
char *a = malloc(n), b = malloc(n);
if(!a||!b){ fprintf(stderr,"alloc fail\n"); return 1; }
memset(a,1,n); memset(b,2,n);
struct timespec t0,t1; clock_gettime(CLOCK_MONOTONIC,&t0);
for(int i=0;i<5;i++) memcpy(b,a,n);
clock_gettime(CLOCK_MONOTONIC,&t1);
volatile char sink = b[0] + b[n-1];
double s=(t1.tv_sec-t0.tv_sec)+(t1.tv_nsec-t0.tv_nsec)1e-9;
printf("memcpy: %.0f GB/s (read+write) sink=%d\n", 2.0n5/s/1e9, (int)sink);
return 0;
}

returns

memcpy: 189 GB/s (read+write) sink=2

[ghostnumber7]

I have been doing some testing on this same topic over the last couple days, so here's my 2 cents:

Box: i9-14900KF (single socket) · 2×48 GB DDR5-4800, dual-channel · 1× RTX 5090 32G VRAM.
GLM-5.2 744B int4, greedy decode.

DRAM bandwidth — @ZacharyZcR membw.c, run on this box (had to drop -O3, it dead-stored the copy):

memcpy: 32 GB/s (read+write)

That's single-thread — one core. Across all 32 cores the same memcpy aggregates to ~63 GB/s (read+write), which is the controller ceiling. The ~2× gap between one core and all of them is the headroom a second concurrent reader — the GPU pulling over PCIe — can actually use, which is the whole reason the streaming path buys anything here.

all-cores variant of @ZacharyZcR's memcpy, split across OpenMP threads

// gcc -fopenmp membw_mt.c -o membw_mt && ./membw_mt   (no -O3 — it dead-stores the copy)
#include <stdio.h>
#include <stdlib.h>
#include <string.h>
#include <time.h>
#include <omp.h>
int main(void){
  size_t n = (size_t)8<<30;                     // 8 GB  ((size_t) cast: Windows `long` is 32-bit)
  char *a = malloc(n), *b = malloc(n);
  memset(a,1,n); memset(b,2,n);                 // fault both in
  struct timespec t0,t1; clock_gettime(CLOCK_MONOTONIC,&t0);
  for(int i=0;i<10;i++)
    #pragma omp parallel for schedule(static)
    for(size_t c=0;c<n;c+=(1u<<20))
      memcpy(b+c, a+c, (c+(1u<<20)<=n) ? (1u<<20) : (n-c));
  clock_gettime(CLOCK_MONOTONIC,&t1);
  double s=(t1.tv_sec-t0.tv_sec)+(t1.tv_nsec-t0.tv_nsec)*1e-9;
  printf("memcpy x%d threads: %.0f GB/s (read+write)\n", omp_get_max_threads(), 2.0*n*10/s/1e9);
}

The thing I was testing (branch is not clean, just committed to add as reference)

https://github.com/ghostnumber7/colibri/tree/perf/stream-tier-async

RAM_GB sweep — streaming path off vs on, and where the box crosses from disk to bandwidth-limited (single 128-token runs; I/O/matmul are [PROF] time-shares):

RAM_GB	hit%	I/O time	matmul time	tok/s off	tok/s on
20	73	50%	35%	1.49	1.63
30	73	50%	36%	1.48	1.58
40	70	51%	35%	1.40	1.55
50	76	48%	37%	1.62	1.76
60	83	39%	44%	1.77	1.93
70	84	40%	43%	1.88	2.06

Two things fall out:

1. The streaming path adds a steady ~+9% at every RAM level on this box with this particular config — it moves the CPU expert-matmul onto the GPU, and there's always enough of it to matter on the "same query mostly deterministic" ideal path.
2. The disk→bandwidth crossover is visible: below ~RAM_GB 55 (hit ≤76%) disk I/O is the bigger share (≈50% vs ~35% matmul); by RAM_GB 60 (hit ~83%) it flips and the DRAM-bandwidth-bound expert-matmul (43–44%) overtakes it. The engine's own verdict line still prints "I/O-bound" throughout — it just points at "raise RAM" — but the time-shares are past the crossover by 60.

NOTE: RAM_GB=60 is the value I normally operate at due to Windows 11, WSL2, Docker Desktop, VSCode and Chrome competing for RAM.

A tighter check at RAM_GB=60 (N=15 interleaved, 256-token) puts the streaming gain at +8.4% (1.90 → 2.06, IQRs disjoint). The tok/s difference with 128 get runs is explained because 128 get table is one-shot, and 256 here was taken from median of 15 samples.

Disk stays a co-factor because the model doesn't fit 96 GB (hit tops out ~84%); we keep it from dominating by mirroring across two SSDs (~10 GB/s effective), which needed the Windows O_DIRECT mirror-shard fix — similar to #593.

The flags I use as base, which appear to be the ones to allow 1.9 tok/s at 256 token gen in the first place are as follows:

COLI_MODEL=D:\glm52;COLI_MODEL_MIRROR=C:\glm52;COLI_CUDA=1;DIRECT=1;PIN_GB=18;PILOT_K=6;PILOT_REAL=1;PILOT_EVICT_GUARD=0;CACHE_ROUTE=1;COLI_CUDA_PIPE=2;CUDA_DENSE=1;COLI_GROUP_ASYNC=1;CUDA_RELEASE_HOST=1;COLI_TEMP=0;
PIN=stats256.txt; # from another initial run without a pin

Full engine log — one on-run (STREAM_GB=60, warm)

[GPU] auto-enabled CUDA · NVIDIA GeForce RTX 5090 · 30.7 GB expert tier
[CUDA] device 0: NVIDIA GeForce RTX 5090, 34.2 GB VRAM, sm_120
[CUDA] stream tier: PCIe expert streaming on (STREAM_GB=60.0 pinned budget, pin tier only)
[CUDA] mode: routed experts + resident dense tensors; VRAM experts without host backing
[MIRROR] probe: primary 4.63 GB/s, mirror 4.39 GB/s
[MIRROR] C:\glm52: 144/144 shards | read split primary 51% / mirror 49% (measured)
[CUDA] hot expert tier: 1059/2573 experts, VRAM 20.03 GB (budget 30.7 GB, reserve 2.0 GB)
[PIN] placement: 1059 VRAM + 1514 RAM expert (28.6 GB warm) in 8s
[CUDA] stream tier: 1514 pin experts pre-registered for DMA (28.7 GB pinned)
loaded in 4.56s | resident dense: 9912.75 MB | layers=78 experts=256 | MTP ACTIVE (draft=0)
prompt: 17 tokens | generating up to 256
PROFILE: expert-disk 137.320s service / 2.482s wait | expert-matmul 3.243s | attention 0.639s | lm_head 0.178s | other 3.092s
MIRROR: primary 54.09 GB (11460 reads) | mirror 50.51 GB (10702 reads) — 48% of expert bytes from the mirror
[t=16   RSS 51.22 GB  hit 87%  swap 15%  2.51 tok/s]
[t=64   RSS 51.23 GB  hit 85%  swap 16%  2.29 tok/s]
[t=128  RSS 51.24 GB  hit 84%  swap 17%  2.13 tok/s]
[t=192  RSS 51.25 GB  hit 83%  swap 17%  2.09 tok/s]
[t=256  RSS 51.26 GB  hit 82%  swap 17%  2.06 tok/s]
[CUDA] resident set: 3802 tensors, 29.19 GB VRAM
[CUDA] expert groups: 18512 call, 68583 expert, 70992 righe (3.70 expert/call)
[CUDA] stream tier: 34276 experts streamed (648.62 GB over PCIe)

[monotophic]

Authored by Fable 5 in Claude Code, analysis in partnership with @monotophic (We/our = the joint analysis; me/my = @monotophic's own hardware and commitments).

Great stuff — the tier framing matches what we've measured from the other side of the wall. @RDouglasSharp has the fresh M5 Max PROF run covered; what we can add is the measured decomposition from an extended campaign on my M5 Max (128 GB), plus two refinements to the table and an instrument offer.

1. The M5 Max row needs one correction: the wall relocates, it doesn't disappear

The projected ~27 tok/s assumes the non-VRAM tail streams from unified RAM. On a 128 GB box holding this 358 GB container, only ~31% fits — the tail comes from SSD at a measured ~14 GB/s device wall (sustained, pattern-specific; we verified it's the device ceiling at any queue depth). The general law our campaign converged on: unified memory breaks the DRAM wall only when the working set fits; otherwise the wall relocates to the next tier down. Measured reality here:

quantity	measured (M5 Max, 2026-07, macOS 26.5)
best full-fidelity decode	2.89 tok/s (identical-output contract, hit 74.3%)
decode equation	cold-bytes/token ÷ ~14 GB/s + serial compute
cold bytes at hit h	≈ your 12.2 GB × (1−h) — cache term for your ceiling model
steady-state hit ceiling	~74% for this workload — LRU steady-state and matured pinned profiles converge on the same number, so it looks like a routing-concentration constant of model+workload, useful as a principled placement parameter in your table (your "~20% placement loss" is the same phenomenon from the other side)

Your key insight holds exactly in the fits-in-memory regime — where your ~27 tok/s projection is the right arithmetic for models that fit whole. For 744B-streaming on 128 GB, the storage tier owns the ceiling.

2. On the PILOT/prefetch-depth row: storage-tier caution from instrumented refutation

We built and instrumented three generations of prefetch on this box (issue-time CPU prefetch, compute-phase-held issue, page-cache-filtered demand). All refuted with the instrument on the storage tier: once reads sit at the device wall, prefetch is zero-sum relabeling (our early "disk 98% busy" reading turned out to be occupancy, not saturation — worth guarding against in any tier's telemetry). Full record on #387. Your DDR5-tier version may behave differently — different tier, different physics — but the occupancy-vs-saturation trap is portable.

3. Byte reduction: complementary FR just posted

#594 proposes lossless entropy compression for streamed expert traffic as a transport layer (any container, decompress at cache-fill, kernels untouched): measured 26.2% size reduction on real GLM-5.2 experts with a working codec, and a Metal decode prototype at 24.3 GB/s (above our storage wall). It attacks the storage tier the way your E8 row attacks bytes-in-memory — complementary, not competing.

4. Instrument offer (the "digital twin")

Your §3 ceiling table is the right first-order model. We maintain a measured version of it — per-token serial-sum decomposition from in-tree telemetry (attention / MoE / orchestration / GPU-wait / disk-felt), interleaved-control bracket discipline for thermal drift, and an interaction matrix (including a measured case where CPU↔GPU intervention effects violate serial-sum additivity). Happy to share the methodology + our full M5 Max dataset so the cross-platform table can be populated with decomposed, drift-controlled numbers rather than headline tok/s. A fuller re-baseline on current dev is already scheduled on my side and can feed the table directly.

(On your question 3: Ultra-class Apple Silicon with 512 GB unified already meets the ">200 GB/s over ~300 GB" spec for this model class today — that's the regime-change box your table implies, if anyone in the community has one to bench.)

[DavidePapero]

....

// membw.c — clang -O3 membw.c -o membw && ./membw
#include <stdio.h>
#include <stdlib.h>
#include <string.h>
#include <time.h>
int main(void){
size_t n = 8UL<<30; // 8 GB buffers
char *a = malloc(n), b = malloc(n);
memset(a,1,n); memset(b,2,n); // fault both in
struct timespec t0,t1; clock_gettime(CLOCK_MONOTONIC,&t0);
for(int i=0;i<10;i++) memcpy(b,a,n);
clock_gettime(CLOCK_MONOTONIC,&t1);
double s=(t1.tv_sec-t0.tv_sec)+(t1.tv_nsec-t0.tv_nsec)1e-9;
printf("memcpy: %.0f GB/s (read+write)\n", 2.0n10/s/1e9);
return 0;
}

....

I know I am testing your patience.

Bear in mind that to make a living, I've spent the last 30 years programming first in Visual Basic and later in C#. The universe has been very hard on me, so don't rub it in :)

The example you posted seems very relevant. You are indeed measuring RAM bandwidth, but the code is sub-optimal. Copying one byte at a time fails to exploit the bus width, and by using memcpy you are polluting the processor cache. You should use a byte count that is a multiple of the bandwidth and use "non-temporal store" copy instructions.

You seem to think that to speed up the process, you need to saturate the data bandwidth. I don't think that's the right approach. Your goal should be to maximize the load on the various compute units, and bandwidth is a constraint that prevents you from keeping those compute units fed. I say "compute units" because with a modern, "un-crippled" CPU (Apple Silicon, Intel, and AMD), you have multiple standard cores plus the graphics cores (which are currently unused) with unified memory, and then potentially discrete graphics cards.

I should preface this by saying I analyzed the code with Gemini Flash. Despite what Vincenzo writes, understanding this code is very complex for someone outside this domain like me.

The hardware architecture of a consumer PC with a discrete graphics card is generally as follows:

• 3 memory systems
  • SSD
  • RAM
  • VRAM
• 3 controllers
  • On an Intel/AMD with integrated graphics or on an Apple M-series chip, I have no idea if there's another controller embedded in the silicon.
• Various compute units
  • CPU
  • GPU
• Various levels of integrated cache
  • CPU
    • L1
    • L2
    • L3
  • GPU
    • L1
    • L2

Then there's the "problem" of the minimum read unit. For RAM, it's one byte, but for an SSD, it will always be a few KBs (depending on how it was formatted).
This means that with an SSD, ideally, you should read contiguous memory blocks that are at most slightly smaller than what will actually be read by the controller. Consequently, there are specific edge cases where, to improve SSD performance, you have to reorganize data so that frequently accessed reads are located close together (spatially local).

Why do I talk about maximizing CPU load? At the core of any engine are calculations. A CPU/GPU/TPU will always have to load data into internal registers and then perform the required math. Hence the need to account for internal caches.

While chatting with the AI over a glass of milk about GPUs—a dark art to me—I also discovered the issue of "memory coalescing". It showed me some code, and I admit my limits as a programmer. Then it also told me: "Watch out, even on the H100/H200, Nvidia traded throughput for latency." I had always assumed that access latency was negligible, and this information throws me off.

That being said, if you look at the big picture, you can arrive at counter-intuitive solutions. For instance, if you are moving large blocks of data, it's better to use instructions that bypass the CPU/GPU, thereby bypassing the internal caches. Conversely, if you need to move blocks that will be used immediately, it's probably better to lower the throughput and use a memcpy that pollutes the L3 cache (or the GPU equivalent to pollute the L2 cache). This is because the data will have to be moved there anyway to perform the calculations (in reality, it will probably traverse L3->L2->L1->CPU, but let's not overcomplicate things).

You'll have to excuse me @ZacharyZcR , but unfortunately for you, I find these discussions about big-picture architecture very stimulating.

Davide

[ZacharyZcR]

2026-07-27 follow-up: the remaining software levers were tested on the 6x5090 host

We ran the pending directions from the opening post against the current full-residency champion regime. The useful result is a narrower ceiling: one memory-lifetime fix is real, but wider SIMD, SMT, cross-expert OpenMP, extra VRAM occupancy, and heat-balanced device assignment all fail to improve sustained decode.

Controlled baseline

• 6x RTX 5090, dual Xeon Silver 4510, 251 GiB RAM
• GLM-5.2 int4, frozen frozen-general.stats placement
• 512/8/8, DRAFT=0, TEMP=0, REPIN=32
• COLI_NUMA=1, 24 physical cores, OMP_PROC_BIND=spread, OMP_PLACES=cores
• COLI_GROUP_ASYNC=1, CUDA_RESERVE_GB=2
• 20-token prompt, 256-token decode, 100% expert hit

The correctly matched baseline is 8.23–8.55 tok/s. The run-to-run spread is mostly CPU-row/background-load variance; the stable profile is approximately:

CPU expert rows:       10.3–11.0 s / 256 tokens
attention:             11.1–11.2 s
expert matmul total:   12.3–13.1 s

Two benchmarking traps were found while reproducing it:

1. COLI_CUDA_PROFILE=1 inserts enough CUDA timing/synchronization to lower decode to ~6.3 tok/s.
2. COLI_CUDA_ATTN=1 is slower on this topology (CPU absorb + separate o-proj: ~11.1 s; fused CUDA attention: ~12.8 s). The exact projection-prefix path does not require COLI_CUDA_ATTN.

Positive result: reclaim detached pin-arena pages

Live GPU REPIN re-attaches the demoted expert's host arena slice, reloads it, uploads the promoted expert, then detaches that promoted host slice. Detaching preserved the virtual arena but left its anonymous pages resident. Each 16-expert REPIN batch therefore grew RSS by roughly 250–300 MiB.

The fix applies MADV_DONTNEED to a detached Linux arena slice after the synchronous H2D upload. The address/capacity remain reusable; expert_host_ensure() reloads the slice before any CPU use.

	RSS at t=16	RSS at t=256	REPIN batch
before	196.30 GB	198.09 GB	82–121 ms
after	196.29 GB	196.29 GB	97–101 ms

Greedy output after stripping telemetry is byte-identical (838 bytes), and the full suite is green (214 tests, 13 skipped). Experimental-fork commits: 527fb8c (exact attention prefix) and 0a346ab (arena-page reclaim).

Negative matrix

Every code experiment below was reverted after measurement.

experiment	result	mechanism-level verdict
AVX-512 MLA absorb dot/AXPY	8.07 tok/s; attention core 4.643 s vs 4.640 s	the loop is memory-bound; doubling SIMD width buys nothing and risks downclock
48 logical threads (OMP_PLACES=threads)	7.69 tok/s; CPU rows 6.48 s/128 vs ~5.47 s	SMT contention exceeds any latency-hiding benefit
CUDA-enabled hybrid XEXP (GPU early issue + RAM experts in one OMP region)	6.26 tok/s; CPU rows 9.76 s/128	only ~2 RAM experts/layer remain; 24 dynamic chunks destroy locality and cost more than fork/join
CUDA_RESERVE_GB=0.5	9,812 vs 9,335 VRAM experts, but 7.91 tok/s	+5.1% resident experts did not reduce the critical CPU-row wall; fuller VRAM is not automatically faster
assign VRAM experts by cumulative routing heat instead of bytes	7.08 tok/s; CPU rows 8.49 s/128	existing per-device historical heat was already within ~4.7%; cross-device regrouping changes FP accumulation, perturbs routing, and reverses the intended gain

The last two also changed greedy output because moving experts between CPU/GPU kernels or regrouping them across devices changes floating-point order. That is important for placement experiments: a different routing trace after the first near-tie can dominate the nominal static-heat improvement.

Updated conclusion

On this host the stable local optimum remains:

24 physical cores
NUMA interleave
COLI_GROUP_ASYNC=1
CUDA_RESERVE_GB=2
REPIN=32
exact q_a/kv_a activation-upload reuse

This sustains roughly 8.2–8.6 tok/s, with occasional short-window peaks around 9.2 tok/s. The opening post's proposed "generalize XEXP to CUDA builds" is now falsified for the current ~50% VRAM-resident regime: after GPU early issue, the RAM subset is too small per layer for block-wide dynamic scheduling.

The remaining path to a material gain is not another CPU micro-kernel or occupancy knob. It must reduce CPU-tier bytes per token without destabilizing routing: a smaller validated expert container (the E8/fmt=6 line), or a placement/substitution method with an explicit quality and token-identity policy. The experiments above tighten the DDR5 mixed-tier software ceiling; they do not change the hardware conclusion of this issue.

[ZacharyZcR]

2026-07-28 follow-up: paper-inspired inference ideas under fixed replay and held-out evaluation

We extended the campaign from local CPU/GPU tuning to several recent inference-system ideas: HybriMoE-style tier balancing, learned expert prefetch, WarpDecode, paged KV storage, speculative/MTP control, and workload-trained placement.

The main result is methodological: expert hit rate, expert count, and kernel throughput are not sufficient optimization targets. On this mixed CPU/6x5090 system, end-to-end latency depends on expert identity, CUDA group shape, per-device packing, NUMA locality, residency, and max(T_cpu, T_gpu). Several mechanisms that looked favorable in an analytic model or in-sample benchmark reversed under fixed-token or held-out evaluation.

Experimental discipline

The controlled replay runs used:

• GLM-5.2 int4, 6x RTX 5090, dual Xeon Silver 4510, 24 physical threads
• fixed teacher-forced token streams (identical 56-token decode inputs)
• NUMA interleave and async CPU/GPU expert-group execution
• all 19,456 experts resident, 100% expert hits, and 0.000 GB fetched during decode
• fresh-process A/B runs, with reversed order where practical

Runs violating residency or I/O invariants were excluded rather than compared.

Placement: large oracle gain, no transfer

A replay-specific oracle ranked experts using the exact replay's route trace. It is deliberately an in-sample upper bound:

placement	decode	CPU rows	GPU rows	CPU / GPU critical
broad Atlas baseline	8.74 tok/s	9,560	24,040	2.670 / 3.067 s
replay oracle, run 1	13.38 tok/s	457	33,143	0.161 / 0.774 s
replay oracle, run 2	12.64 tok/s	457	33,143	0.180 / 0.786 s

The repeated oracle median was 13.01 tok/s (+48.9%). Concentrating reused experts improved CUDA grouping enough that placing 98.6% of routed rows on GPU was faster, contradicting our linear per-selection cost model.

It did not generalize. On an independent fixed replay:

placement	decode	CPU rows	GPU rows
broad Atlas baseline	8.93 tok/s	1,110	3,090
replay oracle	5.54 tok/s	1,404	2,796

That is a 38.0% held-out regression, still with zero disk I/O. The in-sample gain is real, but workload-specific. A placement result without train/test separation can therefore be technically reproducible and still be misleading.

We also tried category-held-out robust frequency ranking on the original 10-category x 3-prompt Expert Atlas. A light variance penalty improved offline mean GPU coverage from 73.29% to 73.70% and the worst category from 68.04% to 69.32%. Real execution nevertheless regressed:

• 56-token replay: 7.48 -> 6.63 tok/s (-11.4%)
• independent short replay: 8.84 -> 6.53 tok/s (-26.1%)

Coverage increased while GPU critical time worsened. Routing coverage is not a latency model.

Tier and device balance

Shrinking the GPU expert budget from 176.6 GB to 171 GB while raising RAM capacity kept every expert resident, but reduced median decode from 7.26 to 6.88 tok/s (-5.2%). Balancing expert counts alone is not useful.

New per-device CUDA telemetry found a 17.7% spread in routed rows and a 14.5% spread in group calls across the six GPUs. An opt-in LPT prototype kept the same 9,335-expert GPU prefix but assigned experts by accumulated Atlas frequency instead of bytes alone:

paired run	byte-balanced	frequency-load assignment	change
1	7.83	7.86 tok/s	+0.4%
2	7.64	8.27 tok/s	+8.2%
3, reverse order	6.57	7.49 tok/s	+14.0%

The three-run median gain was +2.9%, but variance is large and only one replay has been validated. CUDA_EXPERT_LOAD_BALANCE=1 remains experimental and off by default. The per-device telemetry is retained because aggregate GPU counters hid a real imbalance.

Other paper-inspired results

• MTP soft guard: the strongest production-relevant result. Pausing low-yield speculation below a 70% acceptance window (24 proposals) improved the previous negative D1 path by 12.1%, reduced expert work, and preserved the high-acceptance case. It is a loss limiter, not a universal win over DRAFT=0.
• WarpDecode: replayed cleanly on current code, then isolated with eight real GLM-sized W4 experts. Exact output, but median kernel latency increased from 0.248 to 0.326 ms (+31.5% latency / -23.9% effective throughput). Not merged.
• Learned cross-layer prefetch: a legitimate 10-category coupling table produced only six useful nonresident hints; real loading did not improve I/O or decode. Rejected for the full-resident/near-resident regime.
• KV pages: physical CUDA KV pages at 16/32/64/128 tokens were all numerically exact across page boundaries. Page 64 remains the practical Pareto point; smaller pages reduce slack but do not improve steady decode.
• CPU tuning: 20 OpenMP threads lost 3.1% versus 24; separate IDOT execution made expert time 5.1% worse; both were rejected.
• Disaggregation: phase splitting is not deployable on the current topology because one full-resident instance already consumes the available 6-GPU/368 GB expert capacity. The break-even model is retained, but no runtime claim is made.

Updated conclusion

The remaining opportunity is real, but the objective must be measured directly:

minimize held-out end-to-end latency
subject to:
  full residency / bounded miss cost
  max(T_cpu, T_gpu)
  per-device group shape and load
  NUMA locality
  bounded worst-case regression
  explicit token/quality policy

The campaign rejects the shortcut objectives: maximize hit rate, maximize GPU expert count, balance routed-row count, or optimize a kernel in isolation. Future placement work needs category-held-out traces plus real per-device timing; future speculation work needs routed-union cost, not acceptance alone.

The useful outcome is not that the papers were "wrong." Their mechanisms are conditional. Our failures came from transferring results across workload, topology, and system boundaries without carrying those conditions. Fixed replay plus held-out evaluation is now the gate for any claimed default-path improvement.

[ZacharyZcR]

Fresh physical-memory datapoint from the same 6x RTX 5090 / dual Xeon Silver
4510 host, measured 2026-07-28.

The machine has 4x64 GB DDR5-4800 DIMMs configured at 4400 MT/s. A simple
8 GiB parallel sequential-read benchmark, first-touched under the tested NUMA
policy and using 24 physical cores, measured:

placement	threads	best read bandwidth
node 0 local	12	62.84 GB/s
node 1 local	12	62.98 GB/s
dual-node interleave	24	110.78 GB/s

The current Colibri routed-CPU path reaches 67-68 GB/s with selective NUMA
placement, versus 35-40 GB/s without it. Thus the expert kernel converts about
61% of the host's best simple-read bandwidth into useful INT4 expert traffic.
The remainder is not idle memory alone: nibble unpack, per-row scaling,
activation conversion, accumulation, and expert scheduling consume execution
resources.

This also identifies the hardware limitation precisely. The DIMMs provide only
four populated channels across two sockets; their aggregate theoretical
bandwidth at 4400 MT/s is 140.8 GB/s, so 110.8 GB/s is a plausible ~79% of
that physical ceiling. More cores on this exact DIMM population cannot create
the missing channels.

One current-runtime nuance: after NUMA is correct, the measured asynchronous
windows were CPU 2.839 s versus GPU critical 3.290 s for 64 fixed tokens. RAM
was no longer the straggler in that cell. More channels would improve
CPU-heavy routings and permit a larger RAM share, but should not be projected
as a linear end-to-end speedup without rebalancing the GPU tier.

For a RAM-participating next host, a useful procurement target is roughly
220-300 GB/s measured read bandwidth and 140-180 GB/s effective expert
bandwidth. Beyond that, the current INT4 kernel and GPU/attention path are
likely to take over unless core count and kernel structure move with the
memory channels.

[ZacharyZcR]

2026-07-31 closeout: the lossless single-stream software ceiling is now narrow

We completed the remaining paper-inspired and architecture-level checks on the
same 6x RTX 5090 / dual Xeon Silver 4510 / 251 GiB host. The conclusion is now
stronger than the opening estimate: for the original GLM-5.2 744B architecture,
lossless greedy single-stream decode, and this hardware, we are close to the
engine-only ceiling.

Current matched best point

• 8.90 tok/s, 64-token greedy decode
• 9,335 VRAM + 10,121 RAM experts
• 100% expert hit, zero disk traffic
• CPU expert bandwidth: 66.88 GB/s
• time split: approximately 44% expert matmul / 34% attention / 22% other

This is not a claim that 8.90 is a universal model ceiling. It is the best
controlled point for this model/topology/quality regime.

Newly closed directions

1. Post-training cross-expert sharing is absent in the real weights. Across
   layers 3/10/30/50/77 and all 256 experts, shared-mean energy was only
   0.38-0.69% (near the 1/256 random baseline), pairwise cosine means were near
   zero, and 90% energy required roughly 192-224 of 256 bases. A 35% storage
   reduction produced roughly 18-30% sampled reconstruction error. Permutation
   matching was indistinguishable from a random nearest-neighbour control.

2. SERE-style functional substitution does not transfer to GLM-5.2. Five
   diverse prompts, real hidden states, layers 3/30/77, all experts, and 489,600
   comparisons produced mean nearest similarities of
   0.08457/0.02133/0.07281. No expert pair reached 0.5.

3. MTP remains workload-conditional, not a safe default. On the latest
   full-resident explanatory prompt: DRAFT=0 8.90 tok/s; DRAFT=1
   7.51 tok/s (-15.6%, 60% acceptance, +30.4% CPU rows); DRAFT=3
   3.95 tok/s (-55.6%, 23% acceptance, +145% CPU rows). A predictable
   prompt can still win at about 91% acceptance, so the feature stays opt-in;
   acceptance and routed expert-union growth must both be controlled.

4. Existing continuous batching does not multiply aggregate throughput.
   Full-resident mux ABBA medians for 1/2/4/8 active sessions were
   4.84/6.33/8.14/8.30 aggregate tok/s. A short profiled 8-session run
   reached 9.05 tok/s, but CPU expert bandwidth fell to about 27-35 GB/s versus
   66.88 GB/s in the ordinary single-stream path. The CPU side was 5.6-11.1x
   the overlapping GPU critical path. The scheduler/ragged KV machinery is
   useful infrastructure, but the present multi-row INT4 CPU kernel is not a
   throughput win.

5. GPU direct host-memory access is already falsified here. The earlier
   dual-lane v2 in this issue measured only 1.71 tok/s / about 2.8 GB/s for
   demand-fetched zero-copy PCIe reads. Bulk pinned DMA can reach wire speed,
   but the 191 GiB warm store cannot be registered at useful coverage on this
   251 GiB host without RSS inflation or loss of NUMA interleave. This is not a
   remaining unexplored lever on this machine.

These results join the already rejected impact-only placement, learned
prefetch, WarpDecode, fused ANS consumption, CUDA-enabled XEXP, SMT, wider SIMD,
and extra-VRAM-occupancy variants documented above.

Engineering conclusion

The remaining engine-only opportunity is incremental rather than a hidden 2x:
roughly 10-25% aggregate headroom is a reasonable planning range. A stable
9.5-10.5 tok/s is a defensible target; 11-12 tok/s is an exceptional upper
target, not a measured claim.

A material improvement beyond that requires changing a boundary condition:

• more VRAM or substantially higher/coherent memory bandwidth;
• a stronger trained draft that also controls expert-union growth; or
• a retrained locality-aware MoE derivative that reduces RAM/cross-device
  expert bytes per token.

For the current engine, the highest-value work is now stability, automatic
hardware planning, and cross-machine portability. A new optimization proposal
should reopen performance research only if it reduces measured bytes per token
or removes a measured whole-layer critical-path cost under full-residency,
held-out, end-to-end evaluation.