Prefetch: capacity-mode readahead, plus the model and accuracy experiment

[SHTUSIST]

Two separate pieces of work, kept in one PR because the second is what validates the first. Reviewers can read them independently.

---

Part 1 — Capacity-mode readahead: HBFSim actually fetches ahead

This is a system-side prefetcher, not an application-side one. The application does not ask for it and does not know it happened; the only thing the policy has to go on is which pages have already been demanded. It is the analogue of OS page-cache readahead, not of cp.async.

On a demand miss for page N, CapacityPageService queues N+1..N+k; CapacityWorker fetches them only when a slot scan claimed nothing. Off unless switched on, so no existing measurement moves.

Three design decisions, and I got the second one wrong first

1. Queueing happens after the demand is served. The miss path is a synchronous blocking read under one mutex — reading a second page there would double the synchronous work on the path a warp is blocked on.

2. Readahead may take a clean victim. My first version refused to evict at all. Measurement killed it: on a 1024-access MoE stream with 128 frames it fetched 22 pages and skipped 1970, because a warm cache never has a free frame — the feature did nothing in the only regime capacity mode exists for. A dirty victim is still refused and put back, because writing it costs a program on the media, dearer than the read the readahead would save.

3. A readahead page is published unreferenced. HbmCache::publish gained a defaulted referenced parameter (every existing caller unchanged). The policy is CLOCK, so a page nothing ever asks for is now the first candidate to be evicted again rather than the last.

What it is worth

hbf_capacity_readahead_bench drives the real page service and reports media reads, not wall-clock time.

The result that matters most is a cliff. P=8, depth 8, 256 frames, varying how many queued pages the worker drains between two demands:

drain per demand	hit rate	media reads
fully	0.8916	−86.79%
8	0.8916	−86.79%
4	0.8916	−86.79%
2	0.8916	−86.79%
1	0.0635	+14.17%

Draining slower than the queue fills does not degrade gradually — it turns an 86.79% reduction into a 14.17% increase, because queued pages go stale before use while still costing bandwidth and frames. Any deployment must keep drain rate above fill rate.

---

Part 2 — The prefetch model and its accuracy experiment

The figure: x = prefetch accuracy, y = modeled time, one curve per media latency tR (1, 2, 4, 5, 10, 20 µs).

Accuracy alone does not decide the time — a correct prediction still arrives late if issued too near its use, if bandwidth is short, or if the staging buffer dropped it. The fix is more curves, not a different x axis: "how accurate" is the x axis; "did it arrive in time" is the spacing between curves and whether a curve still falls near accuracy 1.

Three policies: None (what HBFSim did), NextPage (no model of the workload), Accuracy (right with a stated probability — how the axis is swept without inventing a predictor). Four knobs, each a real limit: lead distance, staging capacity, media concurrency, and the compute between accesses a prefetch hides behind.

Headline: expert routing being unpredictable is about which expert, not how many bytes

One Qwen3-30B-A3B token = 48 layers × 8 experts × 2304 pages = 884,736 page accesses.

	demand stalls	modeled time	speedup
no prefetch	884,736	12,386,304,000 ns	1.00
naive next-page	374	6,194,648,000 ns	2.00

The 374 are the first page of each expert, against 48 × 8 = 384. The unpredictable part is 0.04% of the traffic. Speedup stops at 2.00 because there the media read is longer than the compute between accesses — bandwidth-bound, not accuracy-bound, which is exactly what the curve family shows.

---

The two parts cross-validate

The model predicts a next-page policy reaches (P−1)/P on an MoE stream. The real implementation reproduces it when the readahead depth covers one expert:

P (pages per expert)	depth	implementation	model (P−1)/P
4	4	0.7559	0.7500
8	8	0.8916	0.8750
16	16	0.9370	0.9375
16	32	0.9443	0.9375

Below that, depth caps the hit rate near D/(D+1): at depth 4, P of 8 and 16 both stop at 0.7500 and 0.7568.

Tests (TDD)

