Skip to content
Closed
Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
36 commits
Select commit Hold shift + click to select a range
0a123fd
fix(hip): zero masked V per row in the wide prompt tile
laudney Nov 6, 2026
b6a612c
feat(hip): name the chips whose dot2 rounds a -0 lane apart
laudney Nov 6, 2026
e5d49e2
fix(hip): keep main's prompt tile off RDNA 3.5, with an edge pass
laudney Nov 6, 2026
6add094
test(hip): contract entries for non-finite values past a row's edge
laudney Nov 6, 2026
4bea033
changelog: #2357 under Unreleased
ashhart Nov 7, 2026
ed7542f
qwen27 metal: a keep that fails part way ends its encoder, so the err…
jschmied Nov 6, 2026
47bcd15
glm cuda: test-glm-snapshots-host's copies module imports core, which…
Bizuayeu Nov 6, 2026
b079937
changelog: #2403 under Unreleased
ashhart Nov 7, 2026
c4dc04b
flashnext cuda: the grouped draft head's lanes carry their page table
Arminova Nov 2, 2026
3147b9c
changelog: #2244 under Unreleased
ashhart Nov 7, 2026
d9b34bb
glm cuda: a prompt step takes two chunks while no other stream decode…
Bizuayeu Oct 26, 2026
97c7d40
glm cuda: a paired step takes rank 0's indexer split for both chunks …
Bizuayeu Oct 27, 2026
951ddd4
glm cuda: glm-split-check takes the pair's Leader.chunk, and a fifth …
Bizuayeu Oct 27, 2026
a32a1a2
glm cuda: a step pairs only where its index rings rebase as a chunk a…
Bizuayeu Oct 27, 2026
7649d99
glm cuda: only an engine with other ranks makes the pair's twin
Bizuayeu Oct 27, 2026
2228fd3
glm cuda: one table gives the twin's fields their size indexes, and a…
Bizuayeu Oct 27, 2026
34cd3c0
glm cuda: glm-split-check's list of prompts is described on one line
Bizuayeu Oct 27, 2026
8cf53e4
glm cuda: the pair's choice and its twin's pieces sit with the pair a…
Bizuayeu Nov 1, 2026
7aa3360
glm cuda: TF_GLM_CHUNK_SOLO=4096 takes the pair's place, and its serv…
Bizuayeu Nov 1, 2026
6d69c5d
docs: the ledger's item 20 lists GLM's paired prompt steps, about 2 %…
Bizuayeu Nov 6, 2026
8026e45
changelog: #1423 under Unreleased
Bizuayeu Nov 6, 2026
1c22fa3
gemma4 host: pin affine_mm.metal's bytes after cf10a6b7d's comment edit
ashhart Nov 7, 2026
a569b24
build: the compile-only Metal contract runners import masked_sources too
ashhart Nov 7, 2026
bc4506b
build: compile-metal-residency-host's native module imports kimi_k3
ashhart Nov 7, 2026
13bdba9
glm metal: the backend fact tests give facts() a host device object
ashhart Nov 7, 2026
f29e504
changelog: #2409 under Unreleased
ashhart Nov 7, 2026
8539b77
cli: run the pull tests' fake hub on std.Io.net so Windows builds them
ElderOrb Nov 5, 2026
c9713f0
ci: drop the workflow's byte-order mark and note the Windows pull tests
ElderOrb Nov 5, 2026
f275d99
changelog: #2284 under Unreleased
ashhart Nov 6, 2026
39fd1ab
build: the DeepSeek server gets the lanes module its /v1/systemone ro…
Defilan Nov 6, 2026
bcd8461
changelog: #2412 under Unreleased
ashhart Nov 7, 2026
91377a4
cli: tensorfold models --json, the listed checkpoints as one document
Whq02 Nov 6, 2026
1fb70c1
engine api: an optional image capability beside speech (#1573)
devopsfortunato Nov 5, 2026
9c3dbb2
moss_image: the transformer gated block by block against float32, on …
BobClawblaw Nov 6, 2026
91fdbe8
changelog: #2294 under Unreleased
BobClawblaw Nov 6, 2026
bee665d
build: the MOSS-Image CUDA tools install from moss_image_tools.zig, s…
ashhart Nov 7, 2026
File filter

Filter by extension

Filter by extension


Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
3 changes: 3 additions & 0 deletions .github/workflows/no-gpu.yml
Original file line number Diff line number Diff line change
Expand Up @@ -176,6 +176,9 @@ jobs:
run: zig build -Dfixed-cpu test-golden
- name: zig build fuzz-cortheon
run: zig build -Dfixed-cpu fuzz-cortheon
# The pull client's framing and path checks: its fake hub runs on std.Io.net, so Windows builds and runs them.
- name: zig build test-pull-security
run: zig build -Dfixed-cpu test-pull-security
- name: Save the build cache
if: ${{ !cancelled() && github.event_name != 'pull_request' && github.ref == 'refs/heads/main' }}
uses: actions/cache/save@0057852bfaa89a56745cba8c7296529d2fc39830 # v4.3.0
Expand Down
19 changes: 19 additions & 0 deletions CHANGELOG.md
Original file line number Diff line number Diff line change
Expand Up @@ -4,13 +4,24 @@ Each release's page on GitHub has its notes and measurements. See [the 1.0.0 rel

## Unreleased

- Build: `zig build compile-dsv41-server` and `dsv41-serve` build again: the DeepSeek server module gets the `lanes` module its `/v1/systemone` route reads (#2412). Thanks to @Defilan.

- Tests: four host steps that no CI job runs build and pass again: `test-gemma4-affine-host` pins `affine_mm.metal`'s current bytes, the compile-only Metal contract runners import `masked_sources`, `compile-metal-residency-host` imports `kimi_k3`, and the GLM Metal fact tests give `facts()` a host device object instead of an undefined one, which crashed them at random (#2409, part of #2031).

- Flash Next on CUDA with `TF_FLASHNEXT_KV_PAGES`: a shared round's grouped draft head reads and writes each stream's own pages, so paged streams' replies equal their runs alone again, and paged decode runs as fast as unpaged; default settings are unchanged (#2244). Thanks to @Arminova.

- Tests: `zig build test-glm-snapshots-host` builds again: its copies module imports `core`, which `cuda_snaps.zig` reads for `TF_GLM_PROMPT_TAIL` since #2226 (#2403). Thanks to @Bizuayeu.

- HIP: on RDNA 3.5 (gfx1150 to gfx1153) a prompt row gets the same bytes wherever its chunk starts. The 64-row prompt tile there now takes each key past a row's causal edge as +0, as RDNA 3.5's dot2 rounds a -0 lane apart from +0, so its prompt bits change, and gfx1151's prefill is up to 0.7% slower at 2k and 8k. Other chips keep main's prompt tile and bits, and gfx1201's prompt speed is main's within noise. On every chip, a non-finite value at a key past a row's edge now adds nothing to that row (#2357). Thanks to @laudney.

- GLM CUDA at TP=2 and up: each rank names its startup phases (NCCL init, peer warm-up, weights, calibration, ready), prints NCCL's init device, version and log path, and warns once a minute while NCCL init or the warm-up is unfinished, with the transport seen so far. It only reports; transfers are unchanged (#2354, part of #2347).

- Docs: the serve flags match the CLI: `--metrics-open` opens `/v1/metrics` too, `--learn-conversations` on CUDA serves the Qwen3.5 hybrids and Flash Next as well as Nemotron, Flash Next shares sessions under an explicit `--parallel` above one and refuses `--recall` only on a Mac, and the Metal Qwen3.5-2B keeps prompt states (#2367).

- Internal: the kernel contract's nvcc 13.0 resource records match what nvcc 13.0.88 builds for 57 CUDA kernels (226 per-SM slots), so the five-SM limits gate passes on main; no kernel, flag or launch changes (#2390).

- Flash Next on CUDA: a prompt-cache mark inside a prompt chunk is kept on the way through the pass instead of cutting the pass there, so a new short prompt's first token comes 19% sooner on a GB10, with replies unchanged; `TF_FLASHNEXT_KEEP_INPASS=0` cuts the pass at the mark as before (#1866). Thanks to @Arminova.
- Tests: `tf-moss-image-blocks` gates MOSS-Image's transformer block by block against the publisher's float32 run, within twice the spread of torch bf16's device and CPU runs, on Metal and CUDA (#2294, part of #1529).

- Cluster: MCDMA bring-up on real hosts, step 2a of #1594 (#2088). Thanks to @ThinkOffApp.

Expand Down Expand Up @@ -59,6 +70,9 @@ Each release's page on GitHub has its notes and measurements. See [the 1.0.0 rel
- Internal: 9 more lean findings fixed in Metal kernel and server comments, with no change in behaviour (#2329).

- Native DeepSeek V4 and V4.1 prompts follow the recorded encoder outputs, except where those outputs drop a tool's namespace, which we keep. Assistant thinking markers, blank-line text joins and tool schemas change; minimal keeps thinking on at budget 50 and xhigh uses 75. Developer instructions render as system, strict tools keep their existing grammar handling, listings and tool-call history keep qualified function names, and schemas retain namespace descriptions. Invalid user null content, malformed tool-call lists or calls without string id/name/arguments or type function, names outside 1-64 ASCII letters/digits/underscore/hyphen, tool replies without tool_call_id and non-function tool envelopes now return HTTP 400. Ultra keeps its existing refusal because no recorded case establishes a budget (#2141).
- GLM-5.3-Flash on CUDA: a prompt step takes two chunks while no other stream decodes, each one's exchanges under the
other's sites, so prompts of 8K to 128K finish about 2% sooner at TP=2, with replies unchanged;
`TF_GLM_PROMPT_PAIR=0` turns it off (#1423).

- TENSORFOLD_PHASE_METRICS=1 adds a bounded request queue-wait histogram and its vLLM alias to native metrics; it is off by default (#2139).

Expand All @@ -79,11 +93,13 @@ Each release's page on GitHub has its notes and measurements. See [the 1.0.0 rel
- Internal: 42 more lean findings fixed in the tests' comments and the tools' docstrings, with no change in behaviour (#2315).

- HIP: the decode round's residual-add and norm tail issues each row's loads before its branches, and a MoE tail loads its slots' rows together, with output bytes unchanged. On a gfx1151, one-stream decode is 2.7% faster on Qwen3.6-35B-A3B and 1.1% on the 27B. `tf_tail` is now under the kernel contract (#2247). Thanks to @AbrahamCete.
- Qwen3.8-27B on Metal: a round's keep that fails part way (an allocation, a refused keep or path) returns its error instead of aborting the process on an unended command encoder, and after a failed allocation the round can keep again; `tf-qwen27-keep-error` fails each of the keep's allocations in turn and refuses a path after the encoder has ended (#2376). Thanks to @jschmied.

- Docs: `docs/optimization-ledger.md` lists which speed and memory optimizations apply to which native family, and where each stands, with links to the PRs that measured them (#2232). Thanks to @jschmied.

- Docs: the pull request pre-flight asks a PR that adds, finishes or drops a family's optimization to update that family's cell in the optimization ledger (#2243). Thanks to @jschmied.
- Server: a refused `--api-key-file` logs the path it was given, the mode the file system reported and which test failed, once at startup and once when a reload is refused (#2365). Thanks to @Whq02.
- CLI: `tensorfold models --json` prints the cached checkpoints as one document (`schema` 1, each checkpoint's `repo`, `family` and `model_type`); the lines are unchanged (#2364). Thanks to @Whq02.

- Server: tool calls completed just before the thinking-close marker are returned when no answer follows, whole and
streamed; calls followed by more reasoning or an answer stay in reasoning (#2313, fixes #2310). Thanks to @AdrianBinDC.
Expand Down Expand Up @@ -126,6 +142,8 @@ Each release's page on GitHub has its notes and measurements. See [the 1.0.0 rel
- Qwen3.8-27B and the other dense Qwen3.5 models on CUDA: streams run on KV pool pages of 256 positions by default, the engine's own stream too, so `--parallel` is now a cap rather than a reservation: a request can wait for pages instead of every stream holding a full-context window. `TF_QWEN_KV_PAGES=0` keeps a window each, as before. Qwen3.6-35B-A3B and the other MoE checkpoints stay on windows by default, with pages opt-in (`TF_QWEN_KV_PAGES=256`). Paged prompt attention stages a chunk's page ids in shared memory, a decode thread looks its page up once a page, and a paged prompt-cache save or restore copies one 2D block a page. Replies are the same tokens on pages and windows (#2011). Thanks to @jschmied.
- Qwen3.5 and Qwen3.6 on CUDA: KV pool pages stay opt-in (`TF_QWEN_KV_PAGES=256`), dense and MoE alike; unset keeps a window each, as before, and `TF_QWEN_KV_PAGES_ALL` unset still leaves the engine's own stream a window. With pages set, `--parallel` is a cap rather than a reservation: a request can wait for pages instead of every stream holding a full-context window. Paged prompt attention stages a chunk's page ids in shared memory, a decode thread looks its page up once a page, and a paged prompt-cache save or restore copies one 2D block a page. Replies are the same tokens on pages and windows (#2011). Thanks to @jschmied.
- Qwen3.5 and Qwen3.6 on CUDA: KV pool pages stay opt-in (`TF_QWEN_KV_PAGES=256`), dense and MoE alike; unset keeps a window each, as before, and `TF_QWEN_KV_PAGES_ALL` unset still leaves the engine's own stream a window. With pages set, `--parallel` is a cap rather than a reservation: a request can wait for pages instead of every stream holding a full-context window. Paged prompt attention stages a chunk's page ids in shared memory, a decode thread looks its page up once a page, and a paged prompt-cache save or restore copies one 2D block a page. Replies are the same tokens on pages and windows (#2011, part of #1943). Thanks to @jschmied.
- Tests: the pull client's security tests build and run on Windows, on a portable fake hub, and CI runs them there
(#2284). Thanks to @ElderOrb.

- Metal: a request's finished stats list its prompt pass widths (`tensorfold.prefill_widths`, one entry per completed
target pass) on every Metal text family except Flash Next, whose engine reports none. Telemetry only (#2293).
Expand Down Expand Up @@ -238,6 +256,7 @@ Each release's page on GitHub has its notes and measurements. See [the 1.0.0 rel
- Flash Next on Metal reports a filling prompt's progress: the dashboard shows what's left, and `TENSORFOLD_PREFILL_ORDER=srpt` works from a server's first long prompt (#2194).

- CI: `zig` and `zig-windows` keep main's build cache instead of rebuilding from nothing, every check still reruns, and the quick checks run first.
- The engine API carries an optional image capability beside speech: `Info.image` states what an engine's image host serves and a defaulted `VTable.submit_image` hands back one encoded PNG a request through the caller's own list and allocator, so an engine without the entry is unchanged (#1573).

- HIP: the `tf_sum_partials` records follow ROCm 10.1's hipcc, taking 8 limit problems off the contract (#2233). Thanks to @AbrahamCete.

Expand Down
2 changes: 2 additions & 0 deletions README.md
Original file line number Diff line number Diff line change
Expand Up @@ -234,12 +234,14 @@ On Metal, each of up to 128 streams in the 27B's shared rounds drafts its own DF
```sh
tensorfold pull TensorFold/NVIDIA-Nemotron-3.5-Lightning-30B-A3B-MLX-4bit
tensorfold models
tensorfold models --json
tensorfold info TensorFold/NVIDIA-Nemotron-3.5-Lightning-30B-A3B-MLX-4bit
```

`pull` downloads a checkpoint into the Hugging Face cache and checks what it downloads against the Hub: large files by sha256, small files by their git blob sha1.
Large files come down in pieces over up to 16 connections, and an interrupted pull resumes with the pieces it lacks.
It refuses a checkpoint whose `model_type` has no native family, and fetches the 27B's DFlash2 drafter for `--drafter`. `models` lists the cached checkpoints with a model type it knows, and `info` shows one checkpoint's family, format, context and memory floor.
`models --json` prints the same list as one document for programs: `schema` 1, and each checkpoint's `repo`, `family` and `model_type`; `family` is a display title and `model_type` is the key to match on.
`tensorfold serve REPO` serves a cached checkpoint by its repository name.

## Context compaction
Expand Down
3 changes: 1 addition & 2 deletions docs/kernel-contract.md
Original file line number Diff line number Diff line change
Expand Up @@ -133,8 +133,7 @@ Python family's Triton glue it was ported from. Nemotron's Metal `tf_xup_w3` giv
the NaN: its relu2 is `metal::max(bfloat(acc), 0.0f)`; the other up pipelines give its bytes on the same inputs.
HIP's prompt attention tile (`tf_fa_paged_*`, `tf_fa_wide_*`) writes 0 for a row whose every score is NaN where
float64 gives NaN, as its running max is `fmaxf` and a row at max -inf takes p = 0; and an infinite value at a key
whose fp32 weight is 0 (past the row's causal edge in its 64-key tile, or under fp32's least) gives NaN where float64
gives inf or a finite value, as its P x V takes every key of the tile.
the row sees under fp32's least weight gives NaN where float64 gives inf, as 0 x inf enters its P x V.
Qwen3.8 Flash Next's `tf_sattn_nax` and `tf_attn256_nax` score masked keys -FLT_MAX, the value each row's running
max starts from, so a row whose every score is -Inf weighs its masked keys 1 and gives a finite average where float64
gives NaN. `tf_sattn_nax`'s padding lanes also read key 0 at p = 0, so a non-finite value in key 0's row reaches every
Expand Down
2 changes: 2 additions & 0 deletions docs/optimization-ledger.md
Original file line number Diff line number Diff line change
Expand Up @@ -371,6 +371,7 @@ below; these are estimates for planning a port, not results.
| glm (C) | cold first token −6..−7 % at TP=2, −7..−8 % at TP=3 | [#1823] |
| flashnext (M) | 1,024-token reply beside a 64K prompt 48 → 8 s, prompt's first token 90 → 54 s | [#2167] |
| flashnext (C), one GPU | prompts 15–23 % faster at `--parallel 1`, up to 8,192 rows a pass | [#2246] |
| glm (C), TP=2 | prompts about 2 % sooner at 8K–128K, two chunks a step, replies unchanged | [#1423] |

`TF_GLM_CHUNK_SOLO=4096` is slower with GLM's NVFP4 checkpoint on two GB10s ([#2288]).

Expand Down Expand Up @@ -460,6 +461,7 @@ below; these are estimates for planning a port, not results.
[#1394]: https://github.com/ashhart/TensorFold/pull/1394
[#1411]: https://github.com/ashhart/TensorFold/pull/1411
[#1413]: https://github.com/ashhart/TensorFold/pull/1413
[#1423]: https://github.com/ashhart/TensorFold/pull/1423
[#1426]: https://github.com/ashhart/TensorFold/pull/1426
[#1434]: https://github.com/ashhart/TensorFold/pull/1434
[#1439]: https://github.com/ashhart/TensorFold/pull/1439
Expand Down
3 changes: 2 additions & 1 deletion docs/recipes/glm-5.3-flash-cuda.md
Original file line number Diff line number Diff line change
Expand Up @@ -100,7 +100,8 @@ for each pass to end. A pass that takes more rows reads each routed expert's wei
comes sooner. `TF_GLM_CHUNK_SOLO=4096` lets a pass take two of those chunks when no other stream is decoding as it
starts. With one decoding, or with `prompt_logprobs` or image rows, the pass takes 2,048. Beside `--drafter`, DFlash2
reads a solo pass's taps whole and keeps the prompt's last 2,048 rows as before; its tap buffer holds 4,096 rows, 80 MiB
more on each rank.
more on each rank. Solo passes take the place of `TF_GLM_PROMPT_PAIR`'s paired steps, whose second set of buffers they
do not hold.

The grid does not move, so warm turns resume where they do unset, and streams decoding beside a prompt see the same
passes. Replies are the same bytes either way. Two costs remain:
Expand Down
23 changes: 23 additions & 0 deletions docs/recipes/moss-image-cuda.md
Original file line number Diff line number Diff line change
Expand Up @@ -77,6 +77,16 @@ python3 tools/zig/moss_image_case.py compare case/reference.png case/zig.png cas

The case tool needs diffusers 0.38 or later: the checkpoint's decoder config names `decoder_block_out_channels`.

`tf-moss-image` gates each whole forward element by element against the case's float32 run, and
`tf-moss-image-blocks` gates each block alone against the farther of torch bf16's two runs, both as on Metal
([moss-image.md](moss-image.md), "Run a case"). The block cases measured below were made on a Mac by
`moss_image_case.py blocks` (torch on MPS and on the CPU); a block case runs the whole transformer in float32, which
does not fit an RTX 5090.

```sh
zig-out/bin/tf-moss-image-blocks MODEL_DIR/transformer case/blocks-23-negative.safetensors
```

## Measured

Setup:
Expand Down Expand Up @@ -114,6 +124,19 @@ velocities and float32 runs as recorded on the GB10):
- **Image:** 44.01 dB PSNR from the pipeline's picture for the fox, 34.51 dB for the sticky note. The latents differ
from the GB10's in their low bits: cuBLASLt picks different GEMM kernels on each GPU.
- **Repeat:** identical bytes, as on the GB10.
- **Each block alone**, on the Metal doc's seven block cases (made on the Mac: torch bf16 on MPS and on the CPU): all
305 gates pass, at a mean of 0.50 of the limit and at most 0.60, and a second run repeats the first byte for byte.
At block 0 of the fox's step 1, negative branch, ours is 5.70e-6 from float32 in cosine error and 0.00960 in its
worst row (the negative prompt's seventh row), as torch bf16 on the CPU is (6.58e-6, 0.00960); torch on MPS is
2.28e-6 and 0.00687. Each of our operations there, run from our own inputs, matches torch's bf16 code to within the
order of its sums; that row turns small rounding differences into larger ones, on torch's devices as on ours.
- **Each whole forward**, over each case's 100: our cosine error from float32 is 1.00 times torch bf16's on average
for the fox (0.87 to 1.13) and 1.01 times for the sticky note (0.85 to 1.55); our worst row 1.01 times (0.34 to 1.63)
and 1.02 times (0.66 to 1.98).
- **The teacher gate** (`tf-moss-image`, exit 1, its bound from the GB10's torch bf16 run) fails all 100 branch-steps
of both cases, as on Metal ([moss-image.md](moss-image.md), Measured, where torch's own bf16 run on the CPU fails it
too). Worst element over its bound, min / median / max: fox 161 / 432 / 2946, sticky note 115 / 459 / 7921. Two
runs give identical latents.
- **Speed**, the GPU otherwise idle:
- torch bf16 (torch 2.13.0, the pipeline's own forward on this card): 0.307 s a forward for the fox, 0.309 s for the
sticky note.
Expand Down
Loading
Loading