lanes: readout requests for decision prompts by laudney · Pull Request #2047 · ashhart/TensorFold

laudney · 2026-11-03T16:31:02Z

Lane core, readout requests (Radeon AI PRO R9700)	Before	After
A prompt-only request that reads its rows (a decision)	not a request kind	`readout` requests: prompt rows in pieces under the shared fill policy, then a family completion
Readout rows alone, 4 at once beside a decoding stream, after a cancel (pieces of 1,024)	—	byte-identical in all three, at 16,384 rows (16 pieces) and at 3,000 rows (3 pieces)
Two short embeddings on a `prefill_many` backend	batched admission took both into one `addStreams` pass and never called their completion, so neither finished	every prompt-only job goes through `launch`, which fills and completes it (`fix(lanes)`, first commit; a fake reproduces the hang)
Other routes on a readout engine	—	refused, naming `POST /v1/systemone`

Stacked on #2046 (the BF16 backbone). Part of #1360. This answers the #1361 review ask that decision prompts join the shared rounds instead of a side path.

What

fix(lanes): launch prompt-only jobs alone: on a backend with prefill_many and many_rows above zero, batched admission took two short embeddings into one addStreams pass and delivered them without their completion, so neither request finished. It can land on its own.
feat(lanes): add readout requests: a readout Spec on the request (question and option spans), validated in the lane core, admitted like embeddings, filled in max_rows pieces under fill.run, cancelled like any job, completed by the family when the last piece is in.
feat(server): share the prompt-only wait: embeddings and readouts share one submit-and-wait helper.
feat(qwen3_5) / feat(hip): the HIP Readout interface (open / consume / finish / close), readout prompts through the shared prefill step, and the engine's readouts in its Info (text_generation is off when they are on).
test(qwen35): the gate 2a readout rows check (alone, together, after a cancel).

Only a fake Readout is wired here; the joint schema head that implements it comes in the next PR.

Checks: zig build test -Dhip=true on gfx1201 (ROCm 10.1), zig fmt --check, lean checker at the baseline.

An unquantized HF checkpoint keeps its projections in BF16, and the HIP engine will multiply them as stored through rocBLAS, as torch's Linear does. This adds the bf16 format module (detected only when a config declares no quantization, rank-2 BF16 weights read as mapped) and its tag in the format, decoder and union types, with a refusing arm in every exhaustive switch so the tree builds. The shared read order is unchanged, so no other family sees the new format.

Clef-Flash ships as an HF BF16 checkpoint, which the family must read with the reference's semantics. The family now detects the bf16 format before the shared formats, maps HF's model.language_model and lm_head names, refuses a mixed dtype by name, reads the conv weight (C, 1, K) as (C, K), and loads the trunk, final, q and k norms as 1 + w in fp32 by role (not by the mean heuristic), leaving linear_attn.norm as stored. The shard list of a bf16 checkpoint comes from its index, so the joint head stays out, and textBytes counts only the text tower and head.

A bf16 Qwen3.5 checkpoint keeps its projections as stored, so the HIP engine multiplies them through rocBLAS on its own stream, BF16 in and out with FP32 sums, in a planned 32 MiB workspace that enters planBytes. tf_embed_dense gains a table kind, read through load(), so the gather takes a BF16 table byte for byte; its kernel test, its contract rig (BF16 table, fp32 table bit-equal to the old gather, two planted controls) and its launch record come with it. The engine refuses bf16 with fp16 activations or more than one rank, turns round graphs off, counts only the text tower before loading, and the server refuses a bf16 checkpoint as decision-only until the readout slices land.

Gate 1 compares every final-norm row of a prompt against transformers, and C2's readout pieces need the rows of each piece. Engine.spanRows runs a span like advance but keeps fwd.span's rows, and advance now calls it. tf-qwen35-test hidden writes every prompt row's final-norm row through spanRows, piece by piece at a given cut, as a BF16 .npy; --layers instead writes each layer's residual rows, to find which layers an error grows in.

Gate 1 needs transformers' rows on the same records as the native run. clef_backbone_rows.py writes the 40 seeded synthetic records (10 per length bucket), runs the text model in BF16 on the GPU and in FP32 on the CPU, drives tf-qwen35-test hidden, and reports per bucket the pooled mean absolute error against FP32 for native and BF16, and every teacher-forced last-row top-1 difference with its margins; a non-finite reference voids the run. SOURCES pins torch, transformers and the checkpoint revision.

On a backend with prefill_many and many_rows above zero, the batched admission took two short embeddings into one addStreams pass and then delivered them; nothing called the embedding completion, so neither request ever finished. A prompt-only stream (an embedding for now) now always goes through launch, which fills it and completes it, and a test on such a fake reproduces the hang without the fix.

A readout is a non-generating request whose answer, BF16 logits per option, the backend computes from the prompt's final-norm rows. The lane core gets the Spec and Info types with their validation, a Request and Stream field, one backend hook, readout_finish, and a completion that mirrors the embedding one: logits preset to NaN, non-finite output refused, one release on success, failure or cancel. Embeddings and readouts now share a prompt-only predicate, so both skip the warm pass, the lone driver, batched admission, the prompt cache, the repeat log and the Recall record. An engine that reports readouts refuses every other request.

A readout engine needs the same borrowed-storage wait as embeddings: submit, poll for the terminal event, cancel once when the client is gone and keep waiting until the engine is done with the request. That loop and its sink move from the embedding route into one module both use, with no new timeout. The server now refuses to start an engine that reports readouts beside text generation or embeddings, and such an engine answers every POST route except System One and tokenize with a 400 that names POST /v1/systemone.

Gate 2 compares readout rows across runs that must cut the prompt at the same rows, and the lane rule picks 1,024 or the policy's fill rows depending on what else decodes. Options.cut_rows fixes the piece size for tests and checks (0 keeps the rule), and the engine grows prompt_rows to hold one such piece, so a whole-prompt cut sizes its prompts arena for it.

The lane core calls readout_finish only, so the family owns a readout stream's state: a Readout interface (open, consume, finish, close) that C3's joint schema head will implement, held by the lane backend. begin opens a stream's state and free closes it, so every release path frees it; a readout starts at row 0 with no cuts and never asks the prefix tree. readoutPiece feeds each piece's final-norm rows from spanRows to the state, and reach honours the test-only cut, which a recorder can list. No implementation ships, so the engine still reports no readouts.

A readout stream's answer comes from its rows, not from a token, so the shared prompt step sends each of its pieces, the last one too, to the family's readoutPiece: the last piece runs no LM-head projection, draws nothing and copies no last row, then ends the fill. A cut on the way keeps no state. The vtable gains readout_finish for a family that declares readoutFinish, and both are checked with hasDecl, so Nemotron and Flash Next are unchanged.

The generic HIP adapter builds one lane host Info for Qwen3.5, Nemotron and Flash Next. It now takes readouts from the family's Loaded when that type has the field, and turns text generation off only when they are set, so every other family keeps its Info, which a host test checks for each registered family. Qwen3.5's Loaded reports what its lane backend computes, still none, and the backend refuses a readout on more than one rank, since a follower's final piece projects.

Gate 2a for C2 needs the rows a readout consumes to match whatever runs beside it. tf-qwen35-test readout installs a test-only readout that downloads every piece's final-norm rows, then runs one readout alone, four together with a piece of each in turn, one released after its first piece, and one after that, all at the same fixed cut. It fails unless the pieces cover the prompt from row 0 at that cut and every copy and its pieces equal the first; an error still releases each open lane while the readout's states live. It opens the engine directly, so the server's BF16 refusal stays.

ashhart · 2026-11-03T16:59:28Z

Thanks @laudney. This is a clean middle layer. Every shared lane change is gated on the request kind, so text generation on every existing family takes the same path as before. The HIP Info test that checks every registered family is a nice touch.

Before this lands, could you:

Post the device output behind the "byte-identical in all three" row. That's the exact tf-qwen35-test readout command, the card and ROCm version, the PASS line, and the zig build test -Dhip=true summary.
Pin a readout's piece size in reach (always the 1,024-row chunk unless the test cut is set). Add a host test for it, and add a device run with the serving rule, alone and beside others. rocBLAS picks its BF16 kernel by row count, so today's 1,024 rests on event ordering, and a row's bits could move with it. Readouts never resuming from a kept state is already right and tested. Thanks for that.
Match the PR text's description of the embedding hang to commit 6 (batched admission on a prefill_many backend).

The embedding fix in the first commit can land on its own whenever you like.

rocBLAS picks its BF16 kernel by row count, so a readout's rows could change bits if its pieces followed the lane rule, which shortens a prompt's pieces to the policy's fill rows while another stream decodes. reach now gives a readout the 1,024-row chunk whatever the other lanes do, unless the test cut is set; a host test checks this beside a decoding lane, and tf-qwen35-test readout takes --cut 0 to run the serving rule, with its together phase now beside a stream that decodes.

laudney · 2026-11-03T17:42:01Z

Ask	Answer
1. Device output behind "byte-identical in all three"	below: R9700 (gfx1201), ROCm 10.1.0, hipcc HIP 7.16; PASS, rows wyhash `f712c7554f2127dc` alone, 4 together and after a cancel; `zig build test -Dhip=true`: 526/526 steps
2. Pin a readout's pieces in `reach`	`5b9fa01` `fix(lanes): pin a readout's pieces to the prompt chunk`: a readout keeps the 1,024-row chunk unless the test cut is set; host test in `hip_prompt_test.zig` (a plain prompt beside a decoding lane gets `fill_rows`, a readout gets the chunk). The device check now runs the serving rule (`--cut 0`) beside a decoding stream, and with the pin removed it fails (control below), so the pin is what keeps the bits
3. The embedding hang text	the PR text now says what the commit fixes: batched admission on a `prefill_many` backend took two short embeddings into one `addStreams` pass and never called their completion

Device receipt

# PR #2047: readout pieces pinned to the prompt chunk (fix(lanes): pin a readout's pieces to the prompt chunk)
Card: AMD Radeon AI PRO R9700 (gfx1201). ROCm 10.1.0, hipcc HIP 7.16.26385 (core/bin). Build: zig build tf-qwen35-test -Dhip=true -Dhip-arch=gfx1201.

## The serving rule (no test cut): 1 alone, 4 together beside a decoding stream, 1 cancelled after its first piece, 1 after
$ tf-qwen35-test readout <Cloudflare/clef-flash@fb4417d3> --len 3000 --cut 0 --streams 4
policy matrix=auto,attention=auto,kernels=auto,graphs=auto,prefill_step=1024,fill_rows=512; mtp drafts 3 under 0.3, 8 snapshots 
info: HIP memory budget from the card's memory: 12.9 of 29.9 GiB free, 1.0 GiB reserve
readout gate 2a: 3000 rows x 8192 bytes, the serving rule, piece 1024, pieces { { 0, 1024 }, { 1024, 1024 }, { 2048, 952 } }
PASS 1 vs 4 together beside a decoding stream and after a cancel, the serving rule: rows byte-identical (wyhash f712c7554f2127dc), 20 cuts recorded
rc=0

## The same with the test cut 1024 (the earlier gate 2a run, now beside a decoding stream too)
$ tf-qwen35-test readout <Cloudflare/clef-flash@fb4417d3> --len 3000 --cut 1024 --streams 4
policy matrix=auto,attention=auto,kernels=auto,graphs=auto,prefill_step=1024,fill_rows=512; mtp drafts 3 under 0.3, 8 snapshots 
info: HIP memory budget from the card's memory: 12.9 of 29.9 GiB free, 1.0 GiB reserve
readout gate 2a: 3000 rows x 8192 bytes, a test cut, piece 1024, pieces { { 0, 1024 }, { 1024, 1024 }, { 2048, 952 } }
PASS 1 vs 4 together beside a decoding stream and after a cancel, a test cut: rows byte-identical (wyhash f712c7554f2127dc), 20 cuts recorded
rc=0

## Control: the serving rule with the pin removed (a readout follows the decoding rule): it fails
$ tf-qwen35-test readout <Cloudflare/clef-flash@fb4417d3> --len 3000 --cut 0 --streams 4
policy matrix=auto,attention=auto,kernels=auto,graphs=auto,prefill_step=1024,fill_rows=512; mtp drafts 3 under 0.3, 8 snapshots 
info: HIP memory budget from the card's memory: 12.9 of 29.9 GiB free, 1.0 GiB reserve
readout gate 2a: 3000 rows x 8192 bytes, the serving rule, piece 1024, pieces { { 0, 1024 }, { 1024, 1024 }, { 2048, 952 } }
FAIL readout 1 of 4 together beside a decoding stream differs from the one alone
FAIL readout 2 of 4 together beside a decoding stream differs from the one alone
FAIL readout 3 of 4 together beside a decoding stream differs from the one alone
FAIL readout 4 of 4 together beside a decoding stream differs from the one alone
FAIL RowsDiffer
rc=1

## zig build test -Dhip=true --summary line
Build Summary: 526/526 steps succeeded; 288/294 tests passed (6 skipped)
(zig counts only the test binaries it ran again; the run before this one, with one compile error, rebuilt the suite: 6869/6918 tests passed, 49 skipped)

Blas.open resolved rocblas_set_workspace and the two workspace queries for every caller, so a loader that never sets a workspace, such as GLM-OCR, needed a rocBLAS that has them. The three calls now form their own table, resolved on the first useWorkspace or workspace call, which only the bf16 Qwen3.5 path makes; a rocBLAS without one is refused there by the call's name. A host test checks that each missing call is named and that open no longer asks for them.

ashhart · 2026-11-03T18:18:58Z

Thanks @laudney, that's everything we asked for. The pin in reach plus the control run, where the together readouts differ once the pin is removed, show the pin is what keeps the bits. Moving the rocBLAS workspace calls to the BF16 path also answers the loader point on #2046.

One small thing: the body's table says "16K prompt", but the posted runs are 3,000 rows. Could you say which length the row covers, or add the 16K run? We don't have an AMD GPU, so your gfx1201 readout check and control are the device evidence. This can land after #2046.

laudney · 2026-11-03T18:37:03Z

The table row now names both lengths. The 16K run, on the same build as the receipt above (branch tip 2f0d57326, R9700 gfx1201, ROCm 10.1):

$ tf-qwen35-test readout <Cloudflare/clef-flash@fb4417d3> --len 16384 --cut 0 --streams 4
policy matrix=auto,attention=auto,kernels=auto,graphs=auto,prefill_step=1024,fill_rows=512; mtp drafts 3 under 0.3, 8 snapshots 
info: HIP memory budget from the card's memory: 12.9 of 29.9 GiB free, 1.0 GiB reserve
readout gate 2a: 16384 rows x 8192 bytes, the serving rule, piece 1024, pieces { { 0, 1024 }, { 1024, 1024 }, { 2048, 1024 }, { 3072, 1024 }, { 4096, 1024 }, { 5120, 1024 }, { 6144, 1024 }, { 7168, 1024 }, { 8192, 1024 }, { 9216, 1024 }, { 10240, 1024 }, { 11264, 1024 }, { 12288, 1024 }, { 13312, 1024 }, { 14336, 1024 }, { 15360, 1024 } }
PASS 1 vs 4 together beside a decoding stream and after a cancel, the serving rule: rows byte-identical (wyhash 5492e4d9425b3977), 98 cuts recorded
rc=0

Main's loaded-group receipts (100363c) set `benchmark_ranks` in the HIP lane host's Info literal, which this branch had moved into `hostInfo`; the merge passes the group's world size to `hostInfo` as a parameter so the whole Info stays in one function, and the Info test now also checks a two-rank group.

laudney · 2026-11-03T22:37:33Z

Merged main in cf13b0b: main's loaded-group receipts (100363c) set benchmark_ranks in the HIP lane host's Info, which this branch builds in hostInfo, so hostInfo now takes the group's world size and the Info test also checks two ranks. Lean against main: 0 new; zig build test -Dhip=true on the R9700 (ROCm 10.1): 531/531 steps.

ashhart · 2026-11-03T22:56:47Z

Thanks @laudney. The 16K run and the table naming both lengths close the last point, and the rocBLAS workspace commit now sits in #2046 where it belongs. Your gfx1201 receipts are the device evidence here, since we have no AMD GPU. One small thing for the rebase onto #2046: the hostInfo change from merging main lives in the merge commit, so could you rebase rather than merge, so the series itself carries it? Then this lands right after #2046.

laudney added 13 commits November 3, 2026 16:14

laudney mentioned this pull request Nov 3, 2026

qwen3_5: serve Clef-Flash's BF16 Qwen3.5 backbone on HIP #2046

Open

Navigation Menu

Repository navigation

lanes: readout requests for decision prompts - #2047

lanes: readout requests for decision prompts#2047
laudney wants to merge 16 commits into
ashhart:mainfrom
laudney:pr-clef-2-readouts

laudney commented Nov 3, 2026 •

edited

Loading

Uh oh!

ashhart commented Nov 3, 2026

Uh oh!

laudney commented Nov 3, 2026

Uh oh!

ashhart commented Nov 3, 2026

Uh oh!

laudney commented Nov 3, 2026

Uh oh!

laudney commented Nov 3, 2026

Uh oh!

ashhart commented Nov 3, 2026

Uh oh!

Reviewers

Uh oh!

Assignees

Labels

Projects

Milestone

Development

Uh oh!

2 participants

Conversation

laudney commented Nov 3, 2026 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

What

Uh oh!

ashhart commented Nov 3, 2026

Uh oh!

laudney commented Nov 3, 2026

Uh oh!

ashhart commented Nov 3, 2026

Uh oh!

laudney commented Nov 3, 2026

Uh oh!

laudney commented Nov 3, 2026

Uh oh!

ashhart commented Nov 3, 2026

Uh oh!

Reviewers

Uh oh!

Assignees

Labels

Projects

Milestone

Development

Uh oh!

2 participants

laudney commented Nov 3, 2026 •

edited

Loading