prefetch_model_test.cpp — 12 property assertions, written first. One was wrong on its first run: it demanded zero stall under a perfect prefetcher, and the model was right to disagree, because a prefetcher issuing L ahead cannot cover the first L accesses.

capacity_readahead_test.cpp — off by default; queue not drained on the demand's own path; dirty victim refused; a demanded page survives the readahead after it; a page that became resident between queueing and draining is skipped; a readahead past the end of the store is skipped, not failed.

What this is not

Not a measurement of any predictor, GPU, or device. Artifacts carry "disclaimer": "modeled, not measured on any device or GPU". The readahead experiment drives the host-side page service — media reads avoided, not wall-clock. The MoE stream is synthetic, not captured from a vLLM run. compute_ns_per_access is a parameter, not measured from a workload.

docs/46-预取实验设计.md carries the design; benchmarks/prefetch/README.md lists every file and the exact reproduction commands.

Test status

CPU-only build. Same four tests fail before and after — context_lifecycle, vmem_tuning, run_with_bpftime, mqsim_benchmark — all pre-existing. Count 32 → 34.

🤖 Generated with Claude Code

[zzyuanyi]

Blocking correctness issues found while reproducing the benchmarks. (1) hbf_capacity_readahead_bench labels demand misses as media_reads and computes media_reads_avoided_fraction from that value, but every successful readahead_pages_fetched is also a backing-store media read. Reproduced at P=8/depth=8/256 frames: drain=0 reports 111 reads and -86.79%, but the implementation performed 111 demand reads + 868 readahead reads = 979 versus the 840-read baseline (+16.55%). drain=1 reports 959 and +14.17%, but actually performed 959 + 825 = 1784 (+112.38%). This reverses the headline conclusion. (2) drain_per_demand is passed as an extra printf argument but has no JSON placeholder, producing -Wformat-extra-args and omitting the parameter that determines the result. (3) fill_free_frame_locked completes a clean eviction before the backing read/copy succeeds; an out-of-range or failed speculative read therefore discards a valid resident page, which test case 4 does not detect. (4) no production construction path calls set_readahead_pages; only tests and the benchmark can enable the feature. The five targeted tests pass, but they do not cover these failures. Please correct the accounting/artifacts, make speculative replacement rollback-safe, expose an intentional runtime configuration path, and regenerate the claimed results.

[SHTUSIST]

All four points hold. Each reproduced before fixing; pushed as 84ff0c0.

(1) Accounting — you are right, and it reverses my headline. I reproduced your numbers exactly against the 840-read baseline:

drain	demand	readahead	total	change
0	111	868	979	+16.55%
1	959	825	1784	+112.38%

The benchmark now prints demand_media_reads, readahead_media_reads, total_media_reads, demand_reads_avoided_fraction and total_media_reads_change_fraction, and the README carries both with the reading they require.

The honest conclusion is the more interesting one: the readahead removes 89.16% of the reads a warp waits on and still raises total media traffic by 16.55%. It trades bandwidth for latency — which is the wrong trade whenever the tier is bandwidth-bound, and the array supply rate is precisely what this project argues is binding. That belongs in the paper as a limitation of readahead on this tier, not as a win.

(2) Format. Confirmed with -Wformat-extra-args; my edit had silently failed to add the placeholder. Fixed, compiles clean under -Wall -Wformat -Wformat-extra-args, and drain_per_demand now appears in the JSON.

(3) Rollback safety. Confirmed. The read now happens first and a frame is taken only once the bytes are in hand. New case 4b covers it: four clean demanded pages filling every frame, readahead running off the end of the store — old ordering leaves resident=0111, new ordering leaves 1111. You were also right that case 4 could not catch it, because free frames were still available there.

(4) Production path. Confirmed — only tests and the benchmark could enable it. readahead_pages is now a profile field, optional and defaulting to 0 so no existing profile changes behaviour, parsed in profile.cpp, in the schema, and applied by CapacityRuntime's constructor. Setting it after the worker starts is safe because the queue only fills on a demand miss and no demand can arrive before a range is registered.

Thank you for reproducing the benchmark rather than reading it — (1) would have gone into the paper as a result.

The claimed results in docs/46-预取实现与预取实验说明.md still carry the old framing; I am correcting that document next.

[SHTUSIST]

You are right, and this is the more fundamental of the two rounds. I verified both halves against the code:

• fill_free_frame_locked produces no CapacityMediaPlan at all — it returns a bool and never touches the media field.
• The subsequent demand hit returns .media default-constructed, so CapacityMediaNone.

So the readahead performs a real backing read and a real host-to-device copy, and the simulated timeline never sees either. It contributes no modeled latency, no queueing, no contention, nothing to the thermal path, and nothing to the unified report; and the hit it enables reports zero media time. In simulation terms the speculative read has been deleted from the timeline, which is exactly the shape that would manufacture a favourable result.

I am therefore withdrawing the performance claim rather than patching around it. The traffic counts stand and reproduce — they are a statement about media reads issued, nothing more. docs/46-预取实现与预取实验说明.md will say that the readahead is not yet in the timing model and that no runtime benefit is claimed.

For the real fix I would rather agree the shape with you before writing it, because it is not local. The readahead has no originating request, so there is no slot or ticket to carry a CapacityMediaPlan back through request_dispatcher. The nearest precedent already in the tree is flush(model_program, range_id), which charges a program to the media model through a callback outside the request path. Mirroring that — a ModelRead callback the worker supplies, so a readahead's media read is charged the same way a writeback's program is — looks like the smallest change that puts speculative reads on the timeline and lets them contend with demand traffic.

Does that match how you would do it, or would you rather the readahead go through the normal request path with a synthetic request id? I will not regenerate any performance number until the speculative read is on the timeline.

[SHTUSIST]

Speculative reads are now on the timeline in 54386c8. This has never run against a live engine and I am asking you to test it before it is trusted.

What I chose, and why not the other two

I went with the synthetic-request shape rather than either option I floated earlier.

• A ModelRead callback mirroring flush(model_program, range_id) would follow a hook that has no production consumer today — I grepped, nothing supplies it on the production path — so it would add a second dead hook. Submitting outside the dispatcher also breaks the ticket_by_engine_id_ pairing.
• Carrying the reads on the next demand's spare media-action slot fits the existing std::array<HbfRequest, 2> for depth 1, but the dispatcher publishes only after all of a request's actions complete, so it puts the speculative read back on that demand's critical path — the thing the readahead exists to avoid.

So: submit_speculative queues a read; poll_once drains one only when it popped no demand and nothing is in flight; the read goes to the same engine as a demand, so it occupies the same channels and contends with demand traffic; publish discards its completion because no slot is behind it.

One bug I found by re-reading, already fixed

publish originally erased the speculative group itself. Both callers erase after publish returns, one of them through the very iterator used to reach the call — so erasing inside invalidated that iterator and the caller then erased through it. publish now leaves the erase to the callers.

Two things I could not verify — marked in the code as UNVERIFIED, PLEASE CHECK

1. Ticket namespace. Speculative tickets set bit 63 to stay clear of real tickets, which are the request's own sequence. I could not establish from the producer side that sequence never reaches 2^63. If it can, the namespace has to be separated another way — a flag on DispatchGroup rather than a bit in the ticket. Everything else holds either way.
2. Wrap. next_speculative_ticket_ wraps after 2^63 speculative reads. Not reachable in any run here, but nothing enforces it.

What the test status does and does not say

Suite is 88%, same four pre-existing failures (context_lifecycle, vmem_tuning, run_with_bpftime, mqsim_benchmark). That only says I did not break what was already there. There is no test for the speculative path itself, and no GPU on this machine, so nothing here exercises a live engine.

What I would most like you to check, in order: whether a speculative read can slip ahead of a demand whose next action is pending; whether the discarded completion leaves anything behind in ticket_by_engine_id_; and the bit-63 assumption above.

No performance number will be regenerated until this has run on your side. docs/46-...md and the D7 section both state that no run-time benefit is claimed